Merge branch 'dev' into dev

CMakeLists.txt

@@ -16,7 +16,7 @@
 cmake_minimum_required(VERSION 3.23.1 FATAL_ERROR)
 
 if(NOT EXISTS ${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake)
-    file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-24.10/RAPIDS.cmake
+    file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-25.02/RAPIDS.cmake
          ${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake)
 endif()
 include(${CMAKE_CURRENT_BINARY_DIR}/CUCO_RAPIDS.cmake)

examples/static_set/device_subsets_example.cu

@@ -44,7 +44,7 @@
  */
 
 auto constexpr cg_size     = 8;   ///< A CUDA Cooperative Group of 8 threads to handle each subset
-auto constexpr window_size = 1;   ///< Number of concurrent slots handled by each thread
+auto constexpr bucket_size = 1;   ///< Number of concurrent slots handled by each thread
 auto constexpr N           = 10;  ///< Number of elements to insert and query
 
 using key_type = int;  ///< Key type
@@ -54,7 +54,7 @@ using probing_scheme_type =
                                                                 ///< and probing scheme (linear
                                                                 ///< probing v.s. double hashing)
 /// Type of bulk allocation storage
-using storage_type = cuco::aow_storage<key_type, window_size>;
+using storage_type = cuco::aow_storage<key_type, bucket_size>;
 /// Lightweight non-owning storage ref type
 using storage_ref_type = typename storage_type::ref_type;
 using ref_type         = cuco::static_set_ref<key_type,
@@ -143,23 +143,23 @@ int main()
 
   for (size_t i = 0; i < num; ++i) {
     valid_sizes.emplace_back(
-      static_cast<std::size_t>(cuco::make_window_extent<ref_type>(subset_sizes[i])));
+      static_cast<std::size_t>(cuco::make_bucket_extent<ref_type>(subset_sizes[i])));
   }
 
   std::vector<std::size_t> offsets(num + 1, 0);
 
-  // prefix sum to compute offsets and total number of windows
+  // prefix sum to compute offsets and total number of buckets
   std::size_t current_sum = 0;
   for (std::size_t i = 0; i < valid_sizes.size(); ++i) {
     current_sum += valid_sizes[i];
     offsets[i + 1] = current_sum;
   }
 
-  // total number of windows is located at the back of the offsets array
-  auto const total_num_windows = offsets.back();
+  // total number of buckets is located at the back of the offsets array
+  auto const total_num_buckets = offsets.back();
 
   // Create a single bulk storage used by all subsets
-  auto set_storage = storage_type{total_num_windows};
+  auto set_storage = storage_type{total_num_buckets};
   // Initializes the storage with the given sentinel
   set_storage.initialize(empty_key_sentinel);
 

examples/static_set/shared_memory_example.cu

@@ -24,16 +24,16 @@
  */
 
 template <class SetRef>
