Replace window with bucket instead (#656)

This PR completely deprecates the window concept in cuco hash tables.

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/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/extent/extent.inl

@@ -111,24 +111,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)
 {
@@ -167,15 +155,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)
 {
@@ -185,15 +164,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

include/cuco/static_set_ref.cuh

@@ -77,7 +77,6 @@ class static_set_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

tests/static_map/insert_or_apply_test.cu

@@ -101,7 +101,7 @@ void test_insert_or_apply_shmem(Map& map, size_type num_keys, size_type num_uniq
                                            cuco::storage<1>>;
 
   using shared_map_ref_type    = typename shared_map_type::ref_type<>;
-  auto constexpr window_extent = cuco::make_window_extent<shared_map_ref_type>(extent_type{});
+  auto constexpr bucket_extent = cuco::make_bucket_extent<shared_map_ref_type>(extent_type{});
 
   // Insert pairs
   auto pairs_begin = thrust::make_transform_iterator(
@@ -122,7 +122,7 @@ void test_insert_or_apply_shmem(Map& map, size_type num_keys, size_type num_uniq
                                                              init,
                                                              cuco::reduce::plus{},
                                                              map.ref(cuco::op::insert_or_apply),
-                                                             window_extent);
+                                                             bucket_extent);
 
   REQUIRE(map.size() == num_unique_keys);