Move empty_value_sentinel to OA impl class

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 <bool Dummy = true, typename Enable = std::enable_if_t<has_payload and Dummy>>
+  [[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 <typename Value>
@@ -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 <typename Value, typename Enable = std::enable_if_t<has_payload and sizeof(Value)>>
+  [[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.
    *

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<Key,
   static_map_ref(
     static_map_ref<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, OtherOperators...>&&
       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<Key, T, Scope, KeyEqual, ProbingScheme, StorageRef, Operators...>::
   empty_value_sentinel() const noexcept
 {
-  return empty_value_sentinel_;
+  return impl_.empty_value_sentinel();
 }
 
 template <typename Key,

include/cuco/static_map_ref.cuh

@@ -161,8 +161,7 @@ class static_map_ref
   [[nodiscard]] __host__ __device__ auto with(NewOperators... ops) && noexcept;
 
  private:
-  impl_type impl_;                    ///< Static map ref implementation
-  mapped_type empty_value_sentinel_;  ///< Empty value sentinel
+  impl_type impl_;  ///< Static map ref implementation
 
   // Mixins need to be friends with this class in order to access private members
   template <typename Op, typename Ref>