[tpn]

NVRTC 12.9 and 13.1 reject the generic lambdas introduced by #12195 as unannotated host functions. This breaks C Parallel and Python reduction, partition, and unique-by-key tests, including CI on #12296.

Use int step parameters in both cuda::static_for lambdas. The existing captures and reduction operations are preserved.

Validation

• Full-warp and partial-warp double-sum reproducer: the original fails NVRTC 12.9/13.1; the fix compiles in C++17 and C++20.
• NVCC 13.1 compiles both versions. All emitted kernels return the expected sums on an RTX PRO 6000 Blackwell GPU.
• CI SASS comparison: no changes across 100 CUB and 75 Thrust benchmark targets on SM75/80/90/100/110/120 (1,050 comparisons).
• The two local kernels also produce byte-identical SASS and resource dumps with NVCC 13.1/C++20 on SM75.
• Targeted pre-commit passes. Existing C Parallel/Python tests cover the failing paths; full CI remains pending.

The local NVRTC 12.9 check used CUDA 12.8 headers; the 13.1 check used matching headers. Local execution covers two kernels, not the full type/operator matrix or performance equivalence.

CI status

Full CI is still running. The CUDA 13.4 Python examples job passes all 70 examples, then fails all 125 benchmark fixture setups: pytest-benchmark 5.3.0 accesses Function._nodeid, which is absent in pytest 9.2.0. This occurs before benchmark execution and is unrelated to the CUB change.

Checklist

• Existing tests cover these changes.
• The documentation is up to date with these changes.

[copy-pr-bot]

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

[coderabbitai]

Important

Review skipped

Review was skipped as selected files did not have any reviewable changes.

Run configuration

• Source mode: Auto
• Repository: NVIDIA/cccl
• Branch: main
• File: .coderabbit.yaml
• Commit: 713d911
• Layers: Repository YAML
• Review profile: CHILL
• Plan: Enterprise
• Run ID: cb6fbe86-bd97-4f25-bbd9-ce60fa72cf33

Commits

Reviewing files that changed from the base of the PR and between cd63613 and 35095a1.

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

• 🔍 Trigger review

No actionable comments were generated in the recent review.

Recent review info

Run configuration

• Source mode: Auto
• Repository: NVIDIA/cccl
• Branch: main
• File: .coderabbit.yaml
• Commit: 215868f
• Layers: Repository YAML
• Review profile: CHILL
• Plan: Enterprise
• Run ID: daab30e3-e28b-4912-abcd-d9aec83e34fd

Commits

Reviewing files that changed from the base of the PR and between 215868f and cd63613.

Files selected for processing: 1

• cub/cub/warp/specializations/warp_reduce_shfl.cuh

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.

📝 Summary

Summary by CodeRabbit

• Refactor
  • Updated internal reduction-step parameter declarations; no user-visible behavior changes.

The ShuffleReduce and FullShuffleReduce callbacks in WarpReduceShfl now declare their step parameter as int instead of auto.

Changes: WarpReduce shuffle reductions

Layer / File(s)	Summary
Update reduction callback parameter types cub/cub/warp/specializations/warp_reduce_shfl.cuh	The static_for callbacks in ShuffleReduce and FullShuffleReduce now declare step as int instead of auto.

Possibly related PRs

• NVIDIA/cccl#12195: This PR introduced the cuda::static_for reduction lambdas in the same file; this change updates their step parameter types.

Suggested reviewers: fbusato, gevtushenko

Priority: Normal

Change: Bug fix

Merge Risk: Minimal · previous review cd636

No actionable merge-blocking risk is identified in this focused NVRTC compatibility change.

- [ ] Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

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

[fbusato]

waiting for CI and SASS check. auto here represents an integral_constant not a run-time value

[pylaterreur]

Some data points from reproducing this without a GPU (pip NVRTC wheels, CCCL main f1b40ea):

• Version boundary. Without -default-device, NVRTC 12.0.140, 12.4.127, 12.8.93, 12.9.86,
  13.0.88 and 13.1.115 reject any generic lambda in device code. This holds even when it is called
  directly in a __global__ function with no cuda::static_for involved:
  auto f = [](auto x) { return x + 1; }; *out = f(41);. NVRTC 13.2.86, 13.3.33 and 13.4.92
  accept it. C++20 template lambdas fail the same way. __device__-annotated lambdas need
  --extended-lambda, which NVRTC does not accept. So no generic-lambda form works before 13.2;
  the alternatives are int or a named functor with a templated __device__ operator() (that
  compiles with 12.9).

