[BUG]: Unqualified `warpReduce` in `warpspeed` scan kernel causes ADL collision

[viclafargue]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

CUB

Describe the bug

The warpspeed scan kernel (kernel_scan_warpspeed.cuh) calls warpReduce() unqualified. When ScanOpT is a type from an external namespace that also defines a warpReduce function template, ADL pulls that function into the overload set and causes an ambiguity error.

This affects any downstream code that passes a custom scan operator to thrust::inclusive_scan / cub::DeviceScan when compiling for sm_100+ (where the warpspeed kernel is selected).

Fix: qualify the call as ::cub::detail::scan::warpReduce to prevent ADL.

How to Reproduce

Minimal reproducer: pass a scan operator from a namespace that also contains a two-argument warpReduce function template. In practice, this is triggered by RAFT's raft::add_op + raft::warpReduce.

See rapidsai/cuvs#1954.

Expected behavior

Successful compilation.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[github-actions]

Hi @viclafargue!

Thanks for submitting this issue - the CCCL team has been notified and we'll get back to you as soon as we can!
In the mean time, feel free to add any relevant information to this issue.

[bernhardmgruber]

Thanks for the bug report! I think we should consider again whether to fully qualify all our internal functions. @miscco.

Fix: qualify the call as ::cub::detail::scan::warpReduce to prevent ADL.

Should be CUB_NS_QUALIFIER::detail::scan::warpReduce to allow for user-defined namespaces wrapping the CUB namespace.

Can you open a PR please? thx!

[viclafargue]

Sure, here is the PR : #8236.