From 5a99a56a81d6866b0ae1ecf3a669d64b2916f723 Mon Sep 17 00:00:00 2001 From: Daniel Juenger <2955913+sleeepyjack@users.noreply.github.com> Date: Thu, 5 Oct 2023 23:01:04 +0000 Subject: [PATCH] Move empty_value_sentinel to OA impl class --- .../open_addressing_ref_impl.cuh | 35 ++++++++++++++++++- .../cuco/detail/static_map/static_map_ref.inl | 11 +++--- include/cuco/static_map_ref.cuh | 3 +- 3 files changed, 39 insertions(+), 10 deletions(-) diff --git a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh index 960f741d0..2d6a09d23 100644 --- a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh +++ b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh @@ -146,12 +146,23 @@ class open_addressing_ref_impl { return this->predicate_.empty_sentinel_; } + /** + * @brief Gets the sentinel value used to represent an empty payload slot. + * + * @return The sentinel value used to represent an empty payload slot + */ + template > + [[nodiscard]] __host__ __device__ constexpr auto const& empty_value_sentinel() const noexcept + { + return this->extract_payload(this->empty_slot_sentinel()); + } + /** * @brief Gets the sentinel used to represent an empty slot. * * @return The sentinel value used to represent an empty slot */ - [[nodiscard]] __device__ constexpr value_type const& empty_slot_sentinel() const noexcept + [[nodiscard]] __host__ __device__ constexpr value_type const& empty_slot_sentinel() const noexcept { return empty_slot_sentinel_; } @@ -692,6 +703,10 @@ class open_addressing_ref_impl { /** * @brief Extracts the key from a given value type. * + * @tparam Value Input type which is implicitly convertible to 'value_type' + * + * @param value The input value + * * @return The key */ template @@ -705,6 +720,24 @@ class open_addressing_ref_impl { } } + /** + * @brief Extracts the payload from a given value type. + * + * @note This function is only available if `this->has_payload == true` + * + * @tparam Value Input type which is implicitly convertible to 'value_type' + * + * @param value The input value + * + * @return The payload + */ + template > + [[nodiscard]] __host__ __device__ constexpr auto const& extract_payload( + Value const& value) const noexcept + { + return value.second; + } + /** * @brief Inserts the specified element with one single CAS operation. * diff --git a/include/cuco/detail/static_map/static_map_ref.inl b/include/cuco/detail/static_map/static_map_ref.inl index f842fec06..b8a6350ab 100644 --- a/include/cuco/detail/static_map/static_map_ref.inl +++ b/include/cuco/detail/static_map/static_map_ref.inl @@ -44,11 +44,8 @@ __host__ __device__ constexpr static_map_ref< KeyEqual const& predicate, ProbingScheme const& probing_scheme, StorageRef storage_ref) noexcept - : impl_{cuco::pair{empty_key_sentinel, empty_value_sentinel}, - predicate, - probing_scheme, - storage_ref}, - empty_value_sentinel_{empty_value_sentinel} + : impl_{ + cuco::pair{empty_key_sentinel, empty_value_sentinel}, predicate, probing_scheme, storage_ref} { } @@ -70,7 +67,7 @@ __host__ __device__ constexpr static_map_ref&& other) noexcept - : impl_{std::move(other.impl_)}, empty_value_sentinel_{std::move(other.empty_value_sentinel_)} + : impl_{std::move(other.impl_)} { } @@ -113,7 +110,7 @@ __host__ __device__ constexpr T static_map_ref:: empty_value_sentinel() const noexcept { - return empty_value_sentinel_; + return impl_.empty_value_sentinel(); } template