• int vs auto. Both uses consume the step as a runtime int:
  ReduceStep(..., const int offset) and warp_shuffle_down(..., const int __delta, ...).
  Locally with nvcc 13.4, a TU instantiating both lambdas over several types, warp sizes and
  DeviceReduce::Sum gave byte-identical SASS for auto and int (69 kernels; sm_75, sm_89,
  sm_90).

• Why CI was green. PR CI runs every NVRTC-compiled test with CTK 13.4 only. These are
  cub.test.nvrtc in the CUB GPU jobs, plus C Parallel and Python through pull_request_lite.
  Nightly's CTK 12.X jobs (CUB, C Parallel and Python tests) should hit this until the fix lands.

• Test gap. cub.test.nvrtc already fails on main with NVRTC < 13.2, but only on the
  line-202 lambda (full-warp Sum). Adding a partial-warp Sum makes it reach the line-192
  lambda too (checked: errors at 202 and 192 on main with 12.9; compiles with this PR):

          __syncwarp();
  
          // Test partial warp reduce
          constexpr int valid_items = 16;
          const int partial_sum = warp_reduce_t(warp_reduce_storage).Sum(items[0], valid_items);
          if (threadIdx.x == 0)
          {
            if (partial_sum != (valid_items * (valid_items - 1) / 2))
            {
              atomicAdd(errors, 1);
            }
          }

• Pre-existing instance, not from [CUB] Modernize/Improve warp_reduce_shfl.cuh #12195. warp_scan_smem.cuh:249 ([&](auto step) with
  if constexpr (step == 0)) makes cub::WarpScan<int, 7>::InclusiveScanPartial fail with
  NVRTC 12.9 at 3a1772c already.

🤖 Generated with Claude Code

https://claude.ai/code/session_01GwUvhmzgfkiZ6vo4uGyesN

[pylaterreur]

Ran the int version of both lambdas on a GPU (Tesla T4, sm_75, NVRTC 12.9.86, with #12314's changes): cub.test.nvrtc's kernel reports errors == 0. A kernel with WarpReduce<__nv_bfloat16>::Sum (full warp), WarpReduce<int>::Sum (full warp) and WarpReduce<int>::Sum(v, 16) returns the expected sums (31, 31 and 15 for inputs i % 3). On the SASS question: nvcc 13.4 gives byte-identical SASS for auto and int (69 kernels, sm_75/89/90), see my earlier comment.

🤖 Generated with Claude Code

https://claude.ai/code/session_01GwUvhmzgfkiZ6vo4uGyesN

[pauleonix]

Consider keeping auto and adding _CCCL_DEVICE to the lambda instead. #12262 contains that fix as well.

[pylaterreur]

@pauleonix I tried that on NVRTC 12.9.86 (sm_89, no GPU needed) and it doesn't compile there:

• [&] _CCCL_DEVICE (auto step) { ... } (also with a plain __device__) in a kernel or in a __device__ function: error: __host__ or __device__ annotation on lambda requires --extended-lambda nvcc flag.
• NVRTC 12.9 rejects the flag itself (nvrtc: error: unrecognized option --extended-lambda found), and adding -default-device doesn't change the error.
• Without the annotation, the generic lambda fails with A function without execution space annotations ... is considered a host function (the CI error in [BUG]: cub::WarpScan::InclusiveScanPartial fails to compile with NVRTC < 13.2 (generic lambda in warp_scan_smem.cuh) #12312 / numba-cuda-mlir#358). Only NVRTC 13.2 and later accept it.

So keeping auto and annotating the lambda only works on 13.2+. For warp_reduce_shfl.cuh the step is used as a runtime int by both consumers, so const int works on every NVRTC version; for warp_scan_smem.cuh, where the step is used in constexpr/if constexpr, #12314 does the same with a runtime if (identical SASS).

The diff of #12262 (clang-tidy plugins, 11 files) doesn't touch these headers; is there another PR number with the fix you mean? If it uses a mechanism that works with NVRTC 12.x, I'd like to look at it.

🤖 Generated with Claude Code

https://claude.ai/code/session_01GwUvhmzgfkiZ6vo4uGyesN

[pauleonix]

@pylaterreur Yes, sorry, I was too quick to believe it would work.

[fbusato]

no SASS differences. Approved

[jrhemstad]

/ok to test 35095a1

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 4h 36m: Pass: 100%/210 | Total: 9d 20h | Max: 3h 28m | Hits: 27%/1024301

See results here.