fix(warp_reduce): add explicit add_op overload to resolve CUB template ambiguity on CUDA 13.2+ by zbrad · Pull Request #3050 · NVIDIA/raft

zbrad · 2026-06-08T11:23:30Z

Summary

On CUDA 13.2 (SM 121, DGX Spark), IVF-PQ builds fail with an ambiguous template instantiation error. When CUB scan kernels call warpReduce(val, raft::add_op{}), both raft::warpReduce<T, ReduceLambda> and cub::detail::scan::warpReduce<Tp, ScanOpT&> match, producing a compile error.

Fix: Add an explicit non-template overload in cpp/include/raft/util/reduction.cuh:

template <typename T>
DI T warpReduce(T val, raft::add_op reduce_op)

The explicit overload is preferred by the compiler over the generic ReduceLambda overload, resolving the ambiguity without changing any existing behavior.

cpp/include/raft/util/reduction.cuh — explicit raft::add_op overload for warpReduce
cpp/tests/util/reduction.cu — regression test (WARP_REDUCE_WITH_ADD_OP) to prevent future regressions

Repro / context

Observed on DGX Spark (SM 121) with CUDA 13.2. The upstream CI does not test CUDA 13.2 / SM 121; the new test compiles and passes on all CUDA versions but will catch regressions if CUDA 13.2 support is added to CI.

Error seen without fix:

error: more than one instance of overloaded function "warpReduce" matches the argument list

Test plan

WARP_REDUCE_WITH_ADD_OP regression test added in cpp/tests/util/reduction.cu
Verified fix compiles and tests pass on CUDA 13.2 (SM 121, DGX Spark)
No changes to existing warpReduce behavior — explicit overload only activates for raft::add_op

🤖 Generated with Claude Code

…e ambiguity on CUDA 13.2 On CUDA 13.2 (SM 121, DGX Spark), IVF-PQ builds fail because both raft::warpReduce<T, ReduceLambda> and cub::detail::scan::warpReduce<Tp, ScanOpT&> match when called with raft::add_op{}, causing an ambiguous template instantiation. Add an explicit non-template overload DI T warpReduce(T val, raft::add_op reduce_op) in reduction.cuh. The explicit overload is preferred by the compiler, resolving ambiguity. Also added a regression test WARP_REDUCE_WITH_ADD_OP to prevent future regressions.

copy-pr-bot · 2026-06-08T11:23:33Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

achirkin

Thank you for reporting the issue! Could you please provide the reference/example where the bug is triggered? Normally, I'd assume the fix should be on the user side - just call the function with explicit namespace.

achirkin · 2026-06-08T13:27:12Z

