Skip to content

Fix env-based DeviceReduce kernel-allowlist tests on MSVC + CUDA 13.3#9726

Open
edenfunf wants to merge 2 commits into
NVIDIA:mainfrom
edenfunf:fix/9643-env-kernel-allowlist
Open

Fix env-based DeviceReduce kernel-allowlist tests on MSVC + CUDA 13.3#9726
edenfunf wants to merge 2 commits into
NVIDIA:mainfrom
edenfunf:fix/9643-env-kernel-allowlist

Conversation

@edenfunf

@edenfunf edenfunf commented Jul 7, 2026

Copy link
Copy Markdown
Contributor

Description

closes #9643

The env-based DeviceReduce tests compute their expected kernel pointers from decltype(d_in) inside the [&] allowlist lambdas. With nvcc 13.3 and MSVC as host compiler, decltype() of a variable captured by reference inside a lambda is evaluated as a reference type (a static_assert(!is_reference_v<decltype(d_in)>) inside the lambda fires on 13.3 and not on 12.9). The tests therefore reference Kernel<..., constant_iterator&, ...> while the dispatch launches Kernel<..., constant_iterator, ...> — a different instantiation with its own host stub that is never registered with the runtime (cudaFuncGetName on the allowlist pointer returns cudaErrorInvalidDeviceFunction), so the pointer check at catch2_test_env_launch_helper.h:97 fails. Which lambda is affected shifts with unrelated code changes, hence the instantiation-specific failures in the issue.

Changes:

  • Build the allowlists in a function template (expected_reduce_kernels) shared by both requirements-parametrized tests. Iterator types are now deduced at test scope, so no decltype() of a capture remains, and the two near-identical per-test lambdas are gone.
  • Take the kernel pointers from the same kernel source types the dispatch uses instead of re-spelling the kernel template arguments, so the expected instantiations cannot drift from the launched ones. The RFA dispatch had no kernel source, so this adds DeterministicDeviceReduceKernelSource plus a default_kernel_source_t alias to dispatch_reduce_deterministic.cuh, following the pattern of the other dispatch layers; invoke_single_tile/invoke_passes take the kernels as parameters like DispatchReduce does. No functional change.
  • Document the pitfall next to allowed_kernels() in the launch helper.

Verified on Windows / MSVC 14.50 / C++17 / sm_120:

  • CUDA 13.3.1: cub.test.device.reduce_env.lid_0 reproduced the reported failure before the change and passes after. reduce_env.lid_2, reduce_deterministic, reduce_nondeterministic, and reduce_env_api also pass.
  • CUDA 12.9: reduce_env.lid_0, reduce_deterministic, reduce_nondeterministic pass (regression check).

The decltype() misevaluation itself looks like a compiler regression worth tracking on the nvcc side. A trivial repro does not trigger it; it seems to need the surrounding template test function and a class-type iterator.

Checklist

  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

edenfunf added 2 commits July 7, 2026 16:50
The environment-based DeviceReduce tests computed their expected kernel
pointers with `decltype(d_in)` inside the `[&]` allowlist lambda. With
nvcc 13.3 and MSVC as host compiler, decltype() of a variable captured
by reference inside a lambda is misevaluated as a reference type, so the
tests referenced (and instantiated) `Kernel<..., InputIt&, ...>` instead
of the `Kernel<..., InputIt, ...>` the dispatch launches. The two are
different instantiations with different host stubs, and the referenced
one is never registered with the CUDA runtime, so the white-box
kernel-allowlist check at catch2_test_env_launch_helper.h:97 failed.

* Hoist the iterator type aliases out of the allowlist lambdas.
* Build the allowlists from the same kernel-source types the dispatch
  uses, so the expected instantiations cannot drift from the launched
  ones.
* Add DeterministicDeviceReduceKernelSource to the RFA dispatch,
  mirroring the KernelSource pattern used by the other dispatch layers,
  and pass the kernels into invoke_single_tile/invoke_passes as
  parameters.

Verified on Windows with MSVC 14.50: CUDA 13.3.1 (previously failing)
and CUDA 12.9 (regression check) both pass all 49 test cases.

Fixes NVIDIA#9643
* Factor the near-identical allowlist lambdas of the two
  requirements-parametrized tests into one expected_reduce_kernels
  function template. As a function template, the iterator types are
  deduced at test scope, which also removes the lambda that triggered
  the decltype() misevaluation entirely.
* Expose detail::rfa::default_kernel_source_t and use it both as the
  dispatch default and in the test, so the expected RFA kernel
  instantiations are derived from a single place.
