From 3a65f23e87d5581702704eb62536512c8598e1c1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Wed, 8 Jul 2026 06:43:29 -0700 Subject: [PATCH 1/8] Set fixed_capacity_map CG size default --- cudax/include/cuda/experimental/__cuco/fixed_capacity_map.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cudax/include/cuda/experimental/__cuco/fixed_capacity_map.cuh b/cudax/include/cuda/experimental/__cuco/fixed_capacity_map.cuh index 3b1dd8e91a4..c7f407dd935 100644 --- a/cudax/include/cuda/experimental/__cuco/fixed_capacity_map.cuh +++ b/cudax/include/cuda/experimental/__cuco/fixed_capacity_map.cuh @@ -72,7 +72,7 @@ template , - class _ProbingScheme = linear_probing<1, hash<_Key>>, + class _ProbingScheme = linear_probing<4, hash<_Key>>, int _BucketSize = 1, class _MemoryResource = ::cuda::device_memory_pool_ref> class fixed_capacity_map From 655fe24c23e86b1dc10534f3072ff1baf2984838 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Wed, 8 Jul 2026 06:43:47 -0700 Subject: [PATCH 2/8] Port fixed_capacity_map benchmarks --- .../benchmarks/bench/cuco/common/defaults.cuh | 44 +++ .../bench/cuco/common/key_generator.cuh | 283 ++++++++++++++++++ .../bench/cuco/fixed_capacity_map/contains.cu | 117 ++++++++ .../bench/cuco/fixed_capacity_map/insert.cu | 112 +++++++ 4 files changed, 556 insertions(+) create mode 100644 cudax/benchmarks/bench/cuco/common/defaults.cuh create mode 100644 cudax/benchmarks/bench/cuco/common/key_generator.cuh create mode 100644 cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu create mode 100644 cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu diff --git a/cudax/benchmarks/bench/cuco/common/defaults.cuh b/cudax/benchmarks/bench/cuco/common/defaults.cuh new file mode 100644 index 00000000000..8b946a6e719 --- /dev/null +++ b/cudax/benchmarks/bench/cuco/common/defaults.cuh @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +#include + +#include +#include + +namespace cuda::experimental::cuco::benchmark::defaults +{ +//! Key types covered by the default CUCO benchmark type axes. +using key_type_range = ::nvbench::type_list<::nvbench::int32_t, ::nvbench::int64_t>; +//! Value types covered by the default CUCO benchmark type axes. +using value_type_range = ::nvbench::type_list<::nvbench::int32_t, ::nvbench::int64_t>; + +//! Default number of inputs used when sweeping another benchmark axis. +inline constexpr auto n = ::nvbench::int64_t{100'000'000}; +//! Default fixed-capacity map target occupancy. +inline constexpr auto occupancy = 0.5; +//! Default lookup matching rate for contains-style benchmarks. +inline constexpr auto matching_rate = 1.0; +//! Default deterministic seed used by benchmark data generators. +inline constexpr auto seed = ::cuda::std::uint32_t{42}; + +//! Input-size sweep that remains cacheable for direct comparisons with CUCO benchmarks. +inline const auto n_range_cache = ::std::vector<::nvbench::int64_t>{8'000, 80'000, 800'000, 8'000'000, 80'000'000}; +//! Occupancy sweep used by fixed-capacity container benchmarks. +inline const auto occupancy_range = ::nvbench::range(0.1, 0.9, 0.1); +//! Average multiplicity sweep for duplicate-key distributions. +inline const auto multiplicity_range = ::std::vector<::nvbench::int64_t>{1, 2, 4, 8, 16}; +//! Matching-rate sweep used by contains-style benchmarks. +inline const auto matching_rate_range = ::nvbench::range(0.1, 1.0, 0.1); +} // namespace cuda::experimental::cuco::benchmark::defaults diff --git a/cudax/benchmarks/bench/cuco/common/key_generator.cuh b/cudax/benchmarks/bench/cuco/common/key_generator.cuh new file mode 100644 index 00000000000..87e96492ccf --- /dev/null +++ b/cudax/benchmarks/bench/cuco/common/key_generator.cuh @@ -0,0 +1,283 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include + +#include "defaults.cuh" +#include + +namespace cuda::experimental::cuco::benchmark +{ +namespace distribution +{ +//! Distribution tag for unique keys generated by shuffling the sequence `[0, N)`. +struct unique +{}; + +//! Distribution tag for uniformly sampled keys with controlled average multiplicity. +struct uniform +{ + //! Constructs a uniform distribution tag with the requested average key multiplicity. + //! + //! @param __multiplicity Average number of generated keys that map to each unique key value. + //! @throws std::invalid_argument if `__multiplicity` is not positive. + explicit uniform(::nvbench::int64_t __multiplicity) + : __value{__multiplicity} + { + if (__multiplicity <= 0) + { + throw ::std::invalid_argument{"Multiplicity must be greater than 0"}; + } + } + + //! Average number of generated keys that map to each unique key value. + ::nvbench::int64_t __value; +}; +} // namespace distribution + +namespace detail +{ +template +struct __generate_uniform_fn +{ + _CCCL_HOST_DEVICE_API constexpr __generate_uniform_fn( + ::cuda::std::size_t __num, _Dist __dist, ::cuda::std::size_t __seed) + : __num{__num} + , __dist{__dist} + , __seed{__seed} + {} + + _CCCL_HOST_DEVICE_API constexpr _Tp operator()(::cuda::std::size_t __idx) const noexcept + { + _Rng __rng; + __rng.seed(__seed + __idx * 1664525ull + 1013904223ull); + + const auto __num_unique_keys_unclamped = static_cast<::cuda::std::size_t>( + ::cuda::std::ceil(static_cast(__num) / static_cast(__dist.__value))); + const auto __num_unique_keys = + __num_unique_keys_unclamped < ::cuda::std::size_t{1} ? ::cuda::std::size_t{1} : __num_unique_keys_unclamped; + ::thrust::uniform_int_distribution<_Tp> __key_dist{_Tp{0}, static_cast<_Tp>(__num_unique_keys - 1)}; + return __key_dist(__rng); + } + + ::cuda::std::size_t __num; + _Dist __dist; + ::cuda::std::size_t __seed; +}; + +template +struct __dropout_fn +{ + _CCCL_HOST_DEVICE_API constexpr explicit __dropout_fn(::cuda::std::size_t __num) + : __num{__num} + {} + + _CCCL_HOST_DEVICE_API _Tp operator()(::cuda::std::size_t __seed) const noexcept + { + _Rng __rng; + ::thrust::uniform_int_distribution<_Tp> __dist{static_cast<_Tp>(__num), ::cuda::std::numeric_limits<_Tp>::max()}; + __rng.seed(__seed); + return __dist(__rng); + } + + ::cuda::std::size_t __num; +}; + +template +struct __dropout_pred +{ + _CCCL_HOST_DEVICE_API constexpr explicit __dropout_pred(double __keep_prob) + : __keep_prob{__keep_prob} + {} + + _CCCL_HOST_DEVICE_API bool operator()(::cuda::std::size_t __seed) const noexcept + { + _Rng __rng; + ::thrust::uniform_real_distribution __dist{0.0, 1.0}; + __rng.seed(__seed); + return __dist(__rng) > __keep_prob; + } + + double __keep_prob; +}; +} // namespace detail + +//! Random key generator used by CUCO benchmarks. +//! +//! The generator defaults to `defaults::seed` to keep benchmark data reproducible across runs. +//! +//! @tparam _Rng Pseudo-random number generator type compatible with Thrust random distributions. +template +class key_generator +{ +public: + //! Constructs a key generator with the given seed. + //! + //! @param __seed Seed used to initialize the generator state. + explicit key_generator(::cuda::std::uint32_t __seed = defaults::seed) + : __rng{__seed} + {} + + //! Generates keys according to the given distribution using the default device execution policy. + //! + //! @tparam _Dist Distribution tag type. + //! @tparam _OutputIt Output iterator type whose value type is the generated key type. + //! @param __dist Distribution tag controlling how keys are generated. + //! @param __out_begin Beginning of the output key range. + //! @param __out_end End of the output key range. + //! @throws std::invalid_argument if `_Dist` is not a supported distribution tag. + template + void generate(_Dist __dist, _OutputIt __out_begin, _OutputIt __out_end) + { + generate(__dist, __out_begin, __out_end, ::thrust::device); + } + + //! Generates keys according to the given distribution using the provided execution policy. + //! + //! @tparam _Dist Distribution tag type. + //! @tparam _OutputIt Output iterator type whose value type is the generated key type. + //! @tparam _ExecPolicy Thrust execution policy type. + //! @param __dist Distribution tag controlling how keys are generated. + //! @param __out_begin Beginning of the output key range. + //! @param __out_end End of the output key range. + //! @param __exec_policy Execution policy used for the underlying Thrust algorithms. + //! @throws std::invalid_argument if `_Dist` is not a supported distribution tag. + template + void generate(_Dist __dist, _OutputIt __out_begin, _OutputIt __out_end, _ExecPolicy __exec_policy) + { + using __value_type = typename ::cuda::std::iterator_traits<_OutputIt>::value_type; + + if constexpr (::cuda::std::is_same_v<_Dist, distribution::unique>) + { + ::thrust::sequence(__exec_policy, __out_begin, __out_end, __value_type{0}); + ::thrust::shuffle(__exec_policy, __out_begin, __out_end, __rng); + } + else if constexpr (::cuda::std::is_same_v<_Dist, distribution::uniform>) + { + const auto __num_keys = static_cast<::cuda::std::size_t>(::cuda::std::distance(__out_begin, __out_end)); + const auto __seed = static_cast<::cuda::std::size_t>(__rng()); + + ::thrust::transform( + __exec_policy, + ::cuda::counting_iterator<::cuda::std::size_t>{0}, + ::cuda::counting_iterator<::cuda::std::size_t>{__num_keys}, + __out_begin, + detail::__generate_uniform_fn<__value_type, _Dist, _Rng>{__num_keys, __dist, __seed}); + } + else + { + throw ::std::invalid_argument{"Unexpected distribution type"}; + } + } + + //! Drops keys with probability `1 - keep_prob` using the default device execution policy. + //! + //! Replaced keys are sampled from `[N, max_key]`, where `N` is the number of keys in the range. + //! + //! @tparam _InOutIt Mutable iterator type whose value type is the key type. + //! @param __begin Beginning of the key range to update in place. + //! @param __end End of the key range to update in place. + //! @param __keep_prob Probability of keeping each original key. Must be in `[0, 1]`. + //! @throws std::invalid_argument if `__keep_prob` is outside `[0, 1]`. + template + void dropout(_InOutIt __begin, _InOutIt __end, double __keep_prob) + { + dropout(__begin, __end, __keep_prob, ::thrust::device); + } + + //! Drops keys with probability `1 - keep_prob` using the provided execution policy. + //! + //! Replaced keys are sampled from `[N, max_key]`, where `N` is the number of keys in the range. + //! + //! @tparam _InOutIt Mutable iterator type whose value type is the key type. + //! @tparam _ExecPolicy Thrust execution policy type. + //! @param __begin Beginning of the key range to update in place. + //! @param __end End of the key range to update in place. + //! @param __keep_prob Probability of keeping each original key. Must be in `[0, 1]`. + //! @param __exec_policy Execution policy used for the underlying Thrust algorithms. + //! @throws std::invalid_argument if `__keep_prob` is outside `[0, 1]`. + template + void dropout(_InOutIt __begin, _InOutIt __end, double __keep_prob, _ExecPolicy __exec_policy) + { + using __value_type = typename ::cuda::std::iterator_traits<_InOutIt>::value_type; + + if (__keep_prob < 0.0 || __keep_prob > 1.0) + { + throw ::std::invalid_argument{"Probability needs to be between 0 and 1"}; + } + + if (__keep_prob < 1.0) + { + const auto __num_keys = static_cast<::cuda::std::size_t>(::cuda::std::distance(__begin, __end)); + ::cuda::counting_iterator<::cuda::std::size_t> __seeds{static_cast<::cuda::std::size_t>(__rng())}; + + ::thrust::transform_if( + __exec_policy, + __seeds, + __seeds + __num_keys, + __begin, + detail::__dropout_fn<__value_type, _Rng>{__num_keys}, + detail::__dropout_pred<_Rng>{__keep_prob}); + } + + ::thrust::shuffle(__exec_policy, __begin, __end, __rng); + } + +private: + _Rng __rng; +}; + +//! Constructs the requested distribution tag from NVBench axis values. +//! +//! `distribution::uniform` reads the `Multiplicity` axis from `state`. +//! +//! @tparam _Dist Distribution tag type to construct. +//! @param __state NVBench state containing distribution-specific axis values. +//! @return Distribution tag initialized from the benchmark state. +//! @throws std::invalid_argument if `_Dist` is not a supported distribution tag. +template +_Dist dist_from_state(::nvbench::state const& __state) +{ + if constexpr (::cuda::std::is_same_v<_Dist, distribution::unique>) + { + return _Dist{}; + } + else if constexpr (::cuda::std::is_same_v<_Dist, distribution::uniform>) + { + return _Dist{__state.get_int64("Multiplicity")}; + } + else + { + throw ::std::invalid_argument{"Unexpected distribution type"}; + } +} +} // namespace cuda::experimental::cuco::benchmark + +NVBENCH_DECLARE_TYPE_STRINGS( + ::cuda::experimental::cuco::benchmark::distribution::unique, "UNIQUE", "distribution::unique"); +NVBENCH_DECLARE_TYPE_STRINGS( + ::cuda::experimental::cuco::benchmark::distribution::uniform, "UNIFORM", "distribution::uniform"); diff --git a/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu b/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu new file mode 100644 index 00000000000..a700ffdc741 --- /dev/null +++ b/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu @@ -0,0 +1,117 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#include +#include + +#include +#include +#include +#include +#include +#include + +#include +#include + +#include "../common/defaults.cuh" +#include "../common/key_generator.cuh" +#include + +namespace cudax = ::cuda::experimental; +namespace bench = ::cuda::experimental::cuco::benchmark; + +/** + * @brief A benchmark evaluating `cudax::cuco::fixed_capacity_map::contains_async` performance. + */ +template +void fixed_capacity_map_contains(::nvbench::state& state, ::nvbench::type_list) +{ + if constexpr (sizeof(Key) != sizeof(Value)) + { + state.skip("Key should be the same type as Value."); + } + else + { + using pair_type = ::cuda::std::pair; + using map_type = cudax::cuco::fixed_capacity_map; + + const auto num_keys = state.get_int64("NumInputs"); + const auto occupancy = state.get_float64("Occupancy"); + const auto matching_rate = state.get_float64("MatchingRate"); + + const auto size = static_cast<::cuda::std::size_t>(static_cast(num_keys) / occupancy); + + const auto device = ::cuda::device_ref{0}; + ::cuda::stream stream{device}; + const ::cuda::device_memory_pool_ref memory_resource = ::cuda::device_default_memory_pool(device); + const auto exec_policy = ::thrust::cuda::par_nosync.on(stream.get()); + + auto keys = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); + + bench::key_generator gen{}; + gen.generate(bench::dist_from_state(state), keys.begin(), keys.end(), exec_policy); + + auto pairs = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); + ::thrust::transform(exec_policy, keys.begin(), keys.end(), pairs.begin(), [] __device__(Key const& key) { + return pair_type{key, Value{}}; + }); + + map_type map{ + size, + cudax::cuco::empty_key{Key{-1}}, + cudax::cuco::empty_value{Value{-1}}, + ::cuda::std::equal_to{}, + typename map_type::probing_scheme_type{}, + memory_resource, + stream}; + map.insert(pairs.begin(), pairs.end(), stream); + + gen.dropout(keys.begin(), keys.end(), matching_rate, exec_policy); + + auto result = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); + stream.sync(); + + state.add_element_count(num_keys); + state.exec([&](::nvbench::launch& launch) { + map.contains_async(keys.begin(), keys.end(), result.begin(), {launch.get_stream()}); + }); + } +} + +NVBENCH_BENCH_TYPES(fixed_capacity_map_contains, + NVBENCH_TYPE_AXES(bench::defaults::key_type_range, + bench::defaults::value_type_range, + ::nvbench::type_list)) + .set_name("fixed_capacity_map_contains_unique_capacity") + .set_type_axes_names({"Key", "Value", "Distribution"}) + .add_int64_axis("NumInputs", bench::defaults::n_range_cache) + .add_float64_axis("Occupancy", {bench::defaults::occupancy}) + .add_float64_axis("MatchingRate", {bench::defaults::matching_rate}); + +NVBENCH_BENCH_TYPES(fixed_capacity_map_contains, + NVBENCH_TYPE_AXES(bench::defaults::key_type_range, + bench::defaults::value_type_range, + ::nvbench::type_list)) + .set_name("fixed_capacity_map_contains_unique_occupancy") + .set_type_axes_names({"Key", "Value", "Distribution"}) + .add_int64_axis("NumInputs", {bench::defaults::n}) + .add_float64_axis("Occupancy", bench::defaults::occupancy_range) + .add_float64_axis("MatchingRate", {bench::defaults::matching_rate}); + +NVBENCH_BENCH_TYPES(fixed_capacity_map_contains, + NVBENCH_TYPE_AXES(bench::defaults::key_type_range, + bench::defaults::value_type_range, + ::nvbench::type_list)) + .set_name("fixed_capacity_map_contains_unique_matching_rate") + .set_type_axes_names({"Key", "Value", "Distribution"}) + .add_int64_axis("NumInputs", {bench::defaults::n}) + .add_float64_axis("Occupancy", {bench::defaults::occupancy}) + .add_float64_axis("MatchingRate", bench::defaults::matching_rate_range); diff --git a/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu b/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu new file mode 100644 index 00000000000..773e0bfb026 --- /dev/null +++ b/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu @@ -0,0 +1,112 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#include +#include + +#include +#include +#include +#include +#include +#include + +#include +#include + +#include "../common/defaults.cuh" +#include "../common/key_generator.cuh" +#include + +namespace cudax = ::cuda::experimental; +namespace bench = ::cuda::experimental::cuco::benchmark; + +/** + * @brief A benchmark evaluating `cudax::cuco::fixed_capacity_map::insert_async` performance. + */ +template +void fixed_capacity_map_insert(::nvbench::state& state, ::nvbench::type_list) +{ + if constexpr (sizeof(Key) != sizeof(Value)) + { + state.skip("Key should be the same type as Value."); + } + else + { + using pair_type = ::cuda::std::pair; + using map_type = cudax::cuco::fixed_capacity_map; + + const auto num_keys = state.get_int64("NumInputs"); + const auto occupancy = state.get_float64("Occupancy"); + + const auto size = static_cast<::cuda::std::size_t>(static_cast(num_keys) / occupancy); + + const auto device = ::cuda::device_ref{0}; + ::cuda::stream stream{device}; + const ::cuda::device_memory_pool_ref memory_resource = ::cuda::device_default_memory_pool(device); + const auto exec_policy = ::thrust::cuda::par_nosync.on(stream.get()); + + auto keys = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); + + bench::key_generator gen{}; + gen.generate(bench::dist_from_state(state), keys.begin(), keys.end(), exec_policy); + + auto pairs = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); + ::thrust::transform(exec_policy, keys.begin(), keys.end(), pairs.begin(), [] __device__(Key const& key) { + return pair_type{key, Value{}}; + }); + + map_type map{ + size, + cudax::cuco::empty_key{Key{-1}}, + cudax::cuco::empty_value{Value{-1}}, + ::cuda::std::equal_to{}, + typename map_type::probing_scheme_type{}, + memory_resource, + stream}; + stream.sync(); + + state.add_element_count(num_keys); + state.exec(::nvbench::exec_tag::timer, [&](::nvbench::launch& launch, auto& timer) { + timer.start(); + map.insert_async(pairs.begin(), pairs.end(), {launch.get_stream()}); + timer.stop(); + map.clear_async({launch.get_stream()}); + }); + } +} + +NVBENCH_BENCH_TYPES(fixed_capacity_map_insert, + NVBENCH_TYPE_AXES(bench::defaults::key_type_range, + bench::defaults::value_type_range, + ::nvbench::type_list)) + .set_name("fixed_capacity_map_insert_unique_capacity") + .set_type_axes_names({"Key", "Value", "Distribution"}) + .add_int64_axis("NumInputs", bench::defaults::n_range_cache) + .add_float64_axis("Occupancy", {bench::defaults::occupancy}); + +NVBENCH_BENCH_TYPES(fixed_capacity_map_insert, + NVBENCH_TYPE_AXES(bench::defaults::key_type_range, + bench::defaults::value_type_range, + ::nvbench::type_list)) + .set_name("fixed_capacity_map_insert_unique_occupancy") + .set_type_axes_names({"Key", "Value", "Distribution"}) + .add_int64_axis("NumInputs", {bench::defaults::n}) + .add_float64_axis("Occupancy", bench::defaults::occupancy_range); + +NVBENCH_BENCH_TYPES(fixed_capacity_map_insert, + NVBENCH_TYPE_AXES(bench::defaults::key_type_range, + bench::defaults::value_type_range, + ::nvbench::type_list)) + .set_name("fixed_capacity_map_insert_uniform_multiplicity") + .set_type_axes_names({"Key", "Value", "Distribution"}) + .add_int64_axis("NumInputs", {bench::defaults::n}) + .add_float64_axis("Occupancy", {bench::defaults::occupancy}) + .add_int64_axis("Multiplicity", bench::defaults::multiplicity_range); From 2ea4d34b15c872dfc73411efff59063d1e11c8da Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Wed, 8 Jul 2026 06:48:56 -0700 Subject: [PATCH 3/8] Tidy benchmark namespace aliases --- cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu | 2 +- cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu b/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu index a700ffdc741..8858ffc7da0 100644 --- a/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu +++ b/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu @@ -26,7 +26,7 @@ #include namespace cudax = ::cuda::experimental; -namespace bench = ::cuda::experimental::cuco::benchmark; +namespace bench = cudax::cuco::benchmark; /** * @brief A benchmark evaluating `cudax::cuco::fixed_capacity_map::contains_async` performance. diff --git a/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu b/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu index 773e0bfb026..2f3702fda41 100644 --- a/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu +++ b/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu @@ -26,7 +26,7 @@ #include namespace cudax = ::cuda::experimental; -namespace bench = ::cuda::experimental::cuco::benchmark; +namespace bench = cudax::cuco::benchmark; /** * @brief A benchmark evaluating `cudax::cuco::fixed_capacity_map::insert_async` performance. From dee3726bb96cd40fe75d13c53a68023cf0a16292 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Wed, 8 Jul 2026 09:35:17 -0700 Subject: [PATCH 4/8] Address review comments --- .../bench/cuco/common/key_generator.cuh | 244 +++++++++--------- .../fixed_capacity_map/test_shared_memory.cu | 27 +- 2 files changed, 137 insertions(+), 134 deletions(-) diff --git a/cudax/benchmarks/bench/cuco/common/key_generator.cuh b/cudax/benchmarks/bench/cuco/common/key_generator.cuh index 87e96492ccf..05eaab6c6ad 100644 --- a/cudax/benchmarks/bench/cuco/common/key_generator.cuh +++ b/cudax/benchmarks/bench/cuco/common/key_generator.cuh @@ -43,86 +43,86 @@ struct uniform { //! Constructs a uniform distribution tag with the requested average key multiplicity. //! - //! @param __multiplicity Average number of generated keys that map to each unique key value. - //! @throws std::invalid_argument if `__multiplicity` is not positive. - explicit uniform(::nvbench::int64_t __multiplicity) - : __value{__multiplicity} + //! @param multiplicity Average number of generated keys that map to each unique key value. + //! @throws std::invalid_argument if `multiplicity` is not positive. + explicit uniform(::nvbench::int64_t multiplicity) + : multiplicity{multiplicity} { - if (__multiplicity <= 0) + if (multiplicity <= 0) { throw ::std::invalid_argument{"Multiplicity must be greater than 0"}; } } //! Average number of generated keys that map to each unique key value. - ::nvbench::int64_t __value; + ::nvbench::int64_t multiplicity; }; } // namespace distribution namespace detail { -template -struct __generate_uniform_fn +template +struct generate_uniform_fn { - _CCCL_HOST_DEVICE_API constexpr __generate_uniform_fn( - ::cuda::std::size_t __num, _Dist __dist, ::cuda::std::size_t __seed) - : __num{__num} - , __dist{__dist} - , __seed{__seed} + _CCCL_HOST_DEVICE_API constexpr generate_uniform_fn( + ::cuda::std::size_t num, Distribution dist, ::cuda::std::size_t seed) + : num{num} + , dist{dist} + , seed{seed} {} - _CCCL_HOST_DEVICE_API constexpr _Tp operator()(::cuda::std::size_t __idx) const noexcept + _CCCL_HOST_DEVICE_API constexpr Key operator()(::cuda::std::size_t idx) const noexcept { - _Rng __rng; - __rng.seed(__seed + __idx * 1664525ull + 1013904223ull); - - const auto __num_unique_keys_unclamped = static_cast<::cuda::std::size_t>( - ::cuda::std::ceil(static_cast(__num) / static_cast(__dist.__value))); - const auto __num_unique_keys = - __num_unique_keys_unclamped < ::cuda::std::size_t{1} ? ::cuda::std::size_t{1} : __num_unique_keys_unclamped; - ::thrust::uniform_int_distribution<_Tp> __key_dist{_Tp{0}, static_cast<_Tp>(__num_unique_keys - 1)}; - return __key_dist(__rng); + Rng rng; + rng.seed(seed + idx * 1664525ull + 1013904223ull); + + const auto num_unique_keys_unclamped = static_cast<::cuda::std::size_t>( + ::cuda::std::ceil(static_cast(num) / static_cast(dist.multiplicity))); + const auto num_unique_keys = + num_unique_keys_unclamped < ::cuda::std::size_t{1} ? ::cuda::std::size_t{1} : num_unique_keys_unclamped; + ::thrust::uniform_int_distribution key_dist{Key{0}, static_cast(num_unique_keys - 1)}; + return key_dist(rng); } - ::cuda::std::size_t __num; - _Dist __dist; - ::cuda::std::size_t __seed; + ::cuda::std::size_t num; + Distribution dist; + ::cuda::std::size_t seed; }; -template -struct __dropout_fn +template +struct dropout_fn { - _CCCL_HOST_DEVICE_API constexpr explicit __dropout_fn(::cuda::std::size_t __num) - : __num{__num} + _CCCL_HOST_DEVICE_API constexpr explicit dropout_fn(::cuda::std::size_t num) + : num{num} {} - _CCCL_HOST_DEVICE_API _Tp operator()(::cuda::std::size_t __seed) const noexcept + _CCCL_HOST_DEVICE_API Key operator()(::cuda::std::size_t seed) const noexcept { - _Rng __rng; - ::thrust::uniform_int_distribution<_Tp> __dist{static_cast<_Tp>(__num), ::cuda::std::numeric_limits<_Tp>::max()}; - __rng.seed(__seed); - return __dist(__rng); + Rng rng; + ::thrust::uniform_int_distribution dist{static_cast(num), ::cuda::std::numeric_limits::max()}; + rng.seed(seed); + return dist(rng); } - ::cuda::std::size_t __num; + ::cuda::std::size_t num; }; -template -struct __dropout_pred +template +struct dropout_pred { - _CCCL_HOST_DEVICE_API constexpr explicit __dropout_pred(double __keep_prob) - : __keep_prob{__keep_prob} + _CCCL_HOST_DEVICE_API constexpr explicit dropout_pred(double keep_prob) + : keep_prob{keep_prob} {} - _CCCL_HOST_DEVICE_API bool operator()(::cuda::std::size_t __seed) const noexcept + _CCCL_HOST_DEVICE_API bool operator()(::cuda::std::size_t seed) const noexcept { - _Rng __rng; - ::thrust::uniform_real_distribution __dist{0.0, 1.0}; - __rng.seed(__seed); - return __dist(__rng) > __keep_prob; + Rng rng; + ::thrust::uniform_real_distribution dist{0.0, 1.0}; + rng.seed(seed); + return dist(rng) > keep_prob; } - double __keep_prob; + double keep_prob; }; } // namespace detail @@ -130,63 +130,63 @@ struct __dropout_pred //! //! The generator defaults to `defaults::seed` to keep benchmark data reproducible across runs. //! -//! @tparam _Rng Pseudo-random number generator type compatible with Thrust random distributions. -template +//! @tparam Rng Pseudo-random number generator type compatible with Thrust random distributions. +template class key_generator { public: //! Constructs a key generator with the given seed. //! - //! @param __seed Seed used to initialize the generator state. - explicit key_generator(::cuda::std::uint32_t __seed = defaults::seed) - : __rng{__seed} + //! @param seed Seed used to initialize the generator state. + explicit key_generator(::cuda::std::uint32_t seed = defaults::seed) + : rng{seed} {} //! Generates keys according to the given distribution using the default device execution policy. //! - //! @tparam _Dist Distribution tag type. - //! @tparam _OutputIt Output iterator type whose value type is the generated key type. - //! @param __dist Distribution tag controlling how keys are generated. - //! @param __out_begin Beginning of the output key range. - //! @param __out_end End of the output key range. - //! @throws std::invalid_argument if `_Dist` is not a supported distribution tag. - template - void generate(_Dist __dist, _OutputIt __out_begin, _OutputIt __out_end) + //! @tparam Distribution Distribution tag type. + //! @tparam OutputIt Output iterator type whose value type is the generated key type. + //! @param dist Distribution tag controlling how keys are generated. + //! @param out_begin Beginning of the output key range. + //! @param out_end End of the output key range. + //! @throws std::invalid_argument if `Distribution` is not a supported distribution tag. + template + void generate(Distribution dist, OutputIt out_begin, OutputIt out_end) { - generate(__dist, __out_begin, __out_end, ::thrust::device); + generate(dist, out_begin, out_end, ::thrust::device); } //! Generates keys according to the given distribution using the provided execution policy. //! - //! @tparam _Dist Distribution tag type. - //! @tparam _OutputIt Output iterator type whose value type is the generated key type. - //! @tparam _ExecPolicy Thrust execution policy type. - //! @param __dist Distribution tag controlling how keys are generated. - //! @param __out_begin Beginning of the output key range. - //! @param __out_end End of the output key range. - //! @param __exec_policy Execution policy used for the underlying Thrust algorithms. - //! @throws std::invalid_argument if `_Dist` is not a supported distribution tag. - template - void generate(_Dist __dist, _OutputIt __out_begin, _OutputIt __out_end, _ExecPolicy __exec_policy) + //! @tparam Distribution Distribution tag type. + //! @tparam OutputIt Output iterator type whose value type is the generated key type. + //! @tparam ExecPolicy Thrust execution policy type. + //! @param dist Distribution tag controlling how keys are generated. + //! @param out_begin Beginning of the output key range. + //! @param out_end End of the output key range. + //! @param exec_policy Execution policy used for the underlying Thrust algorithms. + //! @throws std::invalid_argument if `Distribution` is not a supported distribution tag. + template + void generate(Distribution dist, OutputIt out_begin, OutputIt out_end, ExecPolicy exec_policy) { - using __value_type = typename ::cuda::std::iterator_traits<_OutputIt>::value_type; + using value_type = typename ::cuda::std::iterator_traits::value_type; - if constexpr (::cuda::std::is_same_v<_Dist, distribution::unique>) + if constexpr (::cuda::std::is_same_v) { - ::thrust::sequence(__exec_policy, __out_begin, __out_end, __value_type{0}); - ::thrust::shuffle(__exec_policy, __out_begin, __out_end, __rng); + ::thrust::sequence(exec_policy, out_begin, out_end, value_type{0}); + ::thrust::shuffle(exec_policy, out_begin, out_end, rng); } - else if constexpr (::cuda::std::is_same_v<_Dist, distribution::uniform>) + else if constexpr (::cuda::std::is_same_v) { - const auto __num_keys = static_cast<::cuda::std::size_t>(::cuda::std::distance(__out_begin, __out_end)); - const auto __seed = static_cast<::cuda::std::size_t>(__rng()); + const auto num_keys = static_cast<::cuda::std::size_t>(::cuda::std::distance(out_begin, out_end)); + const auto seed = static_cast<::cuda::std::size_t>(rng()); ::thrust::transform( - __exec_policy, + exec_policy, ::cuda::counting_iterator<::cuda::std::size_t>{0}, - ::cuda::counting_iterator<::cuda::std::size_t>{__num_keys}, - __out_begin, - detail::__generate_uniform_fn<__value_type, _Dist, _Rng>{__num_keys, __dist, __seed}); + ::cuda::counting_iterator<::cuda::std::size_t>{num_keys}, + out_begin, + detail::generate_uniform_fn{num_keys, dist, seed}); } else { @@ -198,77 +198,81 @@ public: //! //! Replaced keys are sampled from `[N, max_key]`, where `N` is the number of keys in the range. //! - //! @tparam _InOutIt Mutable iterator type whose value type is the key type. - //! @param __begin Beginning of the key range to update in place. - //! @param __end End of the key range to update in place. - //! @param __keep_prob Probability of keeping each original key. Must be in `[0, 1]`. - //! @throws std::invalid_argument if `__keep_prob` is outside `[0, 1]`. - template - void dropout(_InOutIt __begin, _InOutIt __end, double __keep_prob) + //! The full range is shuffled afterward, even when all keys are kept. + //! + //! @tparam InOutIt Mutable iterator type whose value type is the key type. + //! @param begin Beginning of the key range to update in place. + //! @param end End of the key range to update in place. + //! @param keep_prob Probability of keeping each original key. Must be in `[0, 1]`. + //! @throws std::invalid_argument if `keep_prob` is outside `[0, 1]`. + template + void dropout(InOutIt begin, InOutIt end, double keep_prob) { - dropout(__begin, __end, __keep_prob, ::thrust::device); + dropout(begin, end, keep_prob, ::thrust::device); } //! Drops keys with probability `1 - keep_prob` using the provided execution policy. //! //! Replaced keys are sampled from `[N, max_key]`, where `N` is the number of keys in the range. //! - //! @tparam _InOutIt Mutable iterator type whose value type is the key type. - //! @tparam _ExecPolicy Thrust execution policy type. - //! @param __begin Beginning of the key range to update in place. - //! @param __end End of the key range to update in place. - //! @param __keep_prob Probability of keeping each original key. Must be in `[0, 1]`. - //! @param __exec_policy Execution policy used for the underlying Thrust algorithms. - //! @throws std::invalid_argument if `__keep_prob` is outside `[0, 1]`. - template - void dropout(_InOutIt __begin, _InOutIt __end, double __keep_prob, _ExecPolicy __exec_policy) + //! The full range is shuffled afterward, even when all keys are kept. + //! + //! @tparam InOutIt Mutable iterator type whose value type is the key type. + //! @tparam ExecPolicy Thrust execution policy type. + //! @param begin Beginning of the key range to update in place. + //! @param end End of the key range to update in place. + //! @param keep_prob Probability of keeping each original key. Must be in `[0, 1]`. + //! @param exec_policy Execution policy used for the underlying Thrust algorithms. + //! @throws std::invalid_argument if `keep_prob` is outside `[0, 1]`. + template + void dropout(InOutIt begin, InOutIt end, double keep_prob, ExecPolicy exec_policy) { - using __value_type = typename ::cuda::std::iterator_traits<_InOutIt>::value_type; + using value_type = typename ::cuda::std::iterator_traits::value_type; - if (__keep_prob < 0.0 || __keep_prob > 1.0) + if (keep_prob < 0.0 || keep_prob > 1.0) { throw ::std::invalid_argument{"Probability needs to be between 0 and 1"}; } - if (__keep_prob < 1.0) + if (keep_prob < 1.0) { - const auto __num_keys = static_cast<::cuda::std::size_t>(::cuda::std::distance(__begin, __end)); - ::cuda::counting_iterator<::cuda::std::size_t> __seeds{static_cast<::cuda::std::size_t>(__rng())}; + const auto num_keys = static_cast<::cuda::std::size_t>(::cuda::std::distance(begin, end)); + ::cuda::counting_iterator<::cuda::std::size_t> seeds{static_cast<::cuda::std::size_t>(rng())}; ::thrust::transform_if( - __exec_policy, - __seeds, - __seeds + __num_keys, - __begin, - detail::__dropout_fn<__value_type, _Rng>{__num_keys}, - detail::__dropout_pred<_Rng>{__keep_prob}); + exec_policy, + seeds, + seeds + num_keys, + begin, + detail::dropout_fn{num_keys}, + detail::dropout_pred{keep_prob}); } - ::thrust::shuffle(__exec_policy, __begin, __end, __rng); + ::thrust::shuffle(exec_policy, begin, end, rng); } private: - _Rng __rng; + Rng rng; }; //! Constructs the requested distribution tag from NVBench axis values. //! //! `distribution::uniform` reads the `Multiplicity` axis from `state`. //! -//! @tparam _Dist Distribution tag type to construct. -//! @param __state NVBench state containing distribution-specific axis values. +//! @tparam Distribution Distribution tag type to construct. +//! @param state NVBench state containing distribution-specific axis values. //! @return Distribution tag initialized from the benchmark state. -//! @throws std::invalid_argument if `_Dist` is not a supported distribution tag. -template -_Dist dist_from_state(::nvbench::state const& __state) +//! @throws std::invalid_argument if `Distribution` is not a supported distribution tag. +template +Distribution dist_from_state(::nvbench::state const& state) { - if constexpr (::cuda::std::is_same_v<_Dist, distribution::unique>) + if constexpr (::cuda::std::is_same_v) { - return _Dist{}; + return Distribution{}; } - else if constexpr (::cuda::std::is_same_v<_Dist, distribution::uniform>) + else if constexpr (::cuda::std::is_same_v) { - return _Dist{__state.get_int64("Multiplicity")}; + return Distribution{state.get_int64("Multiplicity")}; } else { diff --git a/cudax/test/cuco/fixed_capacity_map/test_shared_memory.cu b/cudax/test/cuco/fixed_capacity_map/test_shared_memory.cu index d9675dd2cd8..234ea2a2d2e 100644 --- a/cudax/test/cuco/fixed_capacity_map/test_shared_memory.cu +++ b/cudax/test/cuco/fixed_capacity_map/test_shared_memory.cu @@ -30,27 +30,26 @@ namespace cudax = cuda::experimental; constexpr int empty_key = -1; constexpr int empty_value = -1; -// A static-capacity map. `_Capacity` must be a valid slot count, so it is computed from the probing -// scheme and bucket size (the map's defaults: linear probing, cg_size 1, bucket_size 1) rather than -// hand-written. -using default_probing = cudax::cuco::linear_probing<1, cudax::cuco::hash>; -inline constexpr int default_bucket = 1; +// A static-capacity map with cg_size 1 so the test can use scalar device inserts. +using probing = cudax::cuco::linear_probing<1, cudax::cuco::hash>; +inline constexpr int bucket = 1; inline constexpr ::cuda::std::size_t static_capacity = - cudax::cuco::make_valid_capacity(::cuda::std::size_t{512}); -using fixed_capacity_map_512_type = cudax::cuco::fixed_capacity_map; + cudax::cuco::make_valid_capacity(::cuda::std::size_t{512}); +using fixed_capacity_map_512_type = cudax::cuco:: + fixed_capacity_map, probing>; -template +template struct iota_pair { - __host__ __device__ _Pair operator()(typename _Pair::first_type __i) const noexcept + __host__ __device__ Pair operator()(typename Pair::first_type key) const noexcept { - return _Pair{__i, __i}; + return Pair{key, key}; } }; // Demonstrates compile-time __shared__ sizing via ref_type::capacity_v. -template -__global__ void insert_shmem_kernel(fixed_capacity_map_512_type::ref_type global_ref, _PairIt pairs, int n) +template +__global__ void insert_shmem_kernel(fixed_capacity_map_512_type::ref_type global_ref, PairIt pairs, int num_keys) { using ref_t = fixed_capacity_map_512_type::ref_type; static_assert(ref_t::capacity_v != ::cuda::std::dynamic_extent, @@ -59,9 +58,9 @@ __global__ void insert_shmem_kernel(fixed_capacity_map_512_type::ref_type global __shared__ ::cuda::__uninitialized_array smem; const auto idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; - smem[threadIdx.x] = (idx < n) ? pairs[idx] : ref_t::value_type{}; + smem[threadIdx.x] = (idx < num_keys) ? pairs[idx] : ref_t::value_type{}; __syncthreads(); - if (idx < n) + if (idx < num_keys) { global_ref.insert(smem[threadIdx.x]); } From 1e75578ed94ae83fc854fc5d554e2a1e7c103649 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Wed, 8 Jul 2026 15:38:59 -0700 Subject: [PATCH 5/8] Temporary workaround for #9755 --- .../bench/cuco/common/key_generator.cuh | 22 +++++++++++++++++-- 1 file changed, 20 insertions(+), 2 deletions(-) diff --git a/cudax/benchmarks/bench/cuco/common/key_generator.cuh b/cudax/benchmarks/bench/cuco/common/key_generator.cuh index 05eaab6c6ad..58c62d703ac 100644 --- a/cudax/benchmarks/bench/cuco/common/key_generator.cuh +++ b/cudax/benchmarks/bench/cuco/common/key_generator.cuh @@ -17,12 +17,14 @@ #include #include #include +#include #include #include #include #include #include +#include #include #include @@ -61,6 +63,22 @@ struct uniform namespace detail { +template +void shuffle(ExecPolicy exec_policy, RandomIt begin, RandomIt end, Rng& rng) +{ + const auto num_items = ::cuda::std::distance(begin, end); + if constexpr (::thrust::is_contiguous_iterator_v) + { + // TODO(NVIDIA/cccl#9755): Temporary workaround for thrust::shuffle's &*it contiguous unwrap. + auto shuffle_begin = ::cuda::std::to_address(begin); + ::thrust::shuffle(exec_policy, shuffle_begin, shuffle_begin + num_items, rng); + } + else + { + ::thrust::shuffle(exec_policy, begin, end, rng); + } +} + template struct generate_uniform_fn { @@ -174,7 +192,7 @@ public: if constexpr (::cuda::std::is_same_v) { ::thrust::sequence(exec_policy, out_begin, out_end, value_type{0}); - ::thrust::shuffle(exec_policy, out_begin, out_end, rng); + detail::shuffle(exec_policy, out_begin, out_end, rng); } else if constexpr (::cuda::std::is_same_v) { @@ -248,7 +266,7 @@ public: detail::dropout_pred{keep_prob}); } - ::thrust::shuffle(exec_policy, begin, end, rng); + detail::shuffle(exec_policy, begin, end, rng); } private: From 63f0a58dfbd036b33ea7d7023b9130e4d903b1bd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Thu, 9 Jul 2026 15:03:18 -0700 Subject: [PATCH 6/8] Address fixed capacity benchmark review --- .../bench/cuco/common/key_generator.cuh | 16 +++++++--------- .../bench/cuco/fixed_capacity_map/contains.cu | 2 +- .../bench/cuco/fixed_capacity_map/insert.cu | 2 +- 3 files changed, 9 insertions(+), 11 deletions(-) diff --git a/cudax/benchmarks/bench/cuco/common/key_generator.cuh b/cudax/benchmarks/bench/cuco/common/key_generator.cuh index 58c62d703ac..5b600bdd5a8 100644 --- a/cudax/benchmarks/bench/cuco/common/key_generator.cuh +++ b/cudax/benchmarks/bench/cuco/common/key_generator.cuh @@ -10,8 +10,6 @@ #pragma once -#include - #include #include #include @@ -22,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -82,14 +81,13 @@ void shuffle(ExecPolicy exec_policy, RandomIt begin, RandomIt end, Rng& rng) template struct generate_uniform_fn { - _CCCL_HOST_DEVICE_API constexpr generate_uniform_fn( - ::cuda::std::size_t num, Distribution dist, ::cuda::std::size_t seed) + __host__ __device__ constexpr generate_uniform_fn(::cuda::std::size_t num, Distribution dist, ::cuda::std::size_t seed) : num{num} , dist{dist} , seed{seed} {} - _CCCL_HOST_DEVICE_API constexpr Key operator()(::cuda::std::size_t idx) const noexcept + __host__ __device__ constexpr Key operator()(::cuda::std::size_t idx) const noexcept { Rng rng; rng.seed(seed + idx * 1664525ull + 1013904223ull); @@ -110,11 +108,11 @@ struct generate_uniform_fn template struct dropout_fn { - _CCCL_HOST_DEVICE_API constexpr explicit dropout_fn(::cuda::std::size_t num) + __host__ __device__ constexpr explicit dropout_fn(::cuda::std::size_t num) : num{num} {} - _CCCL_HOST_DEVICE_API Key operator()(::cuda::std::size_t seed) const noexcept + __host__ __device__ Key operator()(::cuda::std::size_t seed) const noexcept { Rng rng; ::thrust::uniform_int_distribution dist{static_cast(num), ::cuda::std::numeric_limits::max()}; @@ -128,11 +126,11 @@ struct dropout_fn template struct dropout_pred { - _CCCL_HOST_DEVICE_API constexpr explicit dropout_pred(double keep_prob) + __host__ __device__ constexpr explicit dropout_pred(double keep_prob) : keep_prob{keep_prob} {} - _CCCL_HOST_DEVICE_API bool operator()(::cuda::std::size_t seed) const noexcept + __host__ __device__ bool operator()(::cuda::std::size_t seed) const noexcept { Rng rng; ::thrust::uniform_real_distribution dist{0.0, 1.0}; diff --git a/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu b/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu index 60b7320bcd6..b22de952d58 100644 --- a/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu +++ b/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu @@ -36,7 +36,7 @@ void fixed_capacity_map_contains(::nvbench::state& state, ::nvbench::type_list Date: Thu, 9 Jul 2026 15:12:57 -0700 Subject: [PATCH 7/8] Use user-facing namespaces in cuco benchmarks --- .../bench/cuco/common/key_generator.cuh | 86 +++++++++---------- .../bench/cuco/fixed_capacity_map/contains.cu | 32 +++---- .../bench/cuco/fixed_capacity_map/insert.cu | 30 +++---- 3 files changed, 74 insertions(+), 74 deletions(-) diff --git a/cudax/benchmarks/bench/cuco/common/key_generator.cuh b/cudax/benchmarks/bench/cuco/common/key_generator.cuh index 5b600bdd5a8..8811bf49752 100644 --- a/cudax/benchmarks/bench/cuco/common/key_generator.cuh +++ b/cudax/benchmarks/bench/cuco/common/key_generator.cuh @@ -46,7 +46,7 @@ struct uniform //! //! @param multiplicity Average number of generated keys that map to each unique key value. //! @throws std::invalid_argument if `multiplicity` is not positive. - explicit uniform(::nvbench::int64_t multiplicity) + explicit uniform(nvbench::int64_t multiplicity) : multiplicity{multiplicity} { if (multiplicity <= 0) @@ -56,7 +56,7 @@ struct uniform } //! Average number of generated keys that map to each unique key value. - ::nvbench::int64_t multiplicity; + nvbench::int64_t multiplicity; }; } // namespace distribution @@ -65,62 +65,62 @@ namespace detail template void shuffle(ExecPolicy exec_policy, RandomIt begin, RandomIt end, Rng& rng) { - const auto num_items = ::cuda::std::distance(begin, end); - if constexpr (::thrust::is_contiguous_iterator_v) + const auto num_items = cuda::std::distance(begin, end); + if constexpr (thrust::is_contiguous_iterator_v) { // TODO(NVIDIA/cccl#9755): Temporary workaround for thrust::shuffle's &*it contiguous unwrap. - auto shuffle_begin = ::cuda::std::to_address(begin); - ::thrust::shuffle(exec_policy, shuffle_begin, shuffle_begin + num_items, rng); + auto shuffle_begin = cuda::std::to_address(begin); + thrust::shuffle(exec_policy, shuffle_begin, shuffle_begin + num_items, rng); } else { - ::thrust::shuffle(exec_policy, begin, end, rng); + thrust::shuffle(exec_policy, begin, end, rng); } } template struct generate_uniform_fn { - __host__ __device__ constexpr generate_uniform_fn(::cuda::std::size_t num, Distribution dist, ::cuda::std::size_t seed) + __host__ __device__ constexpr generate_uniform_fn(cuda::std::size_t num, Distribution dist, cuda::std::size_t seed) : num{num} , dist{dist} , seed{seed} {} - __host__ __device__ constexpr Key operator()(::cuda::std::size_t idx) const noexcept + __host__ __device__ constexpr Key operator()(cuda::std::size_t idx) const noexcept { Rng rng; rng.seed(seed + idx * 1664525ull + 1013904223ull); - const auto num_unique_keys_unclamped = static_cast<::cuda::std::size_t>( - ::cuda::std::ceil(static_cast(num) / static_cast(dist.multiplicity))); + const auto num_unique_keys_unclamped = static_cast( + cuda::std::ceil(static_cast(num) / static_cast(dist.multiplicity))); const auto num_unique_keys = - num_unique_keys_unclamped < ::cuda::std::size_t{1} ? ::cuda::std::size_t{1} : num_unique_keys_unclamped; - ::thrust::uniform_int_distribution key_dist{Key{0}, static_cast(num_unique_keys - 1)}; + num_unique_keys_unclamped < cuda::std::size_t{1} ? cuda::std::size_t{1} : num_unique_keys_unclamped; + thrust::uniform_int_distribution key_dist{Key{0}, static_cast(num_unique_keys - 1)}; return key_dist(rng); } - ::cuda::std::size_t num; + cuda::std::size_t num; Distribution dist; - ::cuda::std::size_t seed; + cuda::std::size_t seed; }; template struct dropout_fn { - __host__ __device__ constexpr explicit dropout_fn(::cuda::std::size_t num) + __host__ __device__ constexpr explicit dropout_fn(cuda::std::size_t num) : num{num} {} - __host__ __device__ Key operator()(::cuda::std::size_t seed) const noexcept + __host__ __device__ Key operator()(cuda::std::size_t seed) const noexcept { Rng rng; - ::thrust::uniform_int_distribution dist{static_cast(num), ::cuda::std::numeric_limits::max()}; + thrust::uniform_int_distribution dist{static_cast(num), cuda::std::numeric_limits::max()}; rng.seed(seed); return dist(rng); } - ::cuda::std::size_t num; + cuda::std::size_t num; }; template @@ -130,10 +130,10 @@ struct dropout_pred : keep_prob{keep_prob} {} - __host__ __device__ bool operator()(::cuda::std::size_t seed) const noexcept + __host__ __device__ bool operator()(cuda::std::size_t seed) const noexcept { Rng rng; - ::thrust::uniform_real_distribution dist{0.0, 1.0}; + thrust::uniform_real_distribution dist{0.0, 1.0}; rng.seed(seed); return dist(rng) > keep_prob; } @@ -147,14 +147,14 @@ struct dropout_pred //! The generator defaults to `defaults::seed` to keep benchmark data reproducible across runs. //! //! @tparam Rng Pseudo-random number generator type compatible with Thrust random distributions. -template +template class key_generator { public: //! Constructs a key generator with the given seed. //! //! @param seed Seed used to initialize the generator state. - explicit key_generator(::cuda::std::uint32_t seed = defaults::seed) + explicit key_generator(cuda::std::uint32_t seed = defaults::seed) : rng{seed} {} @@ -169,7 +169,7 @@ public: template void generate(Distribution dist, OutputIt out_begin, OutputIt out_end) { - generate(dist, out_begin, out_end, ::thrust::device); + generate(dist, out_begin, out_end, thrust::device); } //! Generates keys according to the given distribution using the provided execution policy. @@ -185,22 +185,22 @@ public: template void generate(Distribution dist, OutputIt out_begin, OutputIt out_end, ExecPolicy exec_policy) { - using value_type = typename ::cuda::std::iterator_traits::value_type; + using value_type = typename cuda::std::iterator_traits::value_type; - if constexpr (::cuda::std::is_same_v) + if constexpr (cuda::std::is_same_v) { - ::thrust::sequence(exec_policy, out_begin, out_end, value_type{0}); + thrust::sequence(exec_policy, out_begin, out_end, value_type{0}); detail::shuffle(exec_policy, out_begin, out_end, rng); } - else if constexpr (::cuda::std::is_same_v) + else if constexpr (cuda::std::is_same_v) { - const auto num_keys = static_cast<::cuda::std::size_t>(::cuda::std::distance(out_begin, out_end)); - const auto seed = static_cast<::cuda::std::size_t>(rng()); + const auto num_keys = static_cast(cuda::std::distance(out_begin, out_end)); + const auto seed = static_cast(rng()); - ::thrust::transform( + thrust::transform( exec_policy, - ::cuda::counting_iterator<::cuda::std::size_t>{0}, - ::cuda::counting_iterator<::cuda::std::size_t>{num_keys}, + cuda::counting_iterator{0}, + cuda::counting_iterator{num_keys}, out_begin, detail::generate_uniform_fn{num_keys, dist, seed}); } @@ -224,7 +224,7 @@ public: template void dropout(InOutIt begin, InOutIt end, double keep_prob) { - dropout(begin, end, keep_prob, ::thrust::device); + dropout(begin, end, keep_prob, thrust::device); } //! Drops keys with probability `1 - keep_prob` using the provided execution policy. @@ -243,7 +243,7 @@ public: template void dropout(InOutIt begin, InOutIt end, double keep_prob, ExecPolicy exec_policy) { - using value_type = typename ::cuda::std::iterator_traits::value_type; + using value_type = typename cuda::std::iterator_traits::value_type; if (keep_prob < 0.0 || keep_prob > 1.0) { @@ -252,10 +252,10 @@ public: if (keep_prob < 1.0) { - const auto num_keys = static_cast<::cuda::std::size_t>(::cuda::std::distance(begin, end)); - ::cuda::counting_iterator<::cuda::std::size_t> seeds{static_cast<::cuda::std::size_t>(rng())}; + const auto num_keys = static_cast(cuda::std::distance(begin, end)); + cuda::counting_iterator seeds{static_cast(rng())}; - ::thrust::transform_if( + thrust::transform_if( exec_policy, seeds, seeds + num_keys, @@ -280,13 +280,13 @@ private: //! @return Distribution tag initialized from the benchmark state. //! @throws std::invalid_argument if `Distribution` is not a supported distribution tag. template -Distribution dist_from_state(::nvbench::state const& state) +Distribution dist_from_state(nvbench::state const& state) { - if constexpr (::cuda::std::is_same_v) + if constexpr (cuda::std::is_same_v) { return Distribution{}; } - else if constexpr (::cuda::std::is_same_v) + else if constexpr (cuda::std::is_same_v) { return Distribution{state.get_int64("Multiplicity")}; } @@ -298,6 +298,6 @@ Distribution dist_from_state(::nvbench::state const& state) } // namespace cuda::experimental::cuco::benchmark NVBENCH_DECLARE_TYPE_STRINGS( - ::cuda::experimental::cuco::benchmark::distribution::unique, "UNIQUE", "distribution::unique"); + cuda::experimental::cuco::benchmark::distribution::unique, "UNIQUE", "distribution::unique"); NVBENCH_DECLARE_TYPE_STRINGS( - ::cuda::experimental::cuco::benchmark::distribution::uniform, "UNIFORM", "distribution::uniform"); + cuda::experimental::cuco::benchmark::distribution::uniform, "UNIFORM", "distribution::uniform"); diff --git a/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu b/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu index b22de952d58..b3453504b46 100644 --- a/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu +++ b/cudax/benchmarks/bench/cuco/fixed_capacity_map/contains.cu @@ -25,14 +25,14 @@ #include "../common/key_generator.cuh" #include -namespace cudax = ::cuda::experimental; +namespace cudax = cuda::experimental; namespace bench = cudax::cuco::benchmark; /** * @brief A benchmark evaluating `cudax::cuco::fixed_capacity_map::contains_async` performance. */ template -void fixed_capacity_map_contains(::nvbench::state& state, ::nvbench::type_list) +void fixed_capacity_map_contains(nvbench::state& state, nvbench::type_list) { if constexpr (sizeof(Key) != sizeof(Value)) { @@ -40,27 +40,27 @@ void fixed_capacity_map_contains(::nvbench::state& state, ::nvbench::type_list; + using pair_type = cuda::std::pair; using map_type = cudax::cuco::fixed_capacity_map; const auto num_keys = state.get_int64("NumInputs"); const auto occupancy = state.get_float64("Occupancy"); const auto matching_rate = state.get_float64("MatchingRate"); - const auto size = static_cast<::cuda::std::size_t>(static_cast(num_keys) / occupancy); + const auto size = static_cast(static_cast(num_keys) / occupancy); - const auto device = ::cuda::device_ref{0}; - ::cuda::stream stream{device}; - const ::cuda::device_memory_pool_ref mr = ::cuda::device_default_memory_pool(device); - const auto exec_policy = ::thrust::cuda::par_nosync.on(stream.get()); + const auto device = cuda::device_ref{0}; + cuda::stream stream{device}; + const cuda::device_memory_pool_ref mr = cuda::device_default_memory_pool(device); + const auto exec_policy = thrust::cuda::par_nosync.on(stream.get()); - auto keys = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); + auto keys = cuda::make_device_buffer(stream, device, num_keys, cuda::no_init); bench::key_generator gen{}; gen.generate(bench::dist_from_state(state), keys.begin(), keys.end(), exec_policy); - auto pairs = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); - ::thrust::transform(exec_policy, keys.begin(), keys.end(), pairs.begin(), [] __device__(Key const& key) { + auto pairs = cuda::make_device_buffer(stream, device, num_keys, cuda::no_init); + thrust::transform(exec_policy, keys.begin(), keys.end(), pairs.begin(), [] __device__(Key const& key) { return pair_type{key, Value{}}; }); @@ -69,11 +69,11 @@ void fixed_capacity_map_contains(::nvbench::state& state, ::nvbench::type_list(stream, device, num_keys, ::cuda::no_init); + auto result = cuda::make_device_buffer(stream, device, num_keys, cuda::no_init); stream.sync(); state.add_element_count(num_keys); - state.exec([&](::nvbench::launch& launch) { + state.exec([&](nvbench::launch& launch) { map.contains_async({launch.get_stream()}, keys.begin(), keys.end(), result.begin()); }); } @@ -82,7 +82,7 @@ void fixed_capacity_map_contains(::nvbench::state& state, ::nvbench::type_list)) + nvbench::type_list)) .set_name("fixed_capacity_map_contains_unique_capacity") .set_type_axes_names({"Key", "Value", "Distribution"}) .add_int64_axis("NumInputs", bench::defaults::n_range_cache) @@ -92,7 +92,7 @@ NVBENCH_BENCH_TYPES(fixed_capacity_map_contains, NVBENCH_BENCH_TYPES(fixed_capacity_map_contains, NVBENCH_TYPE_AXES(bench::defaults::key_type_range, bench::defaults::value_type_range, - ::nvbench::type_list)) + nvbench::type_list)) .set_name("fixed_capacity_map_contains_unique_occupancy") .set_type_axes_names({"Key", "Value", "Distribution"}) .add_int64_axis("NumInputs", {bench::defaults::n}) @@ -102,7 +102,7 @@ NVBENCH_BENCH_TYPES(fixed_capacity_map_contains, NVBENCH_BENCH_TYPES(fixed_capacity_map_contains, NVBENCH_TYPE_AXES(bench::defaults::key_type_range, bench::defaults::value_type_range, - ::nvbench::type_list)) + nvbench::type_list)) .set_name("fixed_capacity_map_contains_unique_matching_rate") .set_type_axes_names({"Key", "Value", "Distribution"}) .add_int64_axis("NumInputs", {bench::defaults::n}) diff --git a/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu b/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu index cc594827bd5..0c3637b6abd 100644 --- a/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu +++ b/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu @@ -25,14 +25,14 @@ #include "../common/key_generator.cuh" #include -namespace cudax = ::cuda::experimental; +namespace cudax = cuda::experimental; namespace bench = cudax::cuco::benchmark; /** * @brief A benchmark evaluating `cudax::cuco::fixed_capacity_map::insert_async` performance. */ template -void fixed_capacity_map_insert(::nvbench::state& state, ::nvbench::type_list) +void fixed_capacity_map_insert(nvbench::state& state, nvbench::type_list) { if constexpr (sizeof(Key) != sizeof(Value)) { @@ -40,26 +40,26 @@ void fixed_capacity_map_insert(::nvbench::state& state, ::nvbench::type_list; + using pair_type = cuda::std::pair; using map_type = cudax::cuco::fixed_capacity_map; const auto num_keys = state.get_int64("NumInputs"); const auto occupancy = state.get_float64("Occupancy"); - const auto size = static_cast<::cuda::std::size_t>(static_cast(num_keys) / occupancy); + const auto size = static_cast(static_cast(num_keys) / occupancy); - const auto device = ::cuda::device_ref{0}; - ::cuda::stream stream{device}; - const ::cuda::device_memory_pool_ref mr = ::cuda::device_default_memory_pool(device); - const auto exec_policy = ::thrust::cuda::par_nosync.on(stream.get()); + const auto device = cuda::device_ref{0}; + cuda::stream stream{device}; + const cuda::device_memory_pool_ref mr = cuda::device_default_memory_pool(device); + const auto exec_policy = thrust::cuda::par_nosync.on(stream.get()); - auto keys = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); + auto keys = cuda::make_device_buffer(stream, device, num_keys, cuda::no_init); bench::key_generator gen{}; gen.generate(bench::dist_from_state(state), keys.begin(), keys.end(), exec_policy); - auto pairs = ::cuda::make_device_buffer(stream, device, num_keys, ::cuda::no_init); - ::thrust::transform(exec_policy, keys.begin(), keys.end(), pairs.begin(), [] __device__(Key const& key) { + auto pairs = cuda::make_device_buffer(stream, device, num_keys, cuda::no_init); + thrust::transform(exec_policy, keys.begin(), keys.end(), pairs.begin(), [] __device__(Key const& key) { return pair_type{key, Value{}}; }); @@ -67,7 +67,7 @@ void fixed_capacity_map_insert(::nvbench::state& state, ::nvbench::type_list)) + nvbench::type_list)) .set_name("fixed_capacity_map_insert_unique_capacity") .set_type_axes_names({"Key", "Value", "Distribution"}) .add_int64_axis("NumInputs", bench::defaults::n_range_cache) @@ -88,7 +88,7 @@ NVBENCH_BENCH_TYPES(fixed_capacity_map_insert, NVBENCH_BENCH_TYPES(fixed_capacity_map_insert, NVBENCH_TYPE_AXES(bench::defaults::key_type_range, bench::defaults::value_type_range, - ::nvbench::type_list)) + nvbench::type_list)) .set_name("fixed_capacity_map_insert_unique_occupancy") .set_type_axes_names({"Key", "Value", "Distribution"}) .add_int64_axis("NumInputs", {bench::defaults::n}) @@ -97,7 +97,7 @@ NVBENCH_BENCH_TYPES(fixed_capacity_map_insert, NVBENCH_BENCH_TYPES(fixed_capacity_map_insert, NVBENCH_TYPE_AXES(bench::defaults::key_type_range, bench::defaults::value_type_range, - ::nvbench::type_list)) + nvbench::type_list)) .set_name("fixed_capacity_map_insert_uniform_multiplicity") .set_type_axes_names({"Key", "Value", "Distribution"}) .add_int64_axis("NumInputs", {bench::defaults::n}) From 994c4acdfde4c9378e61a2ee182e798fbdf864c7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20J=C3=BCnger?= Date: Thu, 9 Jul 2026 15:28:43 -0700 Subject: [PATCH 8/8] Use floating point uniform multiplicity --- cudax/benchmarks/bench/cuco/common/defaults.cuh | 2 +- .../bench/cuco/common/key_generator.cuh | 16 ++++++++-------- .../bench/cuco/fixed_capacity_map/insert.cu | 2 +- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/cudax/benchmarks/bench/cuco/common/defaults.cuh b/cudax/benchmarks/bench/cuco/common/defaults.cuh index 8b946a6e719..72fd69620f7 100644 --- a/cudax/benchmarks/bench/cuco/common/defaults.cuh +++ b/cudax/benchmarks/bench/cuco/common/defaults.cuh @@ -38,7 +38,7 @@ inline const auto n_range_cache = ::std::vector<::nvbench::int64_t>{8'000, 80'00 //! Occupancy sweep used by fixed-capacity container benchmarks. inline const auto occupancy_range = ::nvbench::range(0.1, 0.9, 0.1); //! Average multiplicity sweep for duplicate-key distributions. -inline const auto multiplicity_range = ::std::vector<::nvbench::int64_t>{1, 2, 4, 8, 16}; +inline const auto multiplicity_range = ::std::vector{1.0, 2.0, 4.0, 8.0, 16.0}; //! Matching-rate sweep used by contains-style benchmarks. inline const auto matching_rate_range = ::nvbench::range(0.1, 1.0, 0.1); } // namespace cuda::experimental::cuco::benchmark::defaults diff --git a/cudax/benchmarks/bench/cuco/common/key_generator.cuh b/cudax/benchmarks/bench/cuco/common/key_generator.cuh index 8811bf49752..fba63bfa307 100644 --- a/cudax/benchmarks/bench/cuco/common/key_generator.cuh +++ b/cudax/benchmarks/bench/cuco/common/key_generator.cuh @@ -45,18 +45,18 @@ struct uniform //! Constructs a uniform distribution tag with the requested average key multiplicity. //! //! @param multiplicity Average number of generated keys that map to each unique key value. - //! @throws std::invalid_argument if `multiplicity` is not positive. - explicit uniform(nvbench::int64_t multiplicity) + //! @throws std::invalid_argument if `multiplicity` is not finite or is less than 1.0. + explicit uniform(double multiplicity) : multiplicity{multiplicity} { - if (multiplicity <= 0) + if (!cuda::std::isfinite(multiplicity) || multiplicity < 1.0) { - throw ::std::invalid_argument{"Multiplicity must be greater than 0"}; + throw ::std::invalid_argument{"Multiplicity must be finite and at least 1"}; } } //! Average number of generated keys that map to each unique key value. - nvbench::int64_t multiplicity; + double multiplicity; }; } // namespace distribution @@ -92,8 +92,8 @@ struct generate_uniform_fn Rng rng; rng.seed(seed + idx * 1664525ull + 1013904223ull); - const auto num_unique_keys_unclamped = static_cast( - cuda::std::ceil(static_cast(num) / static_cast(dist.multiplicity))); + const auto num_unique_keys_unclamped = + static_cast(cuda::std::ceil(static_cast(num) / dist.multiplicity)); const auto num_unique_keys = num_unique_keys_unclamped < cuda::std::size_t{1} ? cuda::std::size_t{1} : num_unique_keys_unclamped; thrust::uniform_int_distribution key_dist{Key{0}, static_cast(num_unique_keys - 1)}; @@ -288,7 +288,7 @@ Distribution dist_from_state(nvbench::state const& state) } else if constexpr (cuda::std::is_same_v) { - return Distribution{state.get_int64("Multiplicity")}; + return Distribution{state.get_float64("Multiplicity")}; } else { diff --git a/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu b/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu index 0c3637b6abd..607e914ecec 100644 --- a/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu +++ b/cudax/benchmarks/bench/cuco/fixed_capacity_map/insert.cu @@ -102,4 +102,4 @@ NVBENCH_BENCH_TYPES(fixed_capacity_map_insert, .set_type_axes_names({"Key", "Value", "Distribution"}) .add_int64_axis("NumInputs", {bench::defaults::n}) .add_float64_axis("Occupancy", {bench::defaults::occupancy}) - .add_int64_axis("Multiplicity", bench::defaults::multiplicity_range); + .add_float64_axis("Multiplicity", bench::defaults::multiplicity_range);