Apply sentinel namespace to static_multimap

NVIDIA · May 26, 2022 · 360d8af · 360d8af
1 parent 9d2e198
commit 360d8af
Show file tree

Hide file tree

Showing 2 changed files with 15 additions and 14 deletions.
diff --git a/include/cuco/detail/static_multimap/static_multimap.inl b/include/cuco/detail/static_multimap/static_multimap.inl
@@ -33,14 +33,14 @@ template <typename Key,
           class ProbeSequence>
 static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::static_multimap(
   std::size_t capacity,
-  Key empty_key_sentinel,
-  Value empty_value_sentinel,
+  sentinel::empty_key<Key> empty_key_sentinel,
+  sentinel::empty_value<Value> empty_value_sentinel,
   cudaStream_t stream,
   Allocator const& alloc)
   : capacity_{cuco::detail::get_valid_capacity<cg_size(), vector_width(), uses_vector_load()>(
       capacity)},
-    empty_key_sentinel_{empty_key_sentinel},
-    empty_value_sentinel_{empty_value_sentinel},
+    empty_key_sentinel_{empty_key_sentinel.value},
+    empty_value_sentinel_{empty_value_sentinel.value},
     counter_allocator_{alloc},
     slot_allocator_{alloc},
     delete_counter_{counter_allocator_},
@@ -53,7 +53,7 @@ static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::static_multimap(
   auto const grid_size      = (get_capacity() + stride * block_size - 1) / (stride * block_size);
 
   detail::initialize<atomic_key_type, atomic_mapped_type><<<grid_size, block_size, 0, stream>>>(
-    slots_.get(), empty_key_sentinel, empty_value_sentinel, get_capacity());
+    slots_.get(), empty_key_sentinel_, empty_value_sentinel_, get_capacity());
 }
 
 template <typename Key,

diff --git a/include/cuco/static_multimap.cuh b/include/cuco/static_multimap.cuh
@@ -21,6 +21,7 @@
 #include <cuco/detail/prime.hpp>
 #include <cuco/detail/static_multimap/kernels.cuh>
 #include <cuco/probe_sequences.cuh>
+#include <cuco/sentinel.cuh>
 #include <cuco/traits.hpp>
 
 #include <thrust/functional.h>
@@ -223,8 +224,8 @@ class static_multimap {
    * @param alloc Allocator used for allocating device storage
    */
   static_multimap(std::size_t capacity,
-                  Key empty_key_sentinel,
-                  Value empty_value_sentinel,
+                  sentinel::empty_key<Key> empty_key_sentinel,
+                  sentinel::empty_value<Value> empty_value_sentinel,
                   cudaStream_t stream    = 0,
                   Allocator const& alloc = Allocator{});
 
@@ -576,9 +577,9 @@ class static_multimap {
 
     __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}
+                                         sentinel::empty_key<Key> empty_key_sentinel,
+                                         sentinel::empty_value<Value> empty_value_sentinel) noexcept
+      : impl_{slots, capacity, empty_key_sentinel.value, empty_value_sentinel.value}
     {
     }
 
@@ -677,8 +678,8 @@ class static_multimap {
      */
     __host__ __device__ device_mutable_view(pair_atomic_type* slots,
                                             std::size_t capacity,
-                                            Key empty_key_sentinel,
-                                            Value empty_value_sentinel) noexcept
+                                            sentinel::empty_key<Key> empty_key_sentinel,
+                                            sentinel::empty_value<Value> empty_value_sentinel) noexcept
       : view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel}
     {
     }
@@ -728,8 +729,8 @@ class static_multimap {
      */
     __host__ __device__ device_view(pair_atomic_type* slots,
                                     std::size_t capacity,
-                                    Key empty_key_sentinel,
-                                    Value empty_value_sentinel) noexcept
+                                    sentinel::empty_key<Key> empty_key_sentinel,
+                                    sentinel::empty_value<Value> empty_value_sentinel) noexcept
       : view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel}
     {
     }