* Include <thrust/type_traits/unwrap_contiguous_iterator.h> directly
  instead of relying on transitive inclusion.
* Document the decltype()-inside-lambda pitfall next to
  allowed_kernels() in the launch helper, where future allowlist tests
  will look.
* Drop the redundant reduction_op_t alias and the stale commented-out
  UNSCOPED_CAPTURE block.

Re-verified on Windows with MSVC 14.50 on both CUDA 13.3.1 and 12.9:
reduce_env (lid_0 + lid_2), reduce_deterministic, reduce_nondeterministic,
and reduce_env_api all pass.
@edenfunf edenfunf requested a review from a team as a code owner July 7, 2026 09:40
@edenfunf edenfunf requested a review from bernhardmgruber July 7, 2026 09:40
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Jul 7, 2026
@copy-pr-bot

copy-pr-bot Bot commented Jul 7, 2026

Copy link
Copy Markdown
Contributor

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.

@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jul 7, 2026
@coderabbitai

coderabbitai Bot commented Jul 7, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Summary by CodeRabbit

  • New Features

    • Improved deterministic device reduce handling with a more flexible kernel-launch setup.
    • Expanded test coverage to better validate expected kernel usage across deterministic and non-deterministic paths.
  • Documentation

    • Added guidance for identifying the exact kernel variants launched during dispatch.
  • Bug Fixes

    • Simplified and stabilized reduce-related test expectations for environment-based execution.

Walkthrough

Compact metadata
This PR refactors deterministic device reduce dispatch (dispatch_reduce_deterministic.cuh) to route kernel launches through a KernelSource abstraction (DeterministicDeviceReduceKernelSource, default_kernel_source_t), replacing inline kernel template instantiations. Related tests are updated to derive expected kernel allowlists from this abstraction.

Changes

Kernel Source Refactor

Layer / File(s) Summary
Kernel source struct and includes
cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh
New includes and DeterministicDeviceReduceKernelSource/default_kernel_source_t bundle single-tile, multi-tile, and second single-tile kernel getters.
invoke_single_tile and invoke_passes use getters
cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh
Kernel launches call .doit(...) on provided getter objects; reduce_config.__init uses reduce_kernel getter instead of inline template types.
dispatch signature gains KernelSource
cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh
dispatch adds a KernelSource template parameter (default default_kernel_source_t<...>) and kernel_source parameter, wiring kernel_source.SingleTileKernel()/ReductionKernel()/SingleTileSecondKernel() into launch paths.
Test allowlists derived from kernel source
cub/test/catch2_test_device_reduce_env.cu, cub/test/catch2_test_env_launch_helper.h
New expected_reduce_kernels helper builds allowlists via DeviceReduceKernelSource/rfa::default_kernel_source_t per determinism case; environment and not_guaranteed fallback tests use it; helper doc comment clarifies allowlist computation caveats.

Assessment against linked issues

Objective Addressed Explanation
Fix kernel-allowlist mismatch on MSVC + CUDA 13.3 by deriving expected kernel pointers from the actual kernel source/config instead of recomputing separately (#9643)
Cover run_to_run, gpu_to_gpu (RFA), and not_guaranteed determinism cases consistently (#9643)

Possibly related PRs

  • NVIDIA/cccl#8863: Both PRs modify the same deterministic reduce dispatch/invoke kernel launch path in dispatch_reduce_deterministic.cuh.

Suggested reviewers: bernhardmgruber, griwes, srinivasyadav18

important: verify the KernelSource default template argument doesn't silently change ABI/behavior for existing callers that don't pass kernel_source explicitly.

suggestion: the doc comment addition in catch2_test_env_launch_helper.h should reference the specific decltype-by-reference pitfall example to make it actionable for future test authors.


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

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🧹 Nitpick comments (1)
cub/test/catch2_test_device_reduce_env.cu (1)

505-520: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

suggestion: The sum variant (Lines 550-565) is byte-identical to this block, and both mirror the run_to_run branch of expected_reduce_kernels. Extract a small helper so future kernel-source changes can't drift between the two not_guaranteed tests. Optional.


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 82c18609-29b3-482f-a43b-f810cba21503

📥 Commits

Reviewing files that changed from the base of the PR and between 886c291 and eb895af.

📒 Files selected for processing (3)
  • cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh
  • cub/test/catch2_test_device_reduce_env.cu
  • cub/test/catch2_test_env_launch_helper.h

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Status: In Review

Development

Successfully merging this pull request may close these issues.

[BUG]: CUB environment-based DeviceReduce tests fail kernel-allowlist check on MSVC + CUDA 13.3

1 participant