Refactor window storage (#627)

Closes #621

Based on the offline discussions, this PR replaces the `window` logic
with `bucket` and adds new overloads of `make_bucket_extent` so OA no
longer relies on the `Container` type to determine the bucket extent.

---------

Co-authored-by: Daniel Jünger <[email protected]>

Author: PointKernel

README.md

@@ -15,6 +15,8 @@ Similar to how [Thrust](https://github.com/thrust/thrust) and [CUB](https://gith
 
 ### Major Updates
 
+__11/01/2024__ Refined the term `window` as `bucket`
+
 __01/08/2024__ Deprecated the `experimental` namespace
 
 __01/02/2024__ Moved the legacy `static_map` to `cuco::legacy` namespace
@@ -254,4 +256,4 @@ We plan to add many GPU-accelerated, concurrent data structures to `cuCollection
 `cuco::bloom_filter` implements a Blocked Bloom Filter for approximate set membership queries.
 
 #### Examples:
-- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw))
+- [Host-bulk APIs (Default fingerprinting policy)](https://github.com/NVIDIA/cuCollections/blob/dev/examples/bloom_filter/host_bulk_example.cu) (see [live example in godbolt](https://godbolt.org/clientstate/eJydVmtvGjkU_StXsx8WmuEVbVUJQiSapLtoK5IF2qpaVsjj8TBWBnvqBwRF-e977ZmBgZBqtVRqwL6Pc889vvZzoJnWXAod9P9-Dngc9HthkBGxsmTFgn5AbUyCMNDSKup-d94tBLyDG5nvFF-lBhq0CZfdy99CmHwd345HcHM_fbifjubj-0nb2Xr7z5wyoVkMVsRMgUkZjHJC8U-5E8JXphwQuGx3oeEMFkG5twiaAx9lJy2syQ6ENGA1wzBcQ8IzBuyJstwAF0DlOs84EZTBlpvUpyrjeDjwvQwiI0PQnqBHjr-SuiUQs4fuPqkxeb_T2W63beJht6VadbLCWHc-j2_uJrO7FkLfu30RGTILiv2wXGHh0Q5IjsgoiRBvRrYgFZCVYrhnpEO-VdxwsQpBy8RsiWI-Tsy1UTyy5oi8CifWXzdA-ohA4kYzGM8WAXwczcaz0Mf5Np7_cf9lDt9G0-loMh_fzeB-is2a3I5dq_DXJxhNvsOf48ltCAypw1TsKVeuCoTKHa0sLjicMXYEI5EFLJ0zyhNOoVIQrOSGKYFlQc7UmhdaQ5Cxj5PxNTfE-LVXxflUnYVYiF-4oJmNGVxRS2UnyqRcL7Hvhqk2ten1sY1JldWmQ6UVpu02X23FbIMplhtGjVTnTdgTo9YBW-YSm7Y7b6Wxuwyl1j7FwCU2hZG1X-bCoOK4aGwkj5sL8YyFgVukWLdxHIOw6-Uj22kntiH0ur92u90B7D-dTucKfmeCKWJYuQ3O_nwkkxduw0Pcd9Btvx-UkcbIrjKe64QrbSAlWeLjuWCy3PD0vpFAvErQKjMPXntoG5Xd0uhx6SvzOPCr22xVmygR8M0tszvynKmDBmaXF0I7tgCcBE5eaLL0JkOXfbB3neVk60553auYDDFLiM0MFA12mjwF5Kt3kuv365q7qnJdl_GeaxW-lKm1ift91KCBqys8kR9t9ojAPO-e47fxJFgOU7lyJK6Knru5WMDsY6yFgNoHw_tcTMRZmbtQZ79_pPMaatewRtU5P1v3LpWiG26rHbEVyrYZeo82ZnDfe80yDbFOK_nSWyHvdZdB3QAdvVj2thd1sRRWYh-mcDjaqgIcYBwafF7M5Tz3Ds6wlDOJ40aFIiwzNX_KWiRldu0scRBidxoF8BASkml2zN1ZR1F3PDoptRhlKX9Zpna107efqziFN1xane1KDeGs35eGjnN3EepU2iyGIh34q80oy1q51HixbBj44YHEzB-mw16dFTyq7i7Ur6gJD4VXjf0vcH2ZJiUG8Bbzt7W7RpjwqE6Gizui3N3W9QOhvWOuZEQyvN7wQomJIahzZamxGCushSmjsKeUR9yga8HrSd2fHqbXXTxjOZbl5oUskGArIkSOrHhfXtVWq8oNhmPTlOiUafcEif3MwnrP0yn2dIqSTvEGnUkmkS7HtgM7LI64X2xU8vIXWuNYbGcaVF8qT6yTQbMJnTJgIb9Cu0Xe5H_kPa2kvvSTvKJ5fkA6UeKLBb9XJPhl17dyPTmsnw48xVAVArr48wVfq-4NiK9CdXjUBmJDae_yve3htsxN8eINWhhoSC8ueh-gRRRNh3q9_NCFVgvvLYP_GczB4lZG1pF_Bmc8qsWklGa4uCkerriA9YrH4CWs9vHmONpH7oKXf_y_fwHeCexw))

include/cuco/aow_storage.cuh include/cuco/bucket_storage.cuhinclude/cuco/aow_storage.cuh renamed to include/cuco/bucket_storage.cuh

@@ -16,7 +16,7 @@
 
 #pragma once
 
-#include <cuco/detail/storage/aow_storage_base.cuh>
+#include <cuco/detail/storage/bucket_storage_base.cuh>
 #include <cuco/extent.cuh>
 #include <cuco/utility/allocator.hpp>
 
@@ -29,200 +29,213 @@
 #include <memory>
 
 namespace cuco {
+/// Bucket type alias
+template <typename T, int32_t BucketSize>
+using bucket = detail::bucket<T, BucketSize>;
 
-/// Window type alias
-template <typename T, int32_t WindowSize>
-using window = detail::window<T, WindowSize>;
-
-/// forward declaration
-template <typename T, int32_t WindowSize, typename Extent>
-class aow_storage_ref;
+/// Alias for bucket
+template <typename T, int32_t BucketSize>
+using window = bucket<T, BucketSize>;
 
 /**
- * @brief Array of Window open addressing storage class.
+ * @brief Non-owning array of buckets storage reference type.
  *
- * @tparam T Slot type
- * @tparam WindowSize Number of slots in each window
- * @tparam Extent Type of extent denoting number of windows
- * @tparam Allocator Type of allocator used for device storage (de)allocation
+ * @tparam T Storage element type
+ * @tparam BucketSize Number of slots in each bucket
+ * @tparam Extent Type of extent denoting storage capacity
  */
-template <typename T,
-          int32_t WindowSize,
-          typename Extent    = cuco::extent<std::size_t>,
-          typename Allocator = cuco::cuda_allocator<cuco::window<T, WindowSize>>>
-class aow_storage : public detail::aow_storage_base<T, WindowSize, Extent> {
+template <typename T, int32_t BucketSize, typename Extent = cuco::extent<std::size_t>>
+class bucket_storage_ref : public detail::bucket_storage_base<T, BucketSize, Extent> {
  public:
-  using base_type = detail::aow_storage_base<T, WindowSize, Extent>;  ///< AoW base class type
+  /// Array of buckets base class type
+  using base_type = detail::bucket_storage_base<T, BucketSize, Extent>;
 
-  using base_type::window_size;  ///< Number of elements processed per window
+  using base_type::bucket_size;  ///< Number of elements processed per bucket
 
   using extent_type = typename base_type::extent_type;  ///< Storage extent type
   using size_type   = typename base_type::size_type;    ///< Storage size type
   using value_type  = typename base_type::value_type;   ///< Slot type
-  using window_type = typename base_type::window_type;  ///< Slot window type
+  using bucket_type = typename base_type::bucket_type;  ///< Slot bucket type
 
   using base_type::capacity;
-  using base_type::num_windows;
-
-  /// Type of the allocator to (de)allocate windows
-  using allocator_type =
-    typename std::allocator_traits<Allocator>::template rebind_alloc<window_type>;
-  using window_deleter_type =
-    detail::custom_deleter<size_type, allocator_type>;  ///< Type of window deleter
-  using ref_type = aow_storage_ref<value_type, window_size, extent_type>;  ///< Storage ref type
+  using base_type::num_buckets;
 
   /**
-   * @brief Constructor of AoW storage.
-   *
-   * @note The input `size` should be exclusively determined by the return value of
-   * `make_window_extent` since it depends on the requested low-bound value, the probing scheme, and
-   * the storage.
+   * @brief Constructor of AoS storage ref.
    *
-   * @param size Number of windows to (de)allocate
-   * @param allocator Allocator used for (de)allocating device storage
+   * @param size Number of buckets
+   * @param buckets Pointer to the buckets array
    */
-  explicit constexpr aow_storage(Extent size, Allocator const& allocator = {});
+  __host__ __device__ explicit constexpr bucket_storage_ref(Extent size,
+                                                            bucket_type* buckets) noexcept;
 
-  aow_storage(aow_storage&&) = default;  ///< Move constructor
   /**
-   * @brief Replaces the contents of the storage with another storage.
+   * @brief Custom un-incrementable input iterator for the convenience of `find` operations.
    *
-   * @return Reference of the current storage object
+   * @note This iterator is for read only and NOT incrementable.
    */
-  aow_storage& operator=(aow_storage&&) = default;
-  ~aow_storage()                        = default;  ///< Destructor
-
-  aow_storage(aow_storage const&)            = delete;
-  aow_storage& operator=(aow_storage const&) = delete;
+  struct iterator;
+  using const_iterator = iterator const;  ///< Const forward iterator type
 
   /**
-   * @brief Gets windows array.
+   * @brief Returns an iterator to one past the last slot.
+   *
+   * This is provided for convenience for those familiar with checking
+   * an iterator returned from `find()` against the `end()` iterator.
    *
-   * @return Pointer to the first window
+   * @return An iterator to one past the last slot
    */
-  [[nodiscard]] constexpr window_type* data() const noexcept;
+  [[nodiscard]] __device__ constexpr iterator end() noexcept;
 
   /**
-   * @brief Gets the storage allocator.
+   * @brief Returns a const_iterator to one past the last slot.
    *
-   * @return The storage allocator
+   * This is provided for convenience for those familiar with checking
+   * an iterator returned from `find()` against the `end()` iterator.
+   *
+   * @return A const_iterator to one past the last slot
    */
-  [[nodiscard]] constexpr allocator_type allocator() const noexcept;
+  [[nodiscard]] __device__ constexpr const_iterator end() const noexcept;
 
   /**
-   * @brief Gets window storage reference.
+   * @brief Gets buckets array.
    *
-   * @return Reference of window storage
+   * @return Pointer to the first bucket
    */
-  [[nodiscard]] constexpr ref_type ref() const noexcept;
+  [[nodiscard]] __device__ constexpr bucket_type* data() noexcept;
 
   /**
-   * @brief Initializes each slot in the AoW storage to contain `key`.
+   * @brief Gets bucket array.
    *
-   * @param key Key to which all keys in `slots` are initialized
-   * @param stream Stream used for executing the kernel
+   * @return Pointer to the first bucket
    */
-  void initialize(value_type key, cuda::stream_ref stream = {});
+  [[nodiscard]] __device__ constexpr bucket_type* data() const noexcept;
 
   /**
-   * @brief Asynchronously initializes each slot in the AoW storage to contain `key`.
+   * @brief Returns an array of slots (or a bucket) for a given index.
    *
-   * @param key Key to which all keys in `slots` are initialized
-   * @param stream Stream used for executing the kernel
+   * @param index Index of the bucket
+   * @return An array of slots
    */
-  void initialize_async(value_type key, cuda::stream_ref stream = {}) noexcept;
+  [[nodiscard]] __device__ constexpr bucket_type operator[](size_type index) const noexcept;
 
  private:
-  allocator_type allocator_;            ///< Allocator used to (de)allocate windows
-  window_deleter_type window_deleter_;  ///< Custom windows deleter
-  std::unique_ptr<window_type, window_deleter_type> windows_;  ///< Pointer to AoW storage
+  bucket_type* buckets_;  ///< Pointer to the buckets array
 };
 
 /**
- * @brief Non-owning AoW storage reference type.
+ * @brief Array of buckets open addressing storage class.
  *
- * @tparam T Storage element type
- * @tparam WindowSize Number of slots in each window
- * @tparam Extent Type of extent denoting storage capacity
+ * @tparam T Slot type
+ * @tparam BucketSize Number of slots in each bucket
+ * @tparam Extent Type of extent denoting number of buckets
+ * @tparam Allocator Type of allocator used for device storage (de)allocation
  */
-template <typename T, int32_t WindowSize, typename Extent = cuco::extent<std::size_t>>
-class aow_storage_ref : public detail::aow_storage_base<T, WindowSize, Extent> {
+template <typename T,
+          int32_t BucketSize,
+          typename Extent    = cuco::extent<std::size_t>,
+          typename Allocator = cuco::cuda_allocator<cuco::bucket<T, BucketSize>>>
+class bucket_storage : public detail::bucket_storage_base<T, BucketSize, Extent> {
  public:
-  using base_type = detail::aow_storage_base<T, WindowSize, Extent>;  ///< AoW base class type
+  /// Array of buckets base class type
+  using base_type = detail::bucket_storage_base<T, BucketSize, Extent>;
 
-  using base_type::window_size;  ///< Number of elements processed per window
+  using base_type::bucket_size;  ///< Number of elements processed per bucket
 
   using extent_type = typename base_type::extent_type;  ///< Storage extent type
   using size_type   = typename base_type::size_type;    ///< Storage size type
   using value_type  = typename base_type::value_type;   ///< Slot type
-  using window_type = typename base_type::window_type;  ///< Slot window type
+  using bucket_type = typename base_type::bucket_type;  ///< Slot bucket type
 
   using base_type::capacity;
-  using base_type::num_windows;
+  using base_type::num_buckets;
+
+  /// Type of the allocator to (de)allocate buckets
+  using allocator_type =
+    typename std::allocator_traits<Allocator>::template rebind_alloc<bucket_type>;
+  using bucket_deleter_type =
+    detail::custom_deleter<size_type, allocator_type>;  ///< Type of bucket deleter
+  using ref_type = bucket_storage_ref<value_type, bucket_size, extent_type>;  ///< Storage ref type
 
   /**
-   * @brief Constructor of AoS storage ref.
+   * @brief Constructor of bucket storage.
+   *
+   * @note The input `size` should be exclusively determined by the return value of
+   * `make_bucket_extent` since it depends on the requested low-bound value, the probing scheme, and
+   * the storage.
    *
-   * @param size Number of windows
-   * @param windows Pointer to the windows array
+   * @param size Number of buckets to (de)allocate
+   * @param allocator Allocator used for (de)allocating device storage
    */
-  __host__ __device__ explicit constexpr aow_storage_ref(Extent size,
-                                                         window_type* windows) noexcept;
+  explicit constexpr bucket_storage(Extent size, Allocator const& allocator = {});
 
+  bucket_storage(bucket_storage&&) = default;  ///< Move constructor
   /**
-   * @brief Custom un-incrementable input iterator for the convenience of `find` operations.
+   * @brief Replaces the contents of the storage with another storage.
    *
-   * @note This iterator is for read only and NOT incrementable.
+   * @return Reference of the current storage object
    */
-  struct iterator;
-  using const_iterator = iterator const;  ///< Const forward iterator type
+  bucket_storage& operator=(bucket_storage&&) = default;
+  ~bucket_storage()                           = default;  ///< Destructor
+
+  bucket_storage(bucket_storage const&)            = delete;
+  bucket_storage& operator=(bucket_storage const&) = delete;
 
   /**
-   * @brief Returns an iterator to one past the last slot.
+   * @brief Gets buckets array.
    *
-   * This is provided for convenience for those familiar with checking
-   * an iterator returned from `find()` against the `end()` iterator.
-   *
-   * @return An iterator to one past the last slot
+   * @return Pointer to the first bucket
    */
-  [[nodiscard]] __device__ constexpr iterator end() noexcept;
+  [[nodiscard]] constexpr bucket_type* data() const noexcept;
 
   /**
-   * @brief Returns a const_iterator to one past the last slot.
-   *
-   * This is provided for convenience for those familiar with checking
-   * an iterator returned from `find()` against the `end()` iterator.
+   * @brief Gets the storage allocator.
    *
-   * @return A const_iterator to one past the last slot
+   * @return The storage allocator
    */
-  [[nodiscard]] __device__ constexpr const_iterator end() const noexcept;
+  [[nodiscard]] constexpr allocator_type allocator() const noexcept;
 
   /**
-   * @brief Gets windows array.
+   * @brief Gets bucket storage reference.
    *
-   * @return Pointer to the first window
+   * @return Reference of bucket storage
    */
-  [[nodiscard]] __device__ constexpr window_type* data() noexcept;
+  [[nodiscard]] constexpr ref_type ref() const noexcept;
 
   /**
-   * @brief Gets windows array.
+   * @brief Initializes each slot in the bucket storage to contain `key`.
    *
-   * @return Pointer to the first window
+   * @param key Key to which all keys in `slots` are initialized
+   * @param stream Stream used for executing the kernel
    */
-  [[nodiscard]] __device__ constexpr window_type* data() const noexcept;
+  void initialize(value_type key, cuda::stream_ref stream = {});
 
   /**
-   * @brief Returns an array of slots (or a window) for a given index.
+   * @brief Asynchronously initializes each slot in the bucket storage to contain `key`.
    *
-   * @param index Index of the window
-   * @return An array of slots
+   * @param key Key to which all keys in `slots` are initialized
+   * @param stream Stream used for executing the kernel
    */
-  [[nodiscard]] __device__ constexpr window_type operator[](size_type index) const noexcept;
+  void initialize_async(value_type key, cuda::stream_ref stream = {}) noexcept;
 
  private:
-  window_type* windows_;  ///< Pointer to the windows array
+  allocator_type allocator_;            ///< Allocator used to (de)allocate buckets
+  bucket_deleter_type bucket_deleter_;  ///< Custom buckets deleter
+  /// Pointer to the bucket storage
+  std::unique_ptr<bucket_type, bucket_deleter_type> buckets_;
 };
 
+/// Alias for bucket_storage_ref
+template <typename T, int32_t BucketSize, typename Extent = cuco::extent<std::size_t>>
+using aow_storage_ref = bucket_storage_ref<T, BucketSize, Extent>;
+
+/// Alias for bucket_storage
+template <typename T,
+          int32_t BucketSize,
+          typename Extent    = cuco::extent<std::size_t>,
+          typename Allocator = cuco::cuda_allocator<cuco::bucket<T, BucketSize>>>
+using aow_storage = bucket_storage<T, BucketSize, Extent, Allocator>;
+
 }  // namespace cuco
 
-#include <cuco/detail/storage/aow_storage.inl>
+#include <cuco/detail/storage/bucket_storage.inl>

include/cuco/detail/equal_wrapper.cuh

@@ -81,7 +81,7 @@ struct equal_wrapper {
    *
    * @note This function always compares the right-hand side element against sentinel values first
    * then performs a equality check with the given `equal_` callable, i.e., `equal_(lhs, rhs)`.
-   * @note Container (like set or map) buckets MUST be always on the right-hand side.
+   * @note Container (like set or map) slots MUST be always on the right-hand side.
    *
    * @tparam IsInsert Flag indicating whether it's an insert equality check or not. Insert probing
    * stops when it's an empty or erased slot while query probing stops only when it's empty.