Replace cuco hashers with cuda::hash

README.md

@@ -62,6 +62,7 @@ This will take care of downloading `cuCollections` from GitHub and making the he
 ## Requirements
 - NVCC 12.0 or newer
 - C++17
+- CCCL 3.6.0 or newer
 - GPU Architecture: Volta or newer
     - Pascal is partially supported. Any data structures that require blocking algorithms are not supported. See [libcu++](https://nvidia.github.io/libcudacxx/setup/requirements.html#device-architectures) documentation for more details.
 

benchmarks/benchmark_defaults.hpp

@@ -10,6 +10,8 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/functional>
+
 #include <cstdint>
 #include <vector>
 
@@ -22,12 +24,9 @@ using VALUE_TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t,
 using KEY_TYPE_RANGE   = nvbench::type_list<nvbench::int32_t, nvbench::int64_t>;
 using VALUE_TYPE_RANGE = nvbench::type_list<nvbench::int32_t, nvbench::int64_t>;
 #endif
-using HASH_RANGE = nvbench::type_list<cuco::identity_hash<char>,
-                                      cuco::xxhash_32<char>,
-                                      cuco::xxhash_64<char>,
-                                      cuco::murmurhash3_32<char>>;  //,
-// cuco::murmurhash3_x86_128<char>,
-// cuco::murmurhash3_x64_128<char>>; // TODO handle tuple-like hash value
+using HASH_RANGE = nvbench::type_list<cuda::hash<char, cuda::hash_algorithm::xxhash_32>,
+                                      cuda::hash<char, cuda::hash_algorithm::xxhash_64>,
+                                      cuda::hash<char, cuda::hash_algorithm::murmurhash3_32>>;
 
 auto constexpr N             = 100'000'000;
 auto constexpr OCCUPANCY     = 0.5;

benchmarks/benchmark_utils.hpp

@@ -12,6 +12,7 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/functional>
 #include <cuda/iterator>
 
 #include <nv/target>
@@ -35,13 +36,21 @@ auto dist_from_state(nvbench::state const& state)
 }
 
 template <typename T, typename NewType>
-struct rebind_hasher;
+struct rebind_unary_hasher;
 
 template <template <typename> class Template, typename OldType, typename NewType>
-struct rebind_hasher<Template<OldType>, NewType> {
+struct rebind_unary_hasher<Template<OldType>, NewType> {
   using type = Template<NewType>;
 };
 
+template <typename T, typename NewType>
+struct rebind_hasher : rebind_unary_hasher<T, NewType> {};
+
+template <typename OldType, cuda::hash_algorithm Algorithm, typename NewType>
+struct rebind_hasher<cuda::hash<OldType, Algorithm>, NewType> {
+  using type = cuda::hash<NewType, Algorithm>;
+};
+
 template <typename T, typename NewType>
 using rebind_hasher_t = typename rebind_hasher<T, NewType>::type;
 
@@ -86,20 +95,29 @@ NVBENCH_DECLARE_TYPE_STRINGS(cuco::utility::distribution::gaussian,
                              "GAUSSIAN",
                              "distribution::gaussian");
 
-NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::XXHash_64<char>, "xxhash_64", "cuco::xxhash_64");
-NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::XXHash_32<char>, "xxhash_32", "cuco::xxhash_32");
-NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::MurmurHash3_32<char>,
-                             "murmurhash3_32",
-                             "cuco::murmurhash3_32");
-NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::MurmurHash3_x86_128<char>,
-                             "murmurhash3_x86_128",
-                             "cuco::murmurhash3_x86_128");
-NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::MurmurHash3_x64_128<char>,
-                             "murmurhash3_x64_128",
-                             "cuco::murmurhash3_x64_128");
-NVBENCH_DECLARE_TYPE_STRINGS(cuco::detail::identity_hash<char>,
-                             "identity_hash",
-                             "cuco::identity_hash");
+namespace nvbench {
+
+template <cuda::hash_algorithm Algorithm>
+struct type_strings<cuda::hash<char, Algorithm>> {
+  static std::string input_string()
+  {
+    switch (Algorithm) {
+      case cuda::hash_algorithm::xxhash_32: return "xxhash_32";
+      case cuda::hash_algorithm::xxhash_64: return "xxhash_64";
+      case cuda::hash_algorithm::murmurhash3_32: return "murmurhash3_32";
+      case cuda::hash_algorithm::murmurhash3_x86_128: return "murmurhash3_x86_128";
+      case cuda::hash_algorithm::murmurhash3_x64_128: return "murmurhash3_x64_128";
+    }
+    CUCO_FAIL("Unexpected hash algorithm");
+  }
+
+  static std::string description()
+  {
+    return "cuda::hash<char, cuda::hash_algorithm::" + input_string() + ">";
+  }
+};
+
+}  // namespace nvbench
 
 #if defined(CUCO_HAS_128BIT_ATOMICS)
 NVBENCH_DECLARE_TYPE_STRINGS(__int128_t, "I128", "__int128_t");