+{
+  assert(gridDim.x == 1);
+  int th_val = input[threadIdx.x];
+  th_val     = raft::warpReduce(th_val, raft::add_op{});


Does this really cause an ambiguity without the extra overload in util/reduction.cuh? It's called here with raft namespace, so I doubt CUB overload is ever picked up here.

It showed up for me when trying to do a full source rebuild of cuvs on the dgx spark. When finally debugging my build failure, it traced down to raft, but only showed up when building for arm64.

Oh, I have the cuvs regression test for it that I'm submitting to cuvs, it's at cpp/tests/regression/warp_reduce_add_op.cu

divyegala · 2026-06-08T16:40:02Z

@zbrad can you provide steps to reproduce this bug? I have a DGX Spark, and I have been building CUDA 13.2 without any failures.

zbrad · 2026-06-08T19:43:29Z

some other folks had asked for more background, so I went back and re-created the original failure from compiling cuvs. I've attached the doc and the repro.
warp reduce ambiguity doc
repro

achirkin · 2026-06-09T07:00:52Z

Thanks for the reproducer documentation! I feel uneasy about both suggested workarounds:

raft extra overload: it doesn't protect us from someone else failing exactly the same way on another device operation.
cuvs swap to thrust::plus: because we have a best practice guideline to use the raft operations where possible.

Maybe we could make the raft's warpReduce template itself a bit more restrictive so it wouldn't get in the way of thrust+cub primitives?..

Copilot

Pull request overview

This PR addresses a CUDA 13.2+ (SM 121) compilation failure caused by ambiguous warpReduce overload resolution when CUB scan code interacts with raft::warpReduce(val, raft::add_op{}). It does so by adding a more-specific overload for raft::add_op, and adds a regression test to ensure this call remains unambiguous going forward.

Changes:

Add an explicit warpReduce(T, raft::add_op) overload to disambiguate overload resolution on CUDA 13.2+.
Add a regression kernel + gtest case to compile and validate warpReduce(val, raft::add_op{}) usage.
Update .gitignore to ignore build logs and aarch64 build artifacts.

Reviewed changes

Copilot reviewed 2 out of 3 changed files in this pull request and generated no comments.

File	Description
cpp/include/raft/util/reduction.cuh	Adds an explicit `raft::add_op` overload for `warpReduce` to avoid CUDA 13.2+ ambiguity with CUB.
cpp/tests/util/reduction.cu	Adds a regression test that exercises `warpReduce(val, raft::add_op{})` in device code.
.gitignore	Ignores build log files and aarch64 build output artifacts.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

coderabbitai · 2026-07-25T20:13:24Z

📝 Walkthrough

Summary by CodeRabbit

Bug Fixes
- Improved warp-level reduction compatibility with newer CUDA toolchains when using addition operations.
- Prevented potential ambiguity with CUDA’s built-in reduction utilities.
Tests
- Added coverage verifying correct warp-level sum results across supported data types and configurations.
Chores
- Updated ignore rules for build logs and generated build artifacts.

Walkthrough

Adds an explicit raft::add_op overload for warpReduce, introduces CUDA regression coverage for that overload, and updates .gitignore for build logs and packaged artifacts.

Changes

Warp reduction overload

Layer / File(s)	Summary
Add-operation warpReduce overload `cpp/include/raft/util/reduction.cuh`	Adds a `raft::add_op` overload that forwards to `logicalWarpReduce` to avoid CUDA 13.2+ CUB overload ambiguity.
Add-operation regression tests `cpp/tests/util/reduction.cu`	Adds a CUDA kernel, launch helper, fixture method, and parameterized test covering `warpReduce` with `raft::add_op{}`.

Build artifact ignore rules

Layer / File(s)	Summary
Ignore generated build artifacts `.gitignore`	Adds ignore patterns for build logs, aarch64 build output, and `.tar.bz2` packages.

Estimated code review effort: 2 (Simple) | ~10 minutes

Suggested reviewers: copilot

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly summarizes the main change: adding an explicit `raft::add_op` overload to fix CUDA 13.2 template ambiguity.
Description check	✅ Passed	The description accurately explains the same overload fix and the added regression test, matching the changeset.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

✨ Finishing Touches

🧪 Generate unit tests (beta)

Create PR with unit tests

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

_{Comment @coderabbitai help to get the list of available commands.}

coderabbitai

Actionable comments posted: 1

🤖 Prompt for all review comments with AI agents

Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In @.gitignore:
- Line 2: Update the build log ignore pattern in .gitignore from build_*.log so
it matches the CI-generated telemetry-artifacts/build.log, using either
build*.log or the more explicit telemetry-artifacts/*.log pattern.

🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

Push a commit to this branch (recommended)
Create a new PR with the fixes

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: 2f2a9068-a215-4b39-a05f-66e896c967d6

📥 Commits

Reviewing files that changed from the base of the PR and between aa9a6d5 and 9fdbdf9.

📒 Files selected for processing (3)

.gitignore
cpp/include/raft/util/reduction.cuh
cpp/tests/util/reduction.cu

coderabbitai · 2026-07-25T20:13:27Z

@@ -1,3 +1,10 @@
+## build logs
+build_*.log


🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

Match the build log generated by CI.

.github/workflows/pr.yaml:401-402 writes telemetry-artifacts/build.log, but build_*.log only matches filenames beginning with build_. Use build*.log or the more explicit telemetry-artifacts/*.log.

Proposed fix

-build_*.log +build*.log

📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change

build_*.log

build*.log

🤖 Prompt for AI Agents

Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In @.gitignore at line 2, Update the build log ignore pattern in .gitignore from build_*.log so it matches the CI-generated telemetry-artifacts/build.log, using either build*.log or the more explicit telemetry-artifacts/*.log pattern.

zbrad requested a review from a team as a code owner June 8, 2026 11:23

github-project-automation Bot added this to Unstructured Data Processing Jun 8, 2026

filter artifacts

23bb5ac

achirkin requested changes Jun 8, 2026

View reviewed changes

josephine-wolf-oberholtzer added this to Unstructured Data Processing Jul 1, 2026

aamijar assigned zbrad Jul 23, 2026

Merge branch 'main' into fix/warp-reduce-cub-ambiguity

9fdbdf9

Copilot AI review requested due to automatic review settings July 25, 2026 20:09

zbrad requested a review from a team as a code owner July 25, 2026 20:09

Copilot started reviewing on behalf of zbrad July 25, 2026 20:10 View session

Copilot AI reviewed Jul 25, 2026

View reviewed changes

coderabbitai Bot reviewed Jul 25, 2026

View reviewed changes

Uh oh!

Conversation

zbrad commented Jun 8, 2026

Summary

Repro / context

Test plan

Uh oh!

copy-pr-bot Bot commented Jun 8, 2026

Uh oh!

achirkin left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

achirkin Jun 8, 2026

Choose a reason for hiding this comment

Uh oh!

zbrad Jun 8, 2026

Choose a reason for hiding this comment

Uh oh!

zbrad Jun 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

divyegala commented Jun 8, 2026

Uh oh!

zbrad commented Jun 8, 2026

Uh oh!

achirkin commented Jun 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

coderabbitai Bot commented Jul 25, 2026

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot Jul 25, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

achirkin left a comment •

edited

Loading

zbrad Jun 8, 2026 •

edited

Loading

achirkin commented Jun 9, 2026 •

edited

Loading