From 7a2efcc130cf48c7506e6616438ef08385238291 Mon Sep 17 00:00:00 2001 From: edenfunf Date: Tue, 7 Jul 2026 16:50:19 +0800 Subject: [PATCH 1/2] Fix CUB reduce env kernel-allowlist tests on MSVC + CUDA 13.3 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 #9643 --- .../dispatch_reduce_deterministic.cuh | 115 ++++--- cub/test/catch2_test_device_reduce_env.cu | 301 +++++++----------- 2 files changed, 179 insertions(+), 237 deletions(-) diff --git a/cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh b/cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh index 31292948a98..6eeab543919 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh @@ -24,6 +24,7 @@ #include #include #include +#include #include #include @@ -35,6 +36,7 @@ #include #include #include +#include #include #include @@ -84,13 +86,53 @@ struct deterministic_sum_t template +struct DeterministicDeviceReduceKernelSource +{ + // PolicySelector must be stateless, so we can pass the type to the kernel + static_assert(::cuda::std::is_empty_v); + + CUB_DEFINE_KERNEL_GETTER( + SingleTileKernel, + reduce::DeterministicDeviceReduceSingleTileKernel< + PolicySelector, + InputIteratorT, + OutputIteratorT, + ReductionOpT, + InitValueT, + DeterministicAccumT, + TransformOpT>) + + CUB_DEFINE_KERNEL_GETTER( + ReductionKernel, + reduce:: + DeterministicDeviceReduceKernel) + + // The second single-tile pass reduces the per-block partials and always uses an identity transform + CUB_DEFINE_KERNEL_GETTER( + SingleTileSecondKernel, + reduce::DeterministicDeviceReduceSingleTileKernel< + PolicySelector, + DeterministicAccumT*, + OutputIteratorT, + ReductionOpT, + InitValueT, + DeterministicAccumT>) +}; + +template CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t invoke_single_tile( + SingleTileKernelT single_tile_kernel, void* d_temp_storage, size_t& temp_storage_bytes, InputIteratorT d_in, @@ -122,20 +164,7 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t invok // Invoke single_reduce_sweep_kernel if (const auto error = CubDebug( launcher_factory(1, active_policy.single_tile.threads_per_block, 0, stream) - .doit(detail::reduce::DeterministicDeviceReduceSingleTileKernel< - PolicySelector, - InputIteratorT, - OutputIteratorT, - ReductionOpT, - InitValueT, - DeterministicAccumT, - TransformOpT>, - d_in, - d_out, - static_cast(num_items), - reduction_op, - init, - transform_op))) + .doit(single_tile_kernel, d_in, d_out, static_cast(num_items), reduction_op, init, transform_op))) { return error; } @@ -150,16 +179,19 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t invok return CubDebug(detail::DebugSyncStream(stream)); } -template CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t invoke_passes( + ReduceKernelT reduce_kernel, + SingleTileKernelT single_tile_second_kernel, void* d_temp_storage, size_t& temp_storage_bytes, InputIteratorT d_in, @@ -179,10 +211,7 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t invok } KernelConfig reduce_config; - if (const auto error = CubDebug(reduce_config.__init( - detail::reduce:: - DeterministicDeviceReduceKernel, - active_policy.multi_tile))) + if (const auto error = CubDebug(reduce_config.__init(reduce_kernel, active_policy.multi_tile))) { return error; } @@ -249,11 +278,7 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t invok if (const auto error = CubDebug( launcher_factory(current_grid_size, active_policy.multi_tile.threads_per_block, 0, stream) - .doit(detail::reduce::DeterministicDeviceReduceKernel, + .doit(reduce_kernel, d_in, d_chunk_block_reductions, num_current_items, @@ -295,13 +320,7 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t invok // Invoke DeterministicDeviceReduceSingleTileKernel if (const auto error = CubDebug( launcher_factory(1, active_policy.single_tile.threads_per_block, 0, stream) - .doit(detail::reduce::DeterministicDeviceReduceSingleTileKernel< - PolicySelector, - DeterministicAccumT*, - OutputIteratorT, - ReductionOpT, - InitValueT, - DeterministicAccumT>, + .doit(single_tile_second_kernel, d_block_reductions, d_out, reduce_grid_size, @@ -330,6 +349,14 @@ template , typename PolicySelector = policy_selector_from_types, __determinism_t::__gpu_to_gpu>, + typename KernelSource = DeterministicDeviceReduceKernelSource< + PolicySelector, + THRUST_NS_QUALIFIER::try_unwrap_contiguous_iterator_t, + OutputIteratorT, + deterministic_sum_t, + InitValueT, + typename deterministic_sum_t::DeterministicAcc, + TransformOpT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( void* d_temp_storage, @@ -341,6 +368,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( cudaStream_t stream = {}, TransformOpT transform_op = {}, PolicySelector policy_selector = {}, + KernelSource kernel_source = {}, KernelLauncherFactory launcher_factory = {}) { // Get CC @@ -373,14 +401,8 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( if (num_items <= tile_items) { - return invoke_single_tile( + return invoke_single_tile( + kernel_source.SingleTileKernel(), d_temp_storage, temp_storage_bytes, d_in_unwrapped, @@ -394,14 +416,9 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( launcher_factory); } - return invoke_passes( + return invoke_passes( + kernel_source.ReductionKernel(), + kernel_source.SingleTileSecondKernel(), d_temp_storage, temp_storage_bytes, d_in_unwrapped, diff --git a/cub/test/catch2_test_device_reduce_env.cu b/cub/test/catch2_test_device_reduce_env.cu index 5df99e17abc..5fc19538973 100644 --- a/cub/test/catch2_test_device_reduce_env.cu +++ b/cub/test/catch2_test_device_reduce_env.cu @@ -10,6 +10,7 @@ struct stream_registry_factory_t; #include #include +#include #include #include @@ -346,6 +347,13 @@ C2H_TEST("Device reduce uses environment", "[reduce][device]", requirements) auto d_in = cuda::constant_iterator(1.0f); auto d_out = thrust::device_vector(1); + // Compute the iterator types outside of the allowlist lambda below: with nvcc 13.3 on MSVC, decltype() of a + // variable captured by reference inside a lambda is misevaluated as a reference type, making the test reference + // (and instantiate) different kernels than the ones the dispatch launches. + // See: https://github.com/NVIDIA/cccl/issues/9643 + using input_it_t = decltype(d_in); + using output_it_t = decltype(d_out.begin()); + init_value_t init = 0; size_t expected_bytes_allocated{}; @@ -357,38 +365,15 @@ C2H_TEST("Device reduce uses environment", "[reduce][device]", requirements) cudaSuccess == cub::DeviceReduce::Reduce(nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, op_t{}, init)); - using policy_t = cub::detail::reduce::policy_selector_from_types; + // Build the allowlist from the same kernel source type the dispatch uses, so the expected kernel + // instantiations cannot drift from the launched ones. + using policy_t = cub::detail::reduce::policy_selector_from_types; + using kernel_source_t = cub::detail::reduce:: + DeviceReduceKernelSource; return cuda::std::array{ - reinterpret_cast( - cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - decltype(d_in), - decltype(d_out.begin()), - offset_t, - op_t, - init_value_t, - accumulator_t, - transform_t>), - reinterpret_cast( - cub::detail::reduce::DeviceReduceKernel< - policy_t, - /* StableReductionOrder */ true, - decltype(d_in), - accumulator_t*, - offset_t, - op_t, - accumulator_t, - init_value_t, - transform_t>), - reinterpret_cast( - cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - accumulator_t*, - decltype(d_out.begin()), - int, // always used with int offset - op_t, - init_value_t, - accumulator_t>)}; + reinterpret_cast(kernel_source_t::SingleTileKernel()), + reinterpret_cast(kernel_source_t::ReductionKernel()), + reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; } else if constexpr (cub::detail::is_non_deterministic_v) { @@ -409,17 +394,18 @@ C2H_TEST("Device reduce uses environment", "[reduce][device]", requirements) /* stream */ nullptr, transform_t{})); - return cuda::std::array{reinterpret_cast( - cub::detail::reduce::DeviceReduceKernel< - policy_t, - /* StableReductionOrder */ false, - decltype(d_in), - decltype(raw_ptr), - offset_t, - op_t, - accumulator_t, - init_value_t, - transform_t>)}; + // See the run_to_run branch above for why the kernel pointers must come from the kernel source + using kernel_source_t = cub::detail::reduce::DeviceReduceKernelSource< + policy_t, + input_it_t, + decltype(raw_ptr), + offset_t, + op_t, + init_value_t, + accumulator_t, + transform_t, + /* StableReductionOrder */ false>; + return cuda::std::array{reinterpret_cast(kernel_source_t::ReductionKernel())}; } else { @@ -428,36 +414,29 @@ C2H_TEST("Device reduce uses environment", "[reduce][device]", requirements) using policy_t = cub::detail::reduce:: policy_selector_from_types; using deterministic_accum_t = deterministic_add_t::DeterministicAcc; - using output_it_t = decltype(d_out.begin()); - REQUIRE(cudaSuccess - == cub::detail::rfa:: - dispatch( - nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, init)); + REQUIRE( + cudaSuccess + == cub::detail::rfa::dispatch( + nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, init)); - auto k1 = cub::detail::reduce::DeterministicDeviceReduceSingleTileKernel< - policy_t, - decltype(d_in), - output_it_t, - reduction_op_t, - init_value_t, - deterministic_accum_t, - transform_t>; - auto k2 = cub::detail::reduce:: - DeterministicDeviceReduceKernel; - auto k3 = cub::detail::reduce::DeterministicDeviceReduceSingleTileKernel< + // See the run_to_run branch above for why the kernel pointers must come from the kernel source + using kernel_source_t = cub::detail::rfa::DeterministicDeviceReduceKernelSource< policy_t, - deterministic_accum_t*, + thrust::try_unwrap_contiguous_iterator_t, output_it_t, reduction_op_t, init_value_t, deterministic_accum_t, transform_t>; // TODO(bgruber): enable this when we have Catch2 3.13+ - // UNSCOPED_CAPTURE(c2h::type_name(), c2h::type_name(), - // c2h::type_name()); + // UNSCOPED_CAPTURE(c2h::type_name(), + // c2h::type_name(), + // c2h::type_name()); return cuda::std::array{ - reinterpret_cast(k1), reinterpret_cast(k2), reinterpret_cast(k3)}; + reinterpret_cast(kernel_source_t::SingleTileKernel()), + reinterpret_cast(kernel_source_t::ReductionKernel()), + reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; } }(); @@ -486,6 +465,13 @@ C2H_TEST("Device sum uses environment", "[reduce][device]", requirements) auto d_in = cuda::constant_iterator(1.0f); auto d_out = thrust::device_vector(1); + // Compute the iterator types outside of the allowlist lambda below: with nvcc 13.3 on MSVC, decltype() of a + // variable captured by reference inside a lambda is misevaluated as a reference type, making the test reference + // (and instantiate) different kernels than the ones the dispatch launches. + // See: https://github.com/NVIDIA/cccl/issues/9643 + using input_it_t = decltype(d_in); + using output_it_t = decltype(d_out.begin()); + [[maybe_unused]] init_value_t init = 0; size_t expected_bytes_allocated{}; @@ -495,38 +481,15 @@ C2H_TEST("Device sum uses environment", "[reduce][device]", requirements) { REQUIRE(cudaSuccess == cub::DeviceReduce::Sum(nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items)); - using policy_t = cub::detail::reduce::policy_selector_from_types; + // Build the allowlist from the same kernel source type the dispatch uses, so the expected kernel + // instantiations cannot drift from the launched ones. + using policy_t = cub::detail::reduce::policy_selector_from_types; + using kernel_source_t = cub::detail::reduce:: + DeviceReduceKernelSource; return cuda::std::array{ - reinterpret_cast( - cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - decltype(d_in), - decltype(d_out.begin()), - offset_t, - op_t, - init_value_t, - accumulator_t, - transform_t>), - reinterpret_cast( - cub::detail::reduce::DeviceReduceKernel< - policy_t, - /* StableReductionOrder */ true, - decltype(d_in), - accumulator_t*, - offset_t, - op_t, - accumulator_t, - init_value_t, - transform_t>), - reinterpret_cast( - cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - accumulator_t*, - decltype(d_out.begin()), - int, // always used with int offset - op_t, - init_value_t, - accumulator_t>)}; + reinterpret_cast(kernel_source_t::SingleTileKernel()), + reinterpret_cast(kernel_source_t::ReductionKernel()), + reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; } else if constexpr (cub::detail::is_non_deterministic_v) { @@ -547,17 +510,18 @@ C2H_TEST("Device sum uses environment", "[reduce][device]", requirements) /* stream */ nullptr, transform_t{})); - return cuda::std::array{reinterpret_cast( - cub::detail::reduce::DeviceReduceKernel< - policy_t, - /* StableReductionOrder */ false, - decltype(d_in), - decltype(raw_ptr), - offset_t, - op_t, - accumulator_t, - init_value_t, - transform_t>)}; + // See the run_to_run branch above for why the kernel pointers must come from the kernel source + using kernel_source_t = cub::detail::reduce::DeviceReduceKernelSource< + policy_t, + input_it_t, + decltype(raw_ptr), + offset_t, + op_t, + init_value_t, + accumulator_t, + transform_t, + /* StableReductionOrder */ false>; + return cuda::std::array{reinterpret_cast(kernel_source_t::ReductionKernel())}; } else { @@ -566,36 +530,29 @@ C2H_TEST("Device sum uses environment", "[reduce][device]", requirements) using policy_t = cub::detail::reduce:: policy_selector_from_types; using deterministic_accum_t = deterministic_add_t::DeterministicAcc; - using output_it_t = decltype(d_out.begin()); - REQUIRE(cudaSuccess - == cub::detail::rfa:: - dispatch( - nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, init)); + REQUIRE( + cudaSuccess + == cub::detail::rfa::dispatch( + nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, init)); - auto k1 = cub::detail::reduce::DeterministicDeviceReduceSingleTileKernel< - policy_t, - decltype(d_in), - output_it_t, - reduction_op_t, - init_value_t, - deterministic_accum_t, - transform_t>; - auto k2 = cub::detail::reduce:: - DeterministicDeviceReduceKernel; - auto k3 = cub::detail::reduce::DeterministicDeviceReduceSingleTileKernel< + // See the run_to_run branch above for why the kernel pointers must come from the kernel source + using kernel_source_t = cub::detail::rfa::DeterministicDeviceReduceKernelSource< policy_t, - deterministic_accum_t*, + thrust::try_unwrap_contiguous_iterator_t, output_it_t, reduction_op_t, init_value_t, deterministic_accum_t, transform_t>; // TODO(bgruber): enable this when we have Catch2 3.13+ - // UNSCOPED_CAPTURE(c2h::type_name(), c2h::type_name(), - // c2h::type_name()); + // UNSCOPED_CAPTURE(c2h::type_name(), + // c2h::type_name(), + // c2h::type_name()); return cuda::std::array{ - reinterpret_cast(k1), reinterpret_cast(k2), reinterpret_cast(k3)}; + reinterpret_cast(kernel_source_t::SingleTileKernel()), + reinterpret_cast(kernel_source_t::ReductionKernel()), + reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; } }(); @@ -631,38 +588,22 @@ C2H_TEST("Device reduce not_guaranteed falls back when output type differs from == cub::DeviceReduce::Reduce( nullptr, expected_bytes_allocated, d_in.begin(), d_out.begin(), num_items, op_t{}, init)); - using policy_t = cub::detail::reduce::policy_selector_from_types; - auto kernels = cuda::std::array{ - reinterpret_cast( - cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - decltype(d_in.begin()), - decltype(d_out.begin()), - offset_t, - op_t, - init_value_t, - accumulator_t, - transform_t>), - reinterpret_cast( - cub::detail::reduce::DeviceReduceKernel< - policy_t, - /* StableReductionOrder */ true, - decltype(d_in.begin()), - accumulator_t*, - offset_t, - op_t, - accumulator_t, - init_value_t, - transform_t>), - reinterpret_cast( - cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - accumulator_t*, - decltype(d_out.begin()), - int, // always used with int offset - op_t, - init_value_t, - accumulator_t>)}; + // Build the allowlist from the same kernel source type the dispatch uses, so the expected kernel + // instantiations cannot drift from the launched ones. + using policy_t = cub::detail::reduce::policy_selector_from_types; + using kernel_source_t = cub::detail::reduce::DeviceReduceKernelSource< + policy_t, + decltype(d_in.begin()), + decltype(d_out.begin()), + offset_t, + op_t, + init_value_t, + accumulator_t, + transform_t>; + auto kernels = cuda::std::array{ + reinterpret_cast(kernel_source_t::SingleTileKernel()), + reinterpret_cast(kernel_source_t::ReductionKernel()), + reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; auto env = stdexec::env{cuda::execution::require(cuda::execution::determinism::not_guaranteed), allowed_kernels(kernels), @@ -692,38 +633,22 @@ C2H_TEST("Device sum not_guaranteed falls back when output type differs from acc REQUIRE( cudaSuccess == cub::DeviceReduce::Sum(nullptr, expected_bytes_allocated, d_in.begin(), d_out.begin(), num_items)); - using policy_t = cub::detail::reduce::policy_selector_from_types; - auto kernels = cuda::std::array{ - reinterpret_cast( - cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - decltype(d_in.begin()), - decltype(d_out.begin()), - offset_t, - op_t, - init_value_t, - accumulator_t, - transform_t>), - reinterpret_cast( - cub::detail::reduce::DeviceReduceKernel< - policy_t, - /* StableReductionOrder */ true, - decltype(d_in.begin()), - accumulator_t*, - offset_t, - op_t, - accumulator_t, - init_value_t, - transform_t>), - reinterpret_cast( - cub::detail::reduce::DeviceReduceSingleTileKernel< - policy_t, - accumulator_t*, - decltype(d_out.begin()), - int, // always used with int offset - op_t, - init_value_t, - accumulator_t>)}; + // Build the allowlist from the same kernel source type the dispatch uses, so the expected kernel + // instantiations cannot drift from the launched ones. + using policy_t = cub::detail::reduce::policy_selector_from_types; + using kernel_source_t = cub::detail::reduce::DeviceReduceKernelSource< + policy_t, + decltype(d_in.begin()), + decltype(d_out.begin()), + offset_t, + op_t, + init_value_t, + accumulator_t, + transform_t>; + auto kernels = cuda::std::array{ + reinterpret_cast(kernel_source_t::SingleTileKernel()), + reinterpret_cast(kernel_source_t::ReductionKernel()), + reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; auto env = stdexec::env{cuda::execution::require(cuda::execution::determinism::not_guaranteed), allowed_kernels(kernels), From eb895aff4ac3f9e3a5121ac4fac68cb1b9c3a155 Mon Sep 17 00:00:00 2001 From: edenfunf Date: Tue, 7 Jul 2026 17:37:16 +0800 Subject: [PATCH 2/2] Address review: dedup allowlist construction, add missing include * 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 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. --- .../dispatch_reduce_deterministic.cuh | 50 ++-- cub/test/catch2_test_device_reduce_env.cu | 268 ++++++------------ cub/test/catch2_test_env_launch_helper.h | 5 + 3 files changed, 129 insertions(+), 194 deletions(-) diff --git a/cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh b/cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh index 6eeab543919..b62c48f58e4 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh @@ -27,6 +27,8 @@ #include #include +#include + #include #include #include @@ -123,6 +125,25 @@ struct DeterministicDeviceReduceKernelSource DeterministicAccumT>) }; +//! Kernel source instantiation matching the default used by `rfa::dispatch` below. Tests use this alias to +//! reference the exact kernel instantiations the dispatch launches. +template , + typename PolicySelector = + policy_selector_from_types, __determinism_t::__gpu_to_gpu>> +using default_kernel_source_t = DeterministicDeviceReduceKernelSource< + PolicySelector, + THRUST_NS_QUALIFIER::try_unwrap_contiguous_iterator_t, + OutputIteratorT, + deterministic_sum_t, + InitValueT, + typename deterministic_sum_t::DeterministicAcc, + TransformOpT>; + template , - typename PolicySelector = - policy_selector_from_types, __determinism_t::__gpu_to_gpu>, - typename KernelSource = DeterministicDeviceReduceKernelSource< - PolicySelector, - THRUST_NS_QUALIFIER::try_unwrap_contiguous_iterator_t, - OutputIteratorT, - deterministic_sum_t, - InitValueT, - typename deterministic_sum_t::DeterministicAcc, - TransformOpT>, - typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> +template < + typename InputIteratorT, + typename OutputIteratorT, + typename OffsetT, + typename InitValueT, + typename TransformOpT = ::cuda::std::identity, + typename AccumT = accum_t, + typename PolicySelector = + policy_selector_from_types, __determinism_t::__gpu_to_gpu>, + typename KernelSource = + default_kernel_source_t, + typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t dispatch( void* d_temp_storage, size_t& temp_storage_bytes, diff --git a/cub/test/catch2_test_device_reduce_env.cu b/cub/test/catch2_test_device_reduce_env.cu index 5fc19538973..6cafdbc0a1a 100644 --- a/cub/test/catch2_test_device_reduce_env.cu +++ b/cub/test/catch2_test_device_reduce_env.cu @@ -333,6 +333,88 @@ using requirements = cuda::execution::determinism::run_to_run_t, cuda::execution::determinism::not_guaranteed_t>; +// Builds the kernel allowlist for the requirements-parametrized tests below and computes the expected temporary +// storage size. To check if a given algorithm implementation is used, the tests check if the associated kernels are +// invoked. The allowlists are built from the same kernel source types the dispatch uses, so the expected kernel +// instantiations cannot drift from the launched ones. This is deliberately a function template instead of a lambda +// in the tests: with nvcc 13.3 on MSVC, decltype() of a variable captured by reference inside a lambda is +// misevaluated as a reference type, which would make the allowlist reference (and instantiate) different kernels +// than the ones the dispatch launches. See: https://github.com/NVIDIA/cccl/issues/9643 +template +static auto expected_reduce_kernels( + InputItT d_in, + OutputItT d_out, + AccumulatorT* d_out_ptr, + NumItemsT num_items, + InitT init, + size_t& expected_bytes_allocated, + RunToRunSizeF run_to_run_size) +{ + if constexpr (std::is_same_v) + { + REQUIRE(cudaSuccess == run_to_run_size()); + + using policy_t = cub::detail::reduce::policy_selector_from_types; + using kernel_source_t = cub::detail::reduce:: + DeviceReduceKernelSource; + return cuda::std::array{ + reinterpret_cast(kernel_source_t::SingleTileKernel()), + reinterpret_cast(kernel_source_t::ReductionKernel()), + reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; + } + else if constexpr (cub::detail::is_non_deterministic_v) + { + REQUIRE( + cudaSuccess + == cub::detail::reduce::dispatch( + nullptr, + expected_bytes_allocated, + d_in, + d_out_ptr, + num_items, + OpT{}, + init, + /* stream */ nullptr, + TransformT{})); + + using policy_t = + cub::detail::reduce::policy_selector_from_types; + using kernel_source_t = cub::detail::reduce::DeviceReduceKernelSource< + policy_t, + InputItT, + AccumulatorT*, + OffsetT, + OpT, + InitT, + AccumulatorT, + TransformT, + /* StableReductionOrder */ false>; + return cuda::std::array{reinterpret_cast(kernel_source_t::ReductionKernel())}; + } + else + { + REQUIRE(cudaSuccess + == cub::detail::rfa::dispatch( + nullptr, expected_bytes_allocated, d_in, d_out, num_items, init)); + + using kernel_source_t = + cub::detail::rfa::default_kernel_source_t; + return cuda::std::array{ + reinterpret_cast(kernel_source_t::SingleTileKernel()), + reinterpret_cast(kernel_source_t::ReductionKernel()), + reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; + } +} + C2H_TEST("Device reduce uses environment", "[reduce][device]", requirements) { using determinism_t = c2h::get<0, TestType>; @@ -347,98 +429,13 @@ C2H_TEST("Device reduce uses environment", "[reduce][device]", requirements) auto d_in = cuda::constant_iterator(1.0f); auto d_out = thrust::device_vector(1); - // Compute the iterator types outside of the allowlist lambda below: with nvcc 13.3 on MSVC, decltype() of a - // variable captured by reference inside a lambda is misevaluated as a reference type, making the test reference - // (and instantiate) different kernels than the ones the dispatch launches. - // See: https://github.com/NVIDIA/cccl/issues/9643 - using input_it_t = decltype(d_in); - using output_it_t = decltype(d_out.begin()); - init_value_t init = 0; size_t expected_bytes_allocated{}; - // To check if a given algorithm implementation is used, we check if associated kernels are invoked. - auto kernels = [&]() { - if constexpr (std::is_same_v) - { - REQUIRE( - cudaSuccess - == cub::DeviceReduce::Reduce(nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, op_t{}, init)); - - // Build the allowlist from the same kernel source type the dispatch uses, so the expected kernel - // instantiations cannot drift from the launched ones. - using policy_t = cub::detail::reduce::policy_selector_from_types; - using kernel_source_t = cub::detail::reduce:: - DeviceReduceKernelSource; - return cuda::std::array{ - reinterpret_cast(kernel_source_t::SingleTileKernel()), - reinterpret_cast(kernel_source_t::ReductionKernel()), - reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; - } - else if constexpr (cub::detail::is_non_deterministic_v) - { - using policy_t = - cub::detail::reduce::policy_selector_from_types; - auto* raw_ptr = thrust::raw_pointer_cast(d_out.data()); - - REQUIRE( - cudaSuccess - == cub::detail::reduce::dispatch( - nullptr, - expected_bytes_allocated, - d_in, - raw_ptr, - num_items, - op_t{}, - init, - /* stream */ nullptr, - transform_t{})); - - // See the run_to_run branch above for why the kernel pointers must come from the kernel source - using kernel_source_t = cub::detail::reduce::DeviceReduceKernelSource< - policy_t, - input_it_t, - decltype(raw_ptr), - offset_t, - op_t, - init_value_t, - accumulator_t, - transform_t, - /* StableReductionOrder */ false>; - return cuda::std::array{reinterpret_cast(kernel_source_t::ReductionKernel())}; - } - else - { - using deterministic_add_t = cub::detail::rfa::deterministic_sum_t; - using reduction_op_t = deterministic_add_t; - using policy_t = cub::detail::reduce:: - policy_selector_from_types; - using deterministic_accum_t = deterministic_add_t::DeterministicAcc; - - REQUIRE( - cudaSuccess - == cub::detail::rfa::dispatch( - nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, init)); - - // See the run_to_run branch above for why the kernel pointers must come from the kernel source - using kernel_source_t = cub::detail::rfa::DeterministicDeviceReduceKernelSource< - policy_t, - thrust::try_unwrap_contiguous_iterator_t, - output_it_t, - reduction_op_t, - init_value_t, - deterministic_accum_t, - transform_t>; - // TODO(bgruber): enable this when we have Catch2 3.13+ - // UNSCOPED_CAPTURE(c2h::type_name(), - // c2h::type_name(), - // c2h::type_name()); - return cuda::std::array{ - reinterpret_cast(kernel_source_t::SingleTileKernel()), - reinterpret_cast(kernel_source_t::ReductionKernel()), - reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; - } - }(); + auto kernels = expected_reduce_kernels( + d_in, d_out.begin(), thrust::raw_pointer_cast(d_out.data()), num_items, init, expected_bytes_allocated, [&]() { + return cub::DeviceReduce::Reduce(nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, op_t{}, init); + }); // Equivalent to `cuexec::require(cuexec::determinism::run_to_run)` and // `cuexec::require(cuexec::determinism::not_guaranteed)` @@ -465,96 +462,13 @@ C2H_TEST("Device sum uses environment", "[reduce][device]", requirements) auto d_in = cuda::constant_iterator(1.0f); auto d_out = thrust::device_vector(1); - // Compute the iterator types outside of the allowlist lambda below: with nvcc 13.3 on MSVC, decltype() of a - // variable captured by reference inside a lambda is misevaluated as a reference type, making the test reference - // (and instantiate) different kernels than the ones the dispatch launches. - // See: https://github.com/NVIDIA/cccl/issues/9643 - using input_it_t = decltype(d_in); - using output_it_t = decltype(d_out.begin()); - - [[maybe_unused]] init_value_t init = 0; + init_value_t init = 0; size_t expected_bytes_allocated{}; - // To check if a given algorithm implementation is used, we check if associated kernels are invoked. - auto kernels = [&]() { - if constexpr (std::is_same_v) - { - REQUIRE(cudaSuccess == cub::DeviceReduce::Sum(nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items)); - - // Build the allowlist from the same kernel source type the dispatch uses, so the expected kernel - // instantiations cannot drift from the launched ones. - using policy_t = cub::detail::reduce::policy_selector_from_types; - using kernel_source_t = cub::detail::reduce:: - DeviceReduceKernelSource; - return cuda::std::array{ - reinterpret_cast(kernel_source_t::SingleTileKernel()), - reinterpret_cast(kernel_source_t::ReductionKernel()), - reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; - } - else if constexpr (cub::detail::is_non_deterministic_v) - { - using policy_t = - cub::detail::reduce::policy_selector_from_types; - auto* raw_ptr = thrust::raw_pointer_cast(d_out.data()); - - REQUIRE( - cudaSuccess - == cub::detail::reduce::dispatch( - nullptr, - expected_bytes_allocated, - d_in, - raw_ptr, - num_items, - op_t{}, - init, - /* stream */ nullptr, - transform_t{})); - - // See the run_to_run branch above for why the kernel pointers must come from the kernel source - using kernel_source_t = cub::detail::reduce::DeviceReduceKernelSource< - policy_t, - input_it_t, - decltype(raw_ptr), - offset_t, - op_t, - init_value_t, - accumulator_t, - transform_t, - /* StableReductionOrder */ false>; - return cuda::std::array{reinterpret_cast(kernel_source_t::ReductionKernel())}; - } - else - { - using deterministic_add_t = cub::detail::rfa::deterministic_sum_t; - using reduction_op_t = deterministic_add_t; - using policy_t = cub::detail::reduce:: - policy_selector_from_types; - using deterministic_accum_t = deterministic_add_t::DeterministicAcc; - - REQUIRE( - cudaSuccess - == cub::detail::rfa::dispatch( - nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items, init)); - - // See the run_to_run branch above for why the kernel pointers must come from the kernel source - using kernel_source_t = cub::detail::rfa::DeterministicDeviceReduceKernelSource< - policy_t, - thrust::try_unwrap_contiguous_iterator_t, - output_it_t, - reduction_op_t, - init_value_t, - deterministic_accum_t, - transform_t>; - // TODO(bgruber): enable this when we have Catch2 3.13+ - // UNSCOPED_CAPTURE(c2h::type_name(), - // c2h::type_name(), - // c2h::type_name()); - return cuda::std::array{ - reinterpret_cast(kernel_source_t::SingleTileKernel()), - reinterpret_cast(kernel_source_t::ReductionKernel()), - reinterpret_cast(kernel_source_t::SingleTileSecondKernel())}; - } - }(); + auto kernels = expected_reduce_kernels( + d_in, d_out.begin(), thrust::raw_pointer_cast(d_out.data()), num_items, init, expected_bytes_allocated, [&]() { + return cub::DeviceReduce::Sum(nullptr, expected_bytes_allocated, d_in, d_out.begin(), num_items); + }); // Equivalent to `cuexec::require(cuexec::determinism::run_to_run)` and // `cuexec::require(cuexec::determinism::not_guaranteed)` diff --git a/cub/test/catch2_test_env_launch_helper.h b/cub/test/catch2_test_env_launch_helper.h index 950bd5654eb..ca078719273 100644 --- a/cub/test/catch2_test_env_launch_helper.h +++ b/cub/test/catch2_test_env_launch_helper.h @@ -55,6 +55,11 @@ expected_allocation_size(size_t expected) struct get_allowed_kernels_t {}; +//! The listed kernels must be the exact instantiations the dispatch launches, taken from the dispatch's kernel +//! source. Beware of computing the expected instantiations with decltype() of a variable captured by reference +//! inside a lambda: nvcc 13.3 with MSVC misevaluates such a decltype() as a reference type, silently yielding a +//! different kernel instantiation. Compute the types at function scope instead. +//! See: https://github.com/NVIDIA/cccl/issues/9643 __host__ __device__ static cuda::std::execution::prop> allowed_kernels(cuda::std::span allowed_kernels) {