-__global__ void shmem_set_kernel(typename SetRef::extent_type window_extent,
+__global__ void shmem_set_kernel(typename SetRef::extent_type bucket_extent,
                                  cuco::empty_key<typename SetRef::key_type> empty_key_sentinel)
 {
   // We first allocate the shared memory storage for the `set`.
-  // The storage is comprised of contiguous windows of slots,
+  // The storage is comprised of contiguous buckets of slots,
   // which allow for vectorized loads.
-  __shared__ typename SetRef::window_type windows[window_extent.value()];
+  __shared__ typename SetRef::bucket_type buckets[bucket_extent.value()];
 
   // Next, we construct the actual storage object from the raw array.
-  auto storage = SetRef::storage_ref_type(window_extent, windows);
+  auto storage = SetRef::storage_ref_type(bucket_extent, buckets);
   // Now we can instantiate the set from the storage.
   auto set = SetRef(empty_key_sentinel, {}, {}, {}, storage);
 
@@ -75,7 +75,7 @@ int main(void)
   // Inserting or retrieving this value is UB.
   cuco::empty_key<Key> constexpr empty_key_sentinel{-1};
   // Width of vectorized loads during probing.
-  auto constexpr window_size = 1;
+  auto constexpr bucket_size = 1;
   // Cooperative group size
   auto constexpr cg_size = 1;
 
@@ -93,17 +93,17 @@ int main(void)
                                     thrust::equal_to<Key>,
                                     probing_scheme_type,
                                     cuco::cuda_allocator<Key>,
-                                    cuco::storage<window_size>>;
+                                    cuco::storage<bucket_size>>;
   // Next, we can derive the non-owning reference type from the set type.
   // This is the type we use in the kernel to wrap a raw shared memory array as a `static_set`.
   using set_ref_type = typename set_type::ref_type<>;
 
   // Cuco imposes a number of non-trivial contraints on the capacity value.
   // This function will take the requested capacity (1000) and return the next larger
   // valid extent.
-  auto constexpr window_extent = cuco::make_window_extent<set_ref_type>(extent_type{});
+  auto constexpr bucket_extent = cuco::make_bucket_extent<set_ref_type>(extent_type{});
 
   // Launch the kernel with a single thread block.
-  shmem_set_kernel<set_ref_type><<<1, 128>>>(window_extent, empty_key_sentinel);
+  shmem_set_kernel<set_ref_type><<<1, 128>>>(bucket_extent, empty_key_sentinel);
   cudaDeviceSynchronize();
 }

include/cuco/bloom_filter_policies.cuh

@@ -33,7 +33,7 @@ namespace cuco {
  * By default, cuco::xxhash_64 hasher will be used.
  *
  */
-template <class Key, class XXHash64 = cuco::xxhash_64<Key>>
+template <typename Key, template <typename> class XXHash64 = cuco::xxhash_64>
 using arrow_filter_policy = detail::arrow_filter_policy<Key, XXHash64>;
 
 /**

include/cuco/bucket_storage.cuh

@@ -33,10 +33,6 @@ namespace cuco {
 template <typename T, int32_t BucketSize>
 using bucket = detail::bucket<T, BucketSize>;
 
-/// Alias for bucket
-template <typename T, int32_t BucketSize>
-using window = bucket<T, BucketSize>;
-
 /**
  * @brief Non-owning array of buckets storage reference type.
  *
@@ -56,7 +52,6 @@ class bucket_storage_ref : public detail::bucket_storage_base<T, BucketSize, Ext
   using size_type   = typename base_type::size_type;    ///< Storage size type
   using value_type  = typename base_type::value_type;   ///< Slot type
   using bucket_type = typename base_type::bucket_type;  ///< Slot bucket type
-  using window_type = bucket_type;                      ///< Slot bucket type
 
   using base_type::capacity;
   using base_type::num_buckets;
@@ -147,7 +142,6 @@ class bucket_storage : public detail::bucket_storage_base<T, BucketSize, Extent>
   using size_type   = typename base_type::size_type;    ///< Storage size type
   using value_type  = typename base_type::value_type;   ///< Slot type
   using bucket_type = typename base_type::bucket_type;  ///< Slot bucket type
-  using window_type = bucket_type;                      ///< Slot bucket type
 
   using base_type::capacity;
   using base_type::num_buckets;

include/cuco/detail/bloom_filter/arrow_filter_policy.cuh

@@ -40,8 +40,7 @@ namespace cuco::detail {
  * void bulk_insert_and_eval_arrow_policy_bloom_filter(device_vector<KeyType> const& positive_keys,
  *                                                 device_vector<KeyType> const& negative_keys)
  * {
- *     using xxhash_64 = cuco::xxhash_64<KeyType>;
- *     using policy_type = cuco::arrow_filter_policy<KeyType, xxhash_64>;
+ *     using policy_type = cuco::arrow_filter_policy<KeyType, cuco::xxhash_64>;
  *
  *     // Warn or throw if the number of filter blocks is greater than maximum used by Arrow policy.
  *     static_assert(NUM_FILTER_BLOCKS <= policy_type::max_filter_blocks, "NUM_FILTER_BLOCKS must be
@@ -81,14 +80,13 @@ namespace cuco::detail {
  * @tparam Key The type of the values to generate a fingerprint for.
  * @tparam XXHash64 64-bit XXHash hasher implementation for fingerprint generation.
  */
-template <class Key, class XXHash64>
+template <class Key, template <typename> class XXHash64>
 class arrow_filter_policy {
  public:
-  using hasher             = XXHash64;       ///< 64-bit XXHash hasher for Arrow bloom filter policy
-  using word_type          = std::uint32_t;  ///< uint32_t for Arrow bloom filter policy
-  using hash_argument_type = typename hasher::argument_type;  ///< Hash function input type
-  using hash_result_type   = decltype(std::declval<hasher>()(
-    std::declval<hash_argument_type>()));  ///< hash function output type
+  using hasher          = XXHash64<Key>;  ///< 64-bit XXHash hasher for Arrow bloom filter policy
+  using word_type       = std::uint32_t;  ///< uint32_t for Arrow bloom filter policy
+  using key_type        = Key;            ///< Hash function input type
+  using hash_value_type = std::uint64_t;  ///< hash function output type
 
   static constexpr uint32_t bits_set_per_block = 8;  ///< hardcoded bits set per Arrow filter block
   static constexpr uint32_t words_per_block    = 8;  ///< hardcoded words per Arrow filter block
@@ -135,10 +133,7 @@ class arrow_filter_policy {
    *
    * @return The hash value of the key
    */
-  __device__ constexpr hash_result_type hash(hash_argument_type const& key) const
-  {
-    return hash_(key);
-  }
+  __device__ constexpr hash_value_type hash(key_type const& key) const { return hash_(key); }
 
   /**
    * @brief Determines the filter block a key is added into.
@@ -155,7 +150,7 @@ class arrow_filter_policy {
    * @return The block index for the given key's hash value
    */
   template <class Extent>
-  __device__ constexpr auto block_index(hash_result_type hash, Extent num_blocks) const
+  __device__ constexpr auto block_index(hash_value_type hash, Extent num_blocks) const
   {
     constexpr auto hash_bits = cuda::std::numeric_limits<word_type>::digits;
     // TODO: assert if num_blocks > max_filter_blocks
@@ -173,7 +168,7 @@ class arrow_filter_policy {
    *
    * @return The bit pattern for the word/segment in the filter block
    */
-  __device__ constexpr word_type word_pattern(hash_result_type hash, std::uint32_t word_index) const
+  __device__ constexpr word_type word_pattern(hash_value_type hash, std::uint32_t word_index) const
   {
     // SALT array to calculate bit indexes for the current word
     auto constexpr salt = SALT();

include/cuco/detail/extent/extent.inl

@@ -112,24 +112,12 @@ template <int32_t CGSize, int32_t BucketSize, typename SizeType, std::size_t N>
   }
 }
 
-template <int32_t CGSize, int32_t BucketSize, typename SizeType, std::size_t N>
-[[nodiscard]] auto constexpr make_window_extent(extent<SizeType, N> ext)
-{
-  return make_bucket_extent<CGSize, BucketSize, SizeType, N>(ext);
-}
-
 template <int32_t CGSize, int32_t BucketSize, typename SizeType>
 [[nodiscard]] auto constexpr make_bucket_extent(SizeType size)
 {
   return make_bucket_extent<CGSize, BucketSize, SizeType, dynamic_extent>(extent<SizeType>{size});
 }
 
-template <int32_t CGSize, int32_t BucketSize, typename SizeType>
-[[nodiscard]] auto constexpr make_window_extent(SizeType size)
-{
-  return make_bucket_extent<CGSize, BucketSize, SizeType, dynamic_extent>(extent<SizeType>{size});
-}
-
 template <typename ProbingScheme, typename Storage, typename SizeType, std::size_t N>
 [[nodiscard]] auto constexpr make_bucket_extent(extent<SizeType, N> ext)
 {
@@ -168,15 +156,6 @@ template <typename Container, typename SizeType, std::size_t N>
                             N>(ext);
 }
 
-template <typename Container, typename SizeType, std::size_t N>
-[[nodiscard]] auto constexpr make_window_extent(extent<SizeType, N> ext)
-{
-  return make_bucket_extent<typename Container::probing_scheme_type,
-                            typename Container::storage_ref_type,
-                            SizeType,
-                            N>(ext);
-}
-
 template <typename Container, typename SizeType>
 [[nodiscard]] auto constexpr make_bucket_extent(SizeType size)
 {
@@ -186,15 +165,6 @@ template <typename Container, typename SizeType>
                             dynamic_extent>(extent<SizeType>{size});
 }
 
-template <typename Container, typename SizeType>
-[[nodiscard]] auto constexpr make_window_extent(SizeType size)
-{
-  return make_bucket_extent<typename Container::probing_scheme_type,
-                            typename Container::storage_ref_type,
-                            SizeType,
-                            dynamic_extent>(extent<SizeType, dynamic_extent>{size});
-}
-
 namespace detail {
 
 template <typename...>

include/cuco/extent.cuh

@@ -85,10 +85,6 @@ struct extent<SizeType, dynamic_extent> {
 template <typename SizeType, std::size_t N = dynamic_extent>
 struct bucket_extent;
 
-/// Alias for bucket_extent
-template <typename SizeType, std::size_t N = dynamic_extent>
-using window_extent = bucket_extent<SizeType, N>;
-
 /**
  * @brief Computes valid bucket extent based on given parameters.
  *
@@ -112,25 +108,6 @@ using window_extent = bucket_extent<SizeType, N>;
 template <int32_t CGSize, int32_t BucketSize, typename SizeType, std::size_t N>
 [[nodiscard]] auto constexpr make_bucket_extent(extent<SizeType, N> ext);
 
-/**
- * @brief Computes valid bucket extent based on given parameters.
- *
- * @deprecated Use the equivalent `make_bucket_extent` instead.
- *
- * @tparam CGSize Number of elements handled per CG
- * @tparam BucketSize Number of elements handled per Bucket
- * @tparam SizeType Size type
- * @tparam N Extent
- *
- * @param ext The input extent
- *
- * @throw If the input extent is invalid
- *
- * @return Resulting valid extent
- */
-template <int32_t CGSize, int32_t BucketSize, typename SizeType, std::size_t N>
-[[nodiscard]] auto constexpr make_window_extent(extent<SizeType, N> ext);
-
 /**
  * @brief Computes valid bucket extent/capacity based on given parameters.
  *
@@ -153,24 +130,6 @@ template <int32_t CGSize, int32_t BucketSize, typename SizeType, std::size_t N>
 template <int32_t CGSize, int32_t BucketSize, typename SizeType>
 [[nodiscard]] auto constexpr make_bucket_extent(SizeType size);
 
-/**
- * @brief Computes valid bucket extent/capacity based on given parameters.
- *
- * @deprecated Use the equivalent `make_bucket_extent` instead.
- *
- * @tparam CGSize Number of elements handled per CG
- * @tparam BucketSize Number of elements handled per Bucket
- * @tparam SizeType Size type
- *
- * @param size The input size
- *
- * @throw If the input size is invalid
- *
- * @return Resulting valid extent
- */
-template <int32_t CGSize, int32_t BucketSize, typename SizeType>
-[[nodiscard]] auto constexpr make_window_extent(SizeType size);
-
 template <typename ProbingScheme, typename Storage, typename SizeType, std::size_t N>
 [[nodiscard]] auto constexpr make_bucket_extent(cuco::extent<SizeType, N> ext);
 
@@ -199,24 +158,6 @@ template <typename ProbingScheme, typename Storage, typename SizeType>
 template <typename Container, typename SizeType, std::size_t N>
 [[nodiscard]] auto constexpr make_bucket_extent(extent<SizeType, N> ext);
 
-/**
- * @brief Computes a valid bucket extent/capacity for a given container type.
- *
- * @deprecated Use the equivalent `make_bucket_extent` instead.
- *
- * @tparam Container Container type to compute the extent for
- * @tparam SizeType Size type
- * @tparam N Extent
- *
- * @param ext The input extent
- *
- * @throw If the input extent is invalid
- *
- * @return Resulting valid `bucket extent`
- */
-template <typename Container, typename SizeType, std::size_t N>
-[[nodiscard]] auto constexpr make_window_extent(extent<SizeType, N> ext);
-
 /**
  * @brief Computes a valid capacity for a given container type.
  *
@@ -238,23 +179,6 @@ template <typename Container, typename SizeType, std::size_t N>
 template <typename Container, typename SizeType>
 [[nodiscard]] auto constexpr make_bucket_extent(SizeType size);
 
-/**
- * @brief Computes a valid capacity for a given container type.
- *
- * @deprecated Use the equivalent `make_bucket_extent` instead.
- *
- * @tparam Container Container type to compute the extent for
- * @tparam SizeType Size type
- *
- * @param size The input size
- *
- * @throw If the input size is invalid
- *
- * @return Resulting valid extent
- */
-template <typename Container, typename SizeType>
-[[nodiscard]] auto constexpr make_window_extent(SizeType size);
-
 }  // namespace cuco
 
 #include <cuco/detail/extent/extent.inl>

include/cuco/static_map_ref.cuh

@@ -89,7 +89,6 @@ class static_map_ref
   using hasher              = typename probing_scheme_type::hasher;    ///< Hash function type
   using storage_ref_type    = StorageRef;                              ///< Type of storage ref
   using bucket_type         = typename storage_ref_type::bucket_type;  ///< Bucket type
-  using window_type         = bucket_type;                             ///< Bucket type
   using value_type          = typename storage_ref_type::value_type;   ///< Storage element type
   using extent_type         = typename storage_ref_type::extent_type;  ///< Extent type
   using size_type           = typename storage_ref_type::size_type;    ///< Probing scheme size type

include/cuco/static_multimap_ref.cuh

@@ -86,7 +86,6 @@ class static_multimap_ref
   using hasher              = typename probing_scheme_type::hasher;    ///< Hash function type
   using storage_ref_type    = StorageRef;                              ///< Type of storage ref
   using bucket_type         = typename storage_ref_type::bucket_type;  ///< Bucket type
-  using window_type         = bucket_type;                             ///< Bucket type
   using value_type          = typename storage_ref_type::value_type;   ///< Storage element type
   using extent_type         = typename storage_ref_type::extent_type;  ///< Extent type
   using size_type           = typename storage_ref_type::size_type;    ///< Probing scheme size type

include/cuco/static_multiset_ref.cuh

@@ -77,7 +77,6 @@ class static_multiset_ref
   using hasher              = typename probing_scheme_type::hasher;    ///< Hash function type
   using storage_ref_type    = StorageRef;                              ///< Type of storage ref
   using bucket_type         = typename storage_ref_type::bucket_type;  ///< Bucket type
-  using window_type         = bucket_type;                             ///< Bucket type
   using value_type          = typename storage_ref_type::value_type;   ///< Storage element type
   using extent_type         = typename storage_ref_type::extent_type;  ///< Extent type
   using size_type           = typename storage_ref_type::size_type;    ///< Probing scheme size type