diff --git a/include/cuco/detail/probe_sequence_impl.cuh b/include/cuco/detail/probe_sequence_impl.cuh index 3145e8fa5..549d59df3 100644 --- a/include/cuco/detail/probe_sequence_impl.cuh +++ b/include/cuco/detail/probe_sequence_impl.cuh @@ -108,6 +108,7 @@ class probe_sequence_impl_base { { } + public: /** * @brief Returns the capacity of the hash map. */ @@ -126,6 +127,7 @@ class probe_sequence_impl_base { */ __device__ __forceinline__ const_iterator get_slots() const noexcept { return slots_; } + protected: iterator slots_; ///< Pointer to beginning of the hash map slots const std::size_t capacity_; ///< Total number of slots }; // class probe_sequence_impl_base diff --git a/include/cuco/detail/static_multimap/device_view_impl.inl b/include/cuco/detail/static_multimap/device_view_impl.inl index 3a6afc5e1..5737b7af8 100644 --- a/include/cuco/detail/static_multimap/device_view_impl.inl +++ b/include/cuco/detail/static_multimap/device_view_impl.inl @@ -24,27 +24,6 @@ template class static_multimap::device_view_impl_base { - public: - /** - * @brief Gets the sentinel value used to represent an empty key slot. - * - * @return The sentinel value used to represent an empty key slot - */ - __host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept - { - return empty_key_sentinel_; - } - - /** - * @brief Gets the sentinel value used to represent an empty value slot. - * - * @return The sentinel value used to represent an empty value slot - */ - __host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept - { - return empty_value_sentinel_; - } - protected: // Import member type definitions from `static_multimap` using value_type = value_type; @@ -81,26 +60,6 @@ class static_multimap::device_view_ { } - /** - * @brief Gets slots array. - * - * @return Slots array - */ - __device__ __forceinline__ pair_atomic_type* get_slots() noexcept - { - return probe_sequence_.get_slots(); - } - - /** - * @brief Gets slots array. - * - * @return Slots array - */ - __device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept - { - return probe_sequence_.get_slots(); - } - /** * @brief Returns the initial slot for a given key `k` * @@ -161,16 +120,6 @@ class static_multimap::device_view_ return probe_sequence_.next_slot(s); } - /** - * @brief Gets the maximum number of elements the hash map can hold. - * - * @return The maximum number of elements the hash map can hold - */ - __host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept - { - return probe_sequence_.get_capacity(); - } - /** * @brief Load two key/value pairs from the given slot to the target pair array. * @@ -189,6 +138,57 @@ class static_multimap::device_view_ } } + public: + /** + * @brief Gets the sentinel value used to represent an empty key slot. + * + * @return The sentinel value used to represent an empty key slot + */ + __host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept + { + return empty_key_sentinel_; + } + + /** + * @brief Gets the sentinel value used to represent an empty value slot. + * + * @return The sentinel value used to represent an empty value slot + */ + __host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept + { + return empty_value_sentinel_; + } + + /** + * @brief Gets slots array. + * + * @return Slots array + */ + __device__ __forceinline__ pair_atomic_type* get_slots() noexcept + { + return probe_sequence_.get_slots(); + } + + /** + * @brief Gets slots array. + * + * @return Slots array + */ + __device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept + { + return probe_sequence_.get_slots(); + } + + /** + * @brief Gets the maximum number of elements the hash map can hold. + * + * @return The maximum number of elements the hash map can hold + */ + __host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept + { + return probe_sequence_.get_capacity(); + } + private: probe_sequence_type probe_sequence_; ///< Probe sequence used to probe the hash map Key empty_key_sentinel_{}; ///< Key value that represents an empty slot diff --git a/include/cuco/static_multimap.cuh b/include/cuco/static_multimap.cuh index 16b3ecd21..ab9135144 100644 --- a/include/cuco/static_multimap.cuh +++ b/include/cuco/static_multimap.cuh @@ -559,6 +559,78 @@ class static_multimap { class device_mutable_view_impl; class device_view_impl; + template + class device_view_base { + protected: + // Import member type definitions from `static_multimap` + using value_type = value_type; + using key_type = Key; + using mapped_type = Value; + using pair_atomic_type = pair_atomic_type; + using iterator = pair_atomic_type*; + using const_iterator = pair_atomic_type const*; + using probe_sequence_type = probe_sequence_type; + + __host__ __device__ device_view_base(pair_atomic_type* slots, + std::size_t capacity, + Key empty_key_sentinel, + Value empty_value_sentinel) noexcept + : impl_{slots, capacity, empty_key_sentinel, empty_value_sentinel} + { + } + + public: + /** + * @brief Gets slots array. + * + * @return Slots array + */ + __device__ __forceinline__ pair_atomic_type* get_slots() noexcept { return impl_.get_slots(); } + + /** + * @brief Gets slots array. + * + * @return Slots array + */ + __device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept + { + return impl_.get_slots(); + } + + /** + * @brief Gets the maximum number of elements the hash map can hold. + * + * @return The maximum number of elements the hash map can hold + */ + __host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept + { + return impl_.get_capacity(); + } + + /** + * @brief Gets the sentinel value used to represent an empty key slot. + * + * @return The sentinel value used to represent an empty key slot + */ + __host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept + { + return impl_.get_empty_key_sentinel(); + } + + /** + * @brief Gets the sentinel value used to represent an empty value slot. + * + * @return The sentinel value used to represent an empty value slot + */ + __host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept + { + return impl_.get_empty_value_sentinel(); + } + + protected: + ViewImpl impl_; + }; // class device_view_base + public: /** * @brief Mutable, non-owning view-type that may be used in device code to @@ -580,14 +652,14 @@ class static_multimap { * }); * \endcode */ - class device_mutable_view { + class device_mutable_view : public device_view_base { public: - // Import member type definitions from `static_multimap` - using value_type = value_type; - using key_type = Key; - using mapped_type = Value; - using iterator = pair_atomic_type*; - using const_iterator = pair_atomic_type const*; + using view_base_type = device_view_base; + using value_type = typename view_base_type::value_type; + using key_type = typename view_base_type::key_type; + using mapped_type = typename view_base_type::mapped_type; + using iterator = typename view_base_type::iterator; + using const_iterator = typename view_base_type::const_iterator; /** * @brief Construct a mutable view of the first `capacity` slots of the @@ -604,7 +676,7 @@ class static_multimap { std::size_t capacity, Key empty_key_sentinel, Value empty_value_sentinel) noexcept - : impl_{slots, capacity, empty_key_sentinel, empty_value_sentinel} + : view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel} { } @@ -620,7 +692,7 @@ class static_multimap { value_type const& insert_pair) noexcept; private: - device_mutable_view_impl impl_; + using device_view_base::impl_; }; // class device mutable view /** @@ -631,14 +703,14 @@ class static_multimap { * value. * */ - class device_view { + class device_view : public device_view_base { public: - // Import member type definitions from `static_multimap` - using value_type = value_type; - using key_type = Key; - using mapped_type = Value; - using iterator = pair_atomic_type*; - using const_iterator = pair_atomic_type const*; + using view_base_type = device_view_base; + using value_type = typename view_base_type::value_type; + using key_type = typename view_base_type::key_type; + using mapped_type = typename view_base_type::mapped_type; + using iterator = typename view_base_type::iterator; + using const_iterator = typename view_base_type::const_iterator; /** * @brief Construct a view of the first `capacity` slots of the @@ -655,55 +727,8 @@ class static_multimap { std::size_t capacity, Key empty_key_sentinel, Value empty_value_sentinel) noexcept - : impl_{slots, capacity, empty_key_sentinel, empty_value_sentinel} - { - } - - /** - * @brief Gets slots array. - * - * @return Slots array - */ - __device__ __forceinline__ pair_atomic_type* get_slots() noexcept { return impl_.get_slots(); } - - /** - * @brief Gets slots array. - * - * @return Slots array - */ - __device__ __forceinline__ pair_atomic_type const* get_slots() const noexcept - { - return impl_.get_slots(); - } - - /** - * @brief Gets the maximum number of elements the hash map can hold. - * - * @return The maximum number of elements the hash map can hold - */ - __host__ __device__ __forceinline__ std::size_t get_capacity() const noexcept - { - return impl_.get_capacity(); - } - - /** - * @brief Gets the sentinel value used to represent an empty key slot. - * - * @return The sentinel value used to represent an empty key slot - */ - __host__ __device__ __forceinline__ Key get_empty_key_sentinel() const noexcept + : view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel} { - return impl_.get_empty_key_sentinel(); - } - - /** - * @brief Gets the sentinel value used to represent an empty value slot. - * - * @return The sentinel value used to represent an empty value slot - */ - __host__ __device__ __forceinline__ Value get_empty_value_sentinel() const noexcept - { - return impl_.get_empty_value_sentinel(); } /** @@ -1053,7 +1078,7 @@ class static_multimap { PairEqual pair_equal) noexcept; private: - device_view_impl impl_; + using device_view_base::impl_; }; // class device_view /**