benchmarks/bloom_filter/add_bench.cu

@@ -12,6 +12,7 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/functional>
 #include <cuda/std/limits>
 #include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
@@ -59,7 +60,7 @@ void bloom_filter_add(nvbench::state& state,
     auto constexpr contains_vertical_layout   = words_per_block;
     auto constexpr contains_horizontal_layout = 1;
     using policy_type                         = cuco::bloom_filter_policy<Key,
-                                                                          cuco::xxhash_64<Key>,
+                                                                          cuda::hash<Key, cuda::hash_algorithm::xxhash_64>,
                                                                           WordBytes,
                                                                           words_per_block,
                                                                           PatternBits,

benchmarks/bloom_filter/contains_bench.cu

@@ -12,6 +12,7 @@
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/functional>
 #include <cuda/std/limits>
 #include <thrust/count.h>
 #include <thrust/device_vector.h>
@@ -58,7 +59,7 @@ void bloom_filter_contains(nvbench::state& state,
     auto constexpr add_vertical_layout   = 1;
     auto constexpr add_horizontal_layout = words_per_block;
     using policy_type                    = cuco::bloom_filter_policy<Key,
-                                                                     cuco::xxhash_64<Key>,
+                                                                     cuda::hash<Key, cuda::hash_algorithm::xxhash_64>,
                                                                      WordBytes,
                                                                      words_per_block,
                                                                      PatternBits,

benchmarks/hash_function/hash_function_bench.cu

@@ -5,18 +5,20 @@
 
 #include <benchmark_defaults.hpp>
 
+#include <cuco/detail/utility/cuda.cuh>
 #include <cuco/hash_functions.cuh>
 #include <cuco/utility/key_generator.cuh>
 
 #include <nvbench/nvbench.cuh>
 
+#include <cuda/functional>
 #include <cuda/std/cstddef>
+#include <cuda/std/type_traits>
 #include <thrust/device_vector.h>
 
 #include <memory_stats.hpp>
 
 #include <cstdint>
-#include <type_traits>
 
 // repeat hash computation n times
 static constexpr auto n_repeats = 100;
@@ -34,29 +36,21 @@ struct large_key {
   int32_t data_[Words];
 };
 
-template <typename T>
-constexpr __host__ __device__ void hash_result_aggregate(T& agg, T hash_val)
-{
-  agg += hash_val;
-}
+template <typename Hash>
+struct hash_key {
+  using type = typename Hash::argument_type;
+};
 
-template <>
-constexpr __host__ __device__ void hash_result_aggregate(cuda::std::array<uint64_t, 2>& agg,
-                                                         cuda::std::array<uint64_t, 2> hash_val)
-{
-  agg[0] += hash_val[0];
-  agg[1] += hash_val[1];
-}
+template <typename Key, cuda::hash_algorithm Algorithm>
+struct hash_key<cuda::hash<Key, Algorithm>> {
+  using type = Key;
+};
 
-template <>
-constexpr __host__ __device__ void hash_result_aggregate(cuda::std::array<uint32_t, 4>& agg,
-                                                         cuda::std::array<uint32_t, 4> hash_val)
-{
-  agg[0] += hash_val[0];
-  agg[1] += hash_val[1];
-  agg[2] += hash_val[2];
-  agg[3] += hash_val[3];
-}
+template <typename Hash>
+using hash_key_t = typename hash_key<Hash>::type;
+
+template <typename Hash>
+using hash_result_t = cuda::std::invoke_result_t<Hash, hash_key_t<Hash>>;
 
 template <int32_t BlockSize, typename Hasher, typename OutputIt>
 __global__ void hash_bench_kernel(Hasher hash,
@@ -67,12 +61,12 @@ __global__ void hash_bench_kernel(Hasher hash,
   cuco::detail::index_type const gid         = BlockSize * blockIdx.x + threadIdx.x;
   cuco::detail::index_type const loop_stride = gridDim.x * BlockSize;
   cuco::detail::index_type idx               = gid;
-  typename Hasher::result_type agg           = {};
+  hash_result_t<Hasher> agg                  = {};
 
   while (idx < n) {
-    typename Hasher::argument_type key(idx);
+    hash_key_t<Hasher> key(idx);
     for (int32_t i = 0; i < n_repeats; ++i) {  // execute hash func n times
-      hash_result_aggregate(agg, hash(key));
+      agg += hash(key);
     }
     idx += loop_stride;
   }
@@ -91,8 +85,7 @@ void hash_eval(nvbench::state& state, nvbench::type_list<Hash>)
   auto const num_keys           = state.get_int64("NumInputs");
   auto const grid_size          = (num_keys + block_size * 16 - 1) / block_size * 16;
 
-  thrust::device_vector<typename Hash::result_type> hash_values((materialize_result) ? num_keys
-                                                                                     : 1);
+  thrust::device_vector<hash_result_t<Hash>> hash_values((materialize_result) ? num_keys : 1);
 
   state.add_element_count(num_keys);
 
@@ -112,12 +105,12 @@ __global__ void string_hash_bench_kernel(
   cuco::detail::index_type const gid         = BlockSize * blockIdx.x + threadIdx.x;
   cuco::detail::index_type const loop_stride = gridDim.x * BlockSize;
   cuco::detail::index_type idx               = gid;
-  typename Hasher::result_type agg           = {};
+  hash_result_t<Hasher> agg                  = {};
 
   while (idx < n) {
     auto const key = thrust::raw_reference_cast(*(in + idx));
     for (int32_t i = 0; i < n_repeats; ++i) {  // execute hash func n times
-      hash_result_aggregate(agg, hash.compute_hash(key.data(), key.size()));
+      agg += hash(key);
     }
     idx += loop_stride;
   }
@@ -132,7 +125,7 @@ __global__ void string_hash_bench_kernel(
 template <typename Hash>
 void string_hash_eval(nvbench::state& state, nvbench::type_list<Hash>)
 {
-  static_assert(std::is_same_v<typename Hash::argument_type, cuda::std::byte>,
+  static_assert(cuda::std::is_same_v<hash_key_t<Hash>, cuda::std::byte>,
                 "Argument type must be cuda::std::byte");
 
   bool const materialize_result = false;
@@ -153,8 +146,7 @@ void string_hash_eval(nvbench::state& state, nvbench::type_list<Hash>)
   auto const& keys = sequences.first;
   // auto const& storage = sequences.second;
 
-  thrust::device_vector<typename Hash::result_type> hash_values((materialize_result) ? num_keys
-                                                                                     : 1);
+  thrust::device_vector<hash_result_t<Hash>> hash_values((materialize_result) ? num_keys : 1);
 
   state.add_element_count(num_keys);
   // state.add_global_memory_reads<cuda::std::byte>(storage.size() * n_repeats);
@@ -168,36 +160,44 @@ void string_hash_eval(nvbench::state& state, nvbench::type_list<Hash>)
     mem_stats_logger.peak_memory_usage(), "peak_memory_usage", "peak_memory_usage");
 }
 
-NVBENCH_BENCH_TYPES(
-  hash_eval,
-  NVBENCH_TYPE_AXES(nvbench::type_list<cuco::murmurhash3_32<nvbench::int32_t>,
-                                       cuco::murmurhash3_32<nvbench::int64_t>,
-                                       cuco::murmurhash3_32<large_key<32>>,  // 32*4bytes
-                                       cuco::xxhash_32<nvbench::int32_t>,
-                                       cuco::xxhash_32<nvbench::int64_t>,
-                                       cuco::xxhash_32<large_key<32>>,
-                                       cuco::xxhash_64<nvbench::int32_t>,
-                                       cuco::xxhash_64<nvbench::int64_t>,
-                                       cuco::xxhash_64<large_key<32>>,
-                                       cuco::murmurhash3_fmix_32<nvbench::int32_t>,
-                                       cuco::murmurhash3_fmix_64<nvbench::int64_t>,
-                                       cuco::murmurhash3_x86_128<nvbench::int32_t>,
-                                       cuco::murmurhash3_x86_128<nvbench::int64_t>,
-                                       cuco::murmurhash3_x86_128<large_key<32>>,
-                                       cuco::murmurhash3_x64_128<nvbench::int32_t>,
-                                       cuco::murmurhash3_x64_128<nvbench::int64_t>,
-                                       cuco::murmurhash3_x64_128<large_key<32>>>))
+using hash_types =
+  nvbench::type_list<cuda::hash<nvbench::int32_t, cuda::hash_algorithm::murmurhash3_32>,
+                     cuda::hash<nvbench::int64_t, cuda::hash_algorithm::murmurhash3_32>,
+                     cuda::hash<large_key<32>, cuda::hash_algorithm::murmurhash3_32>,  // 32*4bytes
+                     cuda::hash<nvbench::int32_t, cuda::hash_algorithm::xxhash_32>,
+                     cuda::hash<nvbench::int64_t, cuda::hash_algorithm::xxhash_32>,
+                     cuda::hash<large_key<32>, cuda::hash_algorithm::xxhash_32>,
+                     cuda::hash<nvbench::int32_t, cuda::hash_algorithm::xxhash_64>,
+                     cuda::hash<nvbench::int64_t, cuda::hash_algorithm::xxhash_64>,
+                     cuda::hash<large_key<32>, cuda::hash_algorithm::xxhash_64>
+#if defined(CUCO_HAS_INT128)
+                     ,
+                     cuda::hash<nvbench::int32_t, cuda::hash_algorithm::murmurhash3_x86_128>,
+                     cuda::hash<nvbench::int64_t, cuda::hash_algorithm::murmurhash3_x86_128>,
+                     cuda::hash<large_key<32>, cuda::hash_algorithm::murmurhash3_x86_128>,
+                     cuda::hash<nvbench::int32_t, cuda::hash_algorithm::murmurhash3_x64_128>,
+                     cuda::hash<nvbench::int64_t, cuda::hash_algorithm::murmurhash3_x64_128>,
+                     cuda::hash<large_key<32>, cuda::hash_algorithm::murmurhash3_x64_128>
+#endif
+                     >;
+
+NVBENCH_BENCH_TYPES(hash_eval, NVBENCH_TYPE_AXES(hash_types))
   .set_name("hash_function_eval")
   .set_type_axes_names({"Hash"})
   .add_int64_axis("NumInputs", {cuco::benchmark::defaults::N * 10});
 
-NVBENCH_BENCH_TYPES(
-  string_hash_eval,
-  NVBENCH_TYPE_AXES(nvbench::type_list<cuco::murmurhash3_32<cuda::std::byte>,
-                                       cuco::xxhash_32<cuda::std::byte>,
-                                       cuco::xxhash_64<cuda::std::byte>,
-                                       cuco::murmurhash3_x86_128<cuda::std::byte>,
-                                       cuco::murmurhash3_x64_128<cuda::std::byte>>))
+using string_hash_types =
+  nvbench::type_list<cuda::hash<cuda::std::byte, cuda::hash_algorithm::murmurhash3_32>,
+                     cuda::hash<cuda::std::byte, cuda::hash_algorithm::xxhash_32>,
+                     cuda::hash<cuda::std::byte, cuda::hash_algorithm::xxhash_64>
+#if defined(CUCO_HAS_INT128)
+                     ,
+                     cuda::hash<cuda::std::byte, cuda::hash_algorithm::murmurhash3_x86_128>,
+                     cuda::hash<cuda::std::byte, cuda::hash_algorithm::murmurhash3_x64_128>
+#endif
+                     >;
+
+NVBENCH_BENCH_TYPES(string_hash_eval, NVBENCH_TYPE_AXES(string_hash_types))
   .set_name("string_hash_function_eval")
   .set_type_axes_names({"Hash"})
   .add_int64_axis("NumInputs", {cuco::benchmark::defaults::N / 4})

benchmarks/hyperloglog/hyperloglog_bench.cu

@@ -72,7 +72,7 @@ void hyperloglog_e2e(nvbench::state& state, nvbench::type_list<T, Dist>)
 {
   using estimator_type = cuco::hyperloglog<T,
                                            cuda::thread_scope_device,
-                                           cuco::xxhash_64<T>,
+                                           cuda::hash<T, cuda::hash_algorithm::xxhash_64>,
                                            tracking_allocator<cuda::std::byte>>;
 
   auto const num_items      = state.get_int64("NumInputs");
@@ -118,7 +118,7 @@ void hyperloglog_add(nvbench::state& state, nvbench::type_list<T, Dist>)
 {
   using estimator_type = cuco::hyperloglog<T,
                                            cuda::thread_scope_device,
-                                           cuco::xxhash_64<T>,
+                                           cuda::hash<T, cuda::hash_algorithm::xxhash_64>,
                                            tracking_allocator<cuda::std::byte>>;
 
   auto const num_items      = state.get_int64("NumInputs");

include/cuco/bloom_filter_policy.cuh

@@ -9,6 +9,8 @@
 #include <cuco/hash_functions.cuh>
 #include <cuco/utility/cuda.cuh>
 
+#include <cuda/functional>
+
 #include <cstdint>
 
 namespace cuco {
@@ -23,7 +25,8 @@ namespace cuco {
  * multiply-shift) and lower 32 bits (pattern generation).
  *
  * @tparam Key Key type to hash.
- * @tparam Hash 64-bit hash functor type. Defaults to `cuco::xxhash_64<Key>`.
+ * @tparam Hash 64-bit hash functor type. Defaults to `cuda::hash<Key,
+ * cuda::hash_algorithm::xxhash_64>`.
  * @tparam WordBytes Size in bytes of the underlying word type. Must be `4` or `8`. Defaults to `4`.
  * @tparam WordsPerBlock Words per filter block. Defaults to the number of words that fit in one
  * 32-byte sector.
@@ -55,7 +58,7 @@ namespace cuco {
  * `contains` is about as fast as without it. No effect for blocks smaller than 32 bytes.
  */
 template <class Key,
-          class Hash                        = cuco::xxhash_64<Key>,
+          class Hash                        = cuda::hash<Key, cuda::hash_algorithm::xxhash_64>,
           std::uint32_t WordBytes           = 4,
           std::uint32_t WordsPerBlock       = utility::sector_size_bytes / WordBytes,
           std::uint32_t PatternBits         = WordsPerBlock,

include/cuco/detail/__config

@@ -19,8 +19,8 @@
 #error "Support for extended device lambdas is required (nvcc flag --expt-extended-lambda)"
 #endif
 
-#if !defined(CCCL_VERSION) || (CCCL_VERSION < 3002000)
-#error "CCCL version 3.2.0 or later is required"
+#if !defined(CCCL_VERSION) || (CCCL_VERSION < 3006000)
+#error "CCCL version 3.6.0 or later is required"
 #endif
 
 // WAR for cuCollections/804: work around the nvcc __cuda_is_local / isspacep.local