From cf958d0607d42c679687959393b1c13aa2938dcb Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Thu, 10 Aug 2023 10:57:08 -0700 Subject: [PATCH] Expose `aow_storage` to the public (#349) Closes #348 This PR exposes `cuco::experimental::window` type and `aow_storage` class to the public. --- include/cuco/aow_storage.cuh | 218 ++++++++++ include/cuco/detail/storage/aow_storage.cuh | 372 ------------------ include/cuco/detail/storage/aow_storage.inl | 197 ++++++++++ .../cuco/detail/storage/aow_storage_base.cuh | 106 +++++ include/cuco/detail/storage/storage.cuh | 2 +- include/cuco/detail/storage/storage_base.cuh | 2 +- include/cuco/static_map.cuh | 2 +- include/cuco/static_set.cuh | 2 +- include/cuco/storage.cuh | 22 +- tests/static_map/unique_sequence_test.cu | 2 +- tests/static_set/capacity_test.cu | 2 +- tests/static_set/insert_and_find_test.cu | 2 +- tests/static_set/retrieve_all_test.cu | 2 +- tests/static_set/unique_sequence_test.cu | 2 +- tests/utility/storage_test.cu | 22 +- 15 files changed, 553 insertions(+), 402 deletions(-) create mode 100644 include/cuco/aow_storage.cuh delete mode 100644 include/cuco/detail/storage/aow_storage.cuh create mode 100644 include/cuco/detail/storage/aow_storage.inl create mode 100644 include/cuco/detail/storage/aow_storage_base.cuh diff --git a/include/cuco/aow_storage.cuh b/include/cuco/aow_storage.cuh new file mode 100644 index 000000000..fdd970cf4 --- /dev/null +++ b/include/cuco/aow_storage.cuh @@ -0,0 +1,218 @@ +/* + * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +#include +#include + +#include + +#include +#include +#include +#include + +namespace cuco { +namespace experimental { + +/// Window type alias +template +using window = detail::window; + +/// forward declaration +template +class aow_storage_ref; + +/** + * @brief Array of Window open addressing storage class. + * + * @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 + */ +template +class aow_storage : public detail::aow_storage_base { + public: + using base_type = detail::aow_storage_base; ///< AoW base class type + + using base_type::window_size; ///< Number of elements processed per window + + 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 base_type::capacity; + using base_type::num_windows; + + /// Type of the allocator to (de)allocate windows + using allocator_type = typename std::allocator_traits::rebind_alloc; + using window_deleter_type = + detail::custom_deleter; ///< Type of window deleter + using ref_type = aow_storage_ref; ///< Storage ref type + + /** + * @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. + * + * @param size Number of windows to (de)allocate + * @param allocator Allocator used for (de)allocating device storage + */ + explicit constexpr aow_storage(Extent size, Allocator const& allocator) noexcept; + + aow_storage(aow_storage&&) = default; ///< Move constructor + /** + * @brief Replaces the contents of the storage with another storage. + * + * @return Reference of the current storage object + */ + aow_storage& operator=(aow_storage&&) = default; + ~aow_storage() = default; ///< Destructor + + aow_storage(aow_storage const&) = delete; + aow_storage& operator=(aow_storage const&) = delete; + + /** + * @brief Gets windows array. + * + * @return Pointer to the first window + */ + [[nodiscard]] constexpr window_type* data() const noexcept; + + /** + * @brief Gets the storage allocator. + * + * @return The storage allocator + */ + [[nodiscard]] constexpr allocator_type allocator() const noexcept; + + /** + * @brief Gets window storage reference. + * + * @return Reference of window storage + */ + [[nodiscard]] constexpr ref_type ref() const noexcept; + + /** + * @brief Initializes each slot in the AoW storage to contain `key`. + * + * @param key Key to which all keys in `slots` are initialized + * @param stream Stream used for executing the kernel + */ + void initialize(value_type key, cuda_stream_ref stream) noexcept; + + private: + allocator_type allocator_; ///< Allocator used to (de)allocate windows + window_deleter_type window_deleter_; ///< Custom windows deleter + std::unique_ptr windows_; ///< Pointer to AoW storage +}; + +/** + * @brief Non-owning AoW storage reference type. + * + * @tparam T Storage element type + * @tparam WindowSize Number of slots in each window + * @tparam Extent Type of extent denoting storage capacity + */ +template +class aow_storage_ref : public detail::aow_storage_base { + public: + using base_type = detail::aow_storage_base; ///< AoW base class type + + using base_type::window_size; ///< Number of elements processed per window + + 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 base_type::capacity; + using base_type::num_windows; + + /** + * @brief Constructor of AoS storage ref. + * + * @param size Number of windows + * @param windows Pointer to the windows array + */ + __host__ __device__ explicit constexpr aow_storage_ref(Extent size, + window_type* windows) noexcept; + + /** + * @brief Custom un-incrementable input iterator for the convenience of `find` operations. + * + * @note This iterator is for read only and NOT incrementable. + */ + struct iterator; + using const_iterator = iterator const; ///< Const forward iterator type + + /** + * @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 An iterator to one past the last slot + */ + [[nodiscard]] __device__ constexpr iterator end() 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. + * + * @return A const_iterator to one past the last slot + */ + [[nodiscard]] __device__ constexpr const_iterator end() const noexcept; + + /** + * @brief Gets windows array. + * + * @return Pointer to the first window + */ + [[nodiscard]] __device__ constexpr window_type* data() noexcept; + + /** + * @brief Gets windows array. + * + * @return Pointer to the first window + */ + [[nodiscard]] __device__ constexpr window_type* data() const noexcept; + + /** + * @brief Returns an array of slots (or a window) for a given index. + * + * @param index Index of the window + * @return An array of slots + */ + [[nodiscard]] __device__ constexpr window_type operator[](size_type index) const noexcept; + + private: + window_type* windows_; ///< Pointer to the windows array +}; + +} // namespace experimental +} // namespace cuco + +#include diff --git a/include/cuco/detail/storage/aow_storage.cuh b/include/cuco/detail/storage/aow_storage.cuh deleted file mode 100644 index e817e3293..000000000 --- a/include/cuco/detail/storage/aow_storage.cuh +++ /dev/null @@ -1,372 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include -#include -#include -#include - -#include - -#include -#include -#include -#include - -namespace cuco { -namespace experimental { -namespace detail { -/** - * @brief Base class of array of slot windows open addressing storage. - * - * This should NOT be used directly. - * - * @tparam WindowSize Number of elements in each window - * @tparam T Element type - * @tparam Extent Type of extent denoting the number of windows - */ -template -class aow_storage_base : public storage_base { - public: - /** - * @brief The number of elements (slots) processed per window. - */ - static constexpr int32_t window_size = WindowSize; - - using extent_type = typename storage_base::extent_type; ///< Storage extent type - using size_type = typename storage_base::size_type; ///< Storage size type - - using value_type = T; ///< Slot type - using window_type = cuda::std::array; ///< Slot window type - - /** - * @brief Constructor of AoW base storage. - * - * @param size Number of windows to store - */ - explicit constexpr aow_storage_base(Extent size) : storage_base{size} {} - - /** - * @brief Gets the total number of slot windows in the current storage. - * - * @return The total number of slot windows - */ - [[nodiscard]] __host__ __device__ constexpr size_type num_windows() const noexcept - { - return storage_base::capacity(); - } - - /** - * @brief Gets the total number of slots in the current storage. - * - * @return The total number of slots - */ - [[nodiscard]] __host__ __device__ constexpr size_type capacity() const noexcept - { - return storage_base::capacity() * window_size; - } - - /** - * @brief Gets the window extent of the current storage. - * - * @return The window extent. - */ - [[nodiscard]] __host__ __device__ constexpr extent_type window_extent() const noexcept - { - return storage_base::extent(); - } -}; - -/** - * @brief Non-owning AoW storage reference type. - * - * @tparam WindowSize Number of slots in each window - * @tparam T Storage element type - * @tparam Extent Type of extent denoting storage capacity - */ -template -class aow_storage_ref : public aow_storage_base { - public: - using base_type = aow_storage_base; ///< AoW base class type - - using base_type::window_size; ///< Number of elements processed per window - - 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 base_type::capacity; - using base_type::num_windows; - - /** - * @brief Constructor of AoS storage ref. - * - * @param windows Pointer to the windows array - * @param num_windows Number of windows - */ - explicit constexpr aow_storage_ref(Extent num_windows, window_type* windows) noexcept - : aow_storage_base{num_windows}, windows_{windows} - { - } - - /** - * @brief Custom un-incrementable input iterator for the convenience of `find` operations. - * - * @note This iterator is for read only and NOT incrementable. - */ - struct iterator { - public: - using iterator_category = std::input_iterator_tag; ///< iterator category - using reference = value_type&; ///< iterator reference type - - /** - * @brief Constructs a device side input iterator of the given slot. - * - * @param current The slot pointer - */ - __device__ constexpr explicit iterator(value_type* current) noexcept : current_{current} {} - - /** - * @brief Prefix increment operator - * - * @throw This code path should never be chosen. - * - * @return Current iterator - */ - __device__ constexpr iterator& operator++() noexcept - { - static_assert("Un-incrementable input iterator"); - } - - /** - * @brief Postfix increment operator - * - * @throw This code path should never be chosen. - * - * @return Current iterator - */ - __device__ constexpr iterator operator++(int32_t) noexcept - { - static_assert("Un-incrementable input iterator"); - } - - /** - * @brief Dereference operator - * - * @return Reference to the current slot - */ - __device__ constexpr reference operator*() const { return *current_; } - - /** - * @brief Access operator - * - * @return Pointer to the current slot - */ - __device__ constexpr value_type* operator->() const { return current_; } - - /** - * Equality operator - * - * @return True if two iterators are identical - */ - friend __device__ constexpr bool operator==(iterator const& lhs, iterator const& rhs) noexcept - { - return lhs.current_ == rhs.current_; - } - - /** - * Inequality operator - * - * @return True if two iterators are not identical - */ - friend __device__ constexpr bool operator!=(iterator const& lhs, iterator const& rhs) noexcept - { - return not(lhs == rhs); - } - - private: - value_type* current_{}; ///< Pointer to the current slot - }; - using const_iterator = iterator const; ///< Const forward iterator type - - /** - * @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 An iterator to one past the last slot - */ - [[nodiscard]] __device__ constexpr iterator end() noexcept - { - return iterator{reinterpret_cast(this->data() + this->capacity())}; - } - - /** - * @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. - * - * @return A const_iterator to one past the last slot - */ - [[nodiscard]] __device__ constexpr const_iterator end() const noexcept - { - return const_iterator{reinterpret_cast(this->data() + this->capacity())}; - } - - /** - * @brief Gets windows array. - * - * @return Pointer to the first window - */ - [[nodiscard]] __device__ constexpr window_type* data() noexcept { return windows_; } - - /** - * @brief Gets windows array. - * - * @return Pointer to the first window - */ - [[nodiscard]] __device__ constexpr window_type* data() const noexcept { return windows_; } - - /** - * @brief Returns an array of slots (or a window) for a given index. - * - * @param index Index of the window - * @return An array of slots - */ - [[nodiscard]] __device__ constexpr window_type operator[](size_type index) const noexcept - { - return *reinterpret_cast( - __builtin_assume_aligned(this->data() + index, sizeof(value_type) * window_size)); - } - - private: - window_type* windows_; ///< Pointer to the windows array -}; - -/** - * @brief Array of slot Window open addressing storage class. - * - * @tparam WindowSize Number of slots in each window - * @tparam T Slot type - * @tparam Extent Type of extent denoting number of windows - * @tparam Allocator Type of allocator used for device storage (de)allocation - */ -template -class aow_storage : public aow_storage_base { - public: - using base_type = aow_storage_base; ///< AoW base class type - - using base_type::window_size; ///< Number of elements processed per window - - 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 base_type::capacity; - using base_type::num_windows; - - /// Type of the allocator to (de)allocate windows - using allocator_type = typename std::allocator_traits::rebind_alloc; - using window_deleter_type = - custom_deleter; ///< Type of window deleter - using ref_type = aow_storage_ref; ///< Storage ref type - - /** - * @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. - * - * @param size Number of windows to (de)allocate - * @param allocator Allocator used for (de)allocating device storage - */ - explicit constexpr aow_storage(Extent size, Allocator const& allocator) - : aow_storage_base{size}, - allocator_{allocator}, - window_deleter_{capacity(), allocator_}, - windows_{allocator_.allocate(capacity()), window_deleter_} - { - } - - aow_storage(aow_storage&&) = default; ///< Move constructor - /** - * @brief Replaces the contents of the storage with another storage. - * - * @return Reference of the current storage object - */ - aow_storage& operator=(aow_storage&&) = default; - ~aow_storage() = default; ///< Destructor - - aow_storage(aow_storage const&) = delete; - aow_storage& operator=(aow_storage const&) = delete; - - /** - * @brief Gets windows array. - * - * @return Pointer to the first window - */ - [[nodiscard]] constexpr window_type* data() const noexcept { return windows_.get(); } - - /** - * @brief Gets the storage allocator. - * - * @return The storage allocator - */ - [[nodiscard]] constexpr allocator_type allocator() const noexcept { return allocator_; } - - /** - * @brief Gets window storage reference. - * - * @return Reference of window storage - */ - [[nodiscard]] constexpr ref_type ref() const noexcept - { - return ref_type{this->window_extent(), this->data()}; - } - - /** - * @brief Initializes each slot in the AoW storage to contain `key`. - * - * @param key Key to which all keys in `slots` are initialized - * @param stream Stream used for executing the kernel - */ - void initialize(value_type key, cuda_stream_ref stream) noexcept - { - auto constexpr stride = 4; - auto const grid_size = (this->num_windows() + stride * detail::CUCO_DEFAULT_BLOCK_SIZE - 1) / - (stride * detail::CUCO_DEFAULT_BLOCK_SIZE); - - detail::initialize<<>>( - this->data(), this->num_windows(), key); - } - - private: - allocator_type allocator_; ///< Allocator used to (de)allocate windows - window_deleter_type window_deleter_; ///< Custom windows deleter - std::unique_ptr windows_; ///< Pointer to AoW storage -}; - -} // namespace detail -} // namespace experimental -} // namespace cuco diff --git a/include/cuco/detail/storage/aow_storage.inl b/include/cuco/detail/storage/aow_storage.inl new file mode 100644 index 000000000..b4052b2a0 --- /dev/null +++ b/include/cuco/detail/storage/aow_storage.inl @@ -0,0 +1,197 @@ +/* + * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include +#include +#include +#include + +#include + +#include +#include +#include +#include + +namespace cuco { +namespace experimental { + +template +constexpr aow_storage::aow_storage( + Extent size, Allocator const& allocator) noexcept + : detail::aow_storage_base{size}, + allocator_{allocator}, + window_deleter_{capacity(), allocator_}, + windows_{allocator_.allocate(capacity()), window_deleter_} +{ +} + +template +constexpr aow_storage::window_type* +aow_storage::data() const noexcept +{ + return windows_.get(); +} + +template +constexpr aow_storage::allocator_type +aow_storage::allocator() const noexcept +{ + return allocator_; +} + +template +constexpr aow_storage::ref_type +aow_storage::ref() const noexcept +{ + return ref_type{this->window_extent(), this->data()}; +} + +template +void aow_storage::initialize(value_type key, + cuda_stream_ref stream) noexcept +{ + auto constexpr stride = 4; + auto const grid_size = (this->num_windows() + stride * detail::CUCO_DEFAULT_BLOCK_SIZE - 1) / + (stride * detail::CUCO_DEFAULT_BLOCK_SIZE); + + detail::initialize<<>>( + this->data(), this->num_windows(), key); +} + +template +__host__ __device__ constexpr aow_storage_ref::aow_storage_ref( + Extent size, window_type* windows) noexcept + : detail::aow_storage_base{size}, windows_{windows} +{ +} + +template +struct aow_storage_ref::iterator { + public: + using iterator_category = std::input_iterator_tag; ///< iterator category + using reference = value_type&; ///< iterator reference type + + /** + * @brief Constructs a device side input iterator of the given slot. + * + * @param current The slot pointer + */ + __device__ constexpr explicit iterator(value_type* current) noexcept : current_{current} {} + + /** + * @brief Prefix increment operator + * + * @throw This code path should never be chosen. + * + * @return Current iterator + */ + __device__ constexpr iterator& operator++() noexcept + { + static_assert("Un-incrementable input iterator"); + } + + /** + * @brief Postfix increment operator + * + * @throw This code path should never be chosen. + * + * @return Current iterator + */ + __device__ constexpr iterator operator++(int32_t) noexcept + { + static_assert("Un-incrementable input iterator"); + } + + /** + * @brief Dereference operator + * + * @return Reference to the current slot + */ + __device__ constexpr reference operator*() const { return *current_; } + + /** + * @brief Access operator + * + * @return Pointer to the current slot + */ + __device__ constexpr value_type* operator->() const { return current_; } + + /** + * Equality operator + * + * @return True if two iterators are identical + */ + friend __device__ constexpr bool operator==(iterator const& lhs, iterator const& rhs) noexcept + { + return lhs.current_ == rhs.current_; + } + + /** + * Inequality operator + * + * @return True if two iterators are not identical + */ + friend __device__ constexpr bool operator!=(iterator const& lhs, iterator const& rhs) noexcept + { + return not(lhs == rhs); + } + + private: + value_type* current_{}; ///< Pointer to the current slot +}; + +template +__device__ constexpr aow_storage_ref::iterator +aow_storage_ref::end() noexcept +{ + return iterator{reinterpret_cast(this->data() + this->capacity())}; +} + +template +__device__ constexpr aow_storage_ref::const_iterator +aow_storage_ref::end() const noexcept +{ + return const_iterator{reinterpret_cast(this->data() + this->capacity())}; +} + +template +__device__ constexpr aow_storage_ref::window_type* +aow_storage_ref::data() noexcept +{ + return windows_; +} + +template +__device__ constexpr aow_storage_ref::window_type* +aow_storage_ref::data() const noexcept +{ + return windows_; +} + +template +__device__ constexpr aow_storage_ref::window_type +aow_storage_ref::operator[](size_type index) const noexcept +{ + return *reinterpret_cast( + __builtin_assume_aligned(this->data() + index, sizeof(value_type) * window_size)); +} + +} // namespace experimental +} // namespace cuco diff --git a/include/cuco/detail/storage/aow_storage_base.cuh b/include/cuco/detail/storage/aow_storage_base.cuh new file mode 100644 index 000000000..5f3d84df4 --- /dev/null +++ b/include/cuco/detail/storage/aow_storage_base.cuh @@ -0,0 +1,106 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +#include + +#include +#include + +namespace cuco { +namespace experimental { +namespace detail { +/** + * @brief Window data structure type + * + * @tparam T Window slot type + * @tparam WindowSize Number of elements per window + */ +template +struct window : public cuda::std::array { + public: + static int32_t constexpr window_size = WindowSize; ///< Number of slots per window +}; + +/** + * @brief Base class of array of slot windows open addressing storage. + * + * @note This should NOT be used directly. + * + * @tparam T Slot type + * @tparam WindowSize Number of slots in each window + * @tparam Extent Type of extent denoting the number of windows + */ +template +class aow_storage_base : public storage_base { + public: + /** + * @brief The number of elements (slots) processed per window. + */ + static constexpr int32_t window_size = WindowSize; + + using extent_type = typename storage_base::extent_type; ///< Storage extent type + using size_type = typename storage_base::size_type; ///< Storage size type + + using value_type = T; ///< Slot type + using window_type = window; ///< Slot window type + + /** + * @brief Constructor of AoW base storage. + * + * @param size Number of windows to store + */ + __host__ __device__ explicit constexpr aow_storage_base(Extent size) : storage_base{size} + { + } + + /** + * @brief Gets the total number of slot windows in the current storage. + * + * @return The total number of slot windows + */ + [[nodiscard]] __host__ __device__ constexpr size_type num_windows() const noexcept + { + return storage_base::capacity(); + } + + /** + * @brief Gets the total number of slots in the current storage. + * + * @return The total number of slots + */ + [[nodiscard]] __host__ __device__ constexpr size_type capacity() const noexcept + { + return storage_base::capacity() * window_size; + } + + /** + * @brief Gets the window extent of the current storage. + * + * @return The window extent. + */ + [[nodiscard]] __host__ __device__ constexpr extent_type window_extent() const noexcept + { + return storage_base::extent(); + } +}; + +} // namespace detail +} // namespace experimental +} // namespace cuco diff --git a/include/cuco/detail/storage/storage.cuh b/include/cuco/detail/storage/storage.cuh index 108aa7f84..b9a00baa2 100644 --- a/include/cuco/detail/storage/storage.cuh +++ b/include/cuco/detail/storage/storage.cuh @@ -16,7 +16,7 @@ #pragma once -#include +#include namespace cuco { namespace experimental { diff --git a/include/cuco/detail/storage/storage_base.cuh b/include/cuco/detail/storage/storage_base.cuh index 15ec30472..98eed6c13 100644 --- a/include/cuco/detail/storage/storage_base.cuh +++ b/include/cuco/detail/storage/storage_base.cuh @@ -71,7 +71,7 @@ class storage_base { * * @param size Number of elements to (de)allocate */ - explicit constexpr storage_base(Extent size) : extent_{size} {} + __host__ __device__ explicit constexpr storage_base(Extent size) : extent_{size} {} /** * @brief Gets the total number of elements in the current storage. diff --git a/include/cuco/static_map.cuh b/include/cuco/static_map.cuh index 38f3b92c9..2df5b2a10 100644 --- a/include/cuco/static_map.cuh +++ b/include/cuco/static_map.cuh @@ -94,7 +94,7 @@ template >, class Allocator = cuco::cuda_allocator>, - class Storage = cuco::experimental::aow_storage<1>> + class Storage = cuco::experimental::storage<1>> class static_map { static_assert(sizeof(Key) <= 8, "Container does not support key types larger than 8 bytes."); diff --git a/include/cuco/static_set.cuh b/include/cuco/static_set.cuh index 0cb558a65..a7eee42ac 100644 --- a/include/cuco/static_set.cuh +++ b/include/cuco/static_set.cuh @@ -87,7 +87,7 @@ template >, class Allocator = cuco::cuda_allocator, - class Storage = cuco::experimental::aow_storage<1>> + class Storage = cuco::experimental::storage<1>> class static_set { using impl_type = detail:: open_addressing_impl; diff --git a/include/cuco/storage.cuh b/include/cuco/storage.cuh index 969b49f37..e34e59c96 100644 --- a/include/cuco/storage.cuh +++ b/include/cuco/storage.cuh @@ -21,26 +21,28 @@ namespace cuco { namespace experimental { /** - * @brief Public Array of slot Windows storage class. + * @brief Public storage class. * - * The window size defines the workload granularity for each CUDA thread, i.e., how many slots a - * thread would concurrently operate on when performing modify or lookup operations. cuCollections - * uses the AoW storage to supersede the raw flat slot storage due to its superior granularity - * control: When window size equals one, AoW performs the same as the flat storage. If the - * underlying operation is more memory bandwidth bound, e.g., high occupancy multimap operations, a - * larger window size can reduce the length of probing sequences thus improve runtime performance. + * @note This is a public interface used to control storage window size. A window consists of a + * number of contiguous slots. The window size defines the workload granularity for each CUDA + * thread, i.e., how many slots a thread would concurrently operate on when performing modify or + * lookup operations. cuCollections uses the AoW storage to supersede the raw flat slot storage due + * to its superior granularity control: When window size equals one, AoW performs the same as the + * flat storage. If the underlying operation is more memory bandwidth bound, e.g., high occupancy + * multimap operations, a larger window size can reduce the length of probing sequences thus improve + * runtime performance. * * @tparam WindowSize Number of elements per window storage */ template -class aow_storage { +class storage { public: - /// Number of elements per window storage + /// Number of slots per window storage static constexpr int32_t window_size = WindowSize; /// Type of implementation details template - using impl = detail::aow_storage; + using impl = aow_storage; }; } // namespace experimental diff --git a/tests/static_map/unique_sequence_test.cu b/tests/static_map/unique_sequence_test.cu index d0581afdb..6a0165cc2 100644 --- a/tests/static_map/unique_sequence_test.cu +++ b/tests/static_map/unique_sequence_test.cu @@ -295,7 +295,7 @@ TEMPLATE_TEST_CASE_SIG( thrust::equal_to, probe, cuco::cuda_allocator, - cuco::experimental::aow_storage<2>>{ + cuco::experimental::storage<2>>{ num_keys, cuco::empty_key{-1}, cuco::empty_value{-1}}; REQUIRE(map.capacity() == gold_capacity); diff --git a/tests/static_set/capacity_test.cu b/tests/static_set/capacity_test.cu index 3b7681e0a..4c66a7ccc 100644 --- a/tests/static_set/capacity_test.cu +++ b/tests/static_set/capacity_test.cu @@ -24,7 +24,7 @@ TEST_CASE("Static set capacity", "") using ProbeT = cuco::experimental::double_hashing<1, cuco::default_hash_function>; using Equal = thrust::equal_to; using AllocatorT = cuco::cuda_allocator; - using StorageT = cuco::experimental::aow_storage<2>; + using StorageT = cuco::experimental::storage<2>; SECTION("zero capacity is allowed.") { diff --git a/tests/static_set/insert_and_find_test.cu b/tests/static_set/insert_and_find_test.cu index 9d0cc057a..278510e08 100644 --- a/tests/static_set/insert_and_find_test.cu +++ b/tests/static_set/insert_and_find_test.cu @@ -104,7 +104,7 @@ TEMPLATE_TEST_CASE_SIG( thrust::equal_to, probe, cuco::cuda_allocator, - cuco::experimental::aow_storage<2>>{ + cuco::experimental::storage<2>>{ num_keys, cuco::empty_key{-1}}; test_insert_and_find(set, num_keys); } diff --git a/tests/static_set/retrieve_all_test.cu b/tests/static_set/retrieve_all_test.cu index 97a489455..616e35138 100644 --- a/tests/static_set/retrieve_all_test.cu +++ b/tests/static_set/retrieve_all_test.cu @@ -86,7 +86,7 @@ TEMPLATE_TEST_CASE_SIG( thrust::equal_to, probe, cuco::cuda_allocator, - cuco::experimental::aow_storage<1>>{ + cuco::experimental::storage<1>>{ num_keys, cuco::empty_key{-1}}; REQUIRE(set.capacity() == gold_capacity); diff --git a/tests/static_set/unique_sequence_test.cu b/tests/static_set/unique_sequence_test.cu index 4c037463a..53ede7524 100644 --- a/tests/static_set/unique_sequence_test.cu +++ b/tests/static_set/unique_sequence_test.cu @@ -143,7 +143,7 @@ TEMPLATE_TEST_CASE_SIG( thrust::equal_to, probe, cuco::cuda_allocator, - cuco::experimental::aow_storage<2>>{ + cuco::experimental::storage<2>>{ num_keys, cuco::empty_key{-1}}; REQUIRE(set.capacity() == gold_capacity); diff --git a/tests/utility/storage_test.cu b/tests/utility/storage_test.cu index afb9848d3..b776f628c 100644 --- a/tests/utility/storage_test.cu +++ b/tests/utility/storage_test.cu @@ -16,7 +16,7 @@ #include -#include +#include #include #include #include @@ -39,11 +39,11 @@ TEMPLATE_TEST_CASE_SIG("Storage tests", SECTION("Allocate array of pairs with AoS storage.") { - auto s = cuco::experimental::detail::aow_storage, - cuco::experimental::extent, - allocator_type>( - cuco::experimental::extent{size}, allocator); + auto s = + cuco::experimental::aow_storage, + window_size, + cuco::experimental::extent, + allocator_type>(cuco::experimental::extent{size}, allocator); auto const num_windows = s.num_windows(); auto const capacity = s.capacity(); @@ -54,8 +54,8 @@ TEMPLATE_TEST_CASE_SIG("Storage tests", SECTION("Allocate array of pairs with AoS storage with static extent.") { using extent_type = cuco::experimental::extent; - auto s = cuco::experimental::detail:: - aow_storage, extent_type, allocator_type>(extent_type{}, + auto s = cuco::experimental:: + aow_storage, window_size, extent_type, allocator_type>(extent_type{}, allocator); auto const num_windows = s.num_windows(); auto const capacity = s.capacity(); @@ -66,8 +66,8 @@ TEMPLATE_TEST_CASE_SIG("Storage tests", SECTION("Allocate array of keys with AoS storage.") { - auto s = cuco::experimental::detail:: - aow_storage, allocator_type>( + auto s = cuco::experimental:: + aow_storage, allocator_type>( cuco::experimental::extent{size}, allocator); auto const num_windows = s.num_windows(); auto const capacity = s.capacity(); @@ -79,7 +79,7 @@ TEMPLATE_TEST_CASE_SIG("Storage tests", SECTION("Allocate array of keys with AoS storage with static extent.") { using extent_type = cuco::experimental::extent; - auto s = cuco::experimental::detail::aow_storage( + auto s = cuco::experimental::aow_storage( extent_type{}, allocator); auto const num_windows = s.num_windows(); auto const capacity = s.capacity();