Merge pull request #163 from m3g4d1v3r/dev

Apply sentinel namespace to other map structures
NVIDIA · May 27, 2022 · b2af504 · b2af504
2 parents 917f1e5 + 4dc7a77
commit b2af504
Show file tree

Hide file tree

Showing 4 changed files with 28 additions and 24 deletions.
diff --git a/include/cuco/detail/dynamic_map.inl b/include/cuco/detail/dynamic_map.inl
@@ -17,12 +17,13 @@
 namespace cuco {
 
 template <typename Key, typename Value, cuda::thread_scope Scope, typename Allocator>
-dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(std::size_t initial_capacity,
-                                                       Key empty_key_sentinel,
-                                                       Value empty_value_sentinel,
-                                                       Allocator const& alloc)
-  : empty_key_sentinel_(empty_key_sentinel),
-    empty_value_sentinel_(empty_value_sentinel),
+dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(
+  std::size_t initial_capacity,
+  sentinel::empty_key<Key> empty_key_sentinel,
+  sentinel::empty_value<Value> empty_value_sentinel,
+  Allocator const& alloc)
+  : empty_key_sentinel_(empty_key_sentinel.value),
+    empty_value_sentinel_(empty_value_sentinel.value),
     size_(0),
     capacity_(initial_capacity),
     min_insert_size_(1E4),

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/dynamic_map.cuh b/include/cuco/dynamic_map.cuh
@@ -18,6 +18,7 @@
 
 #include <cuco/detail/dynamic_map_kernels.cuh>
 #include <cuco/detail/error.hpp>
+#include <cuco/sentinel.cuh>
 #include <cuco/static_map.cuh>
 
 #include <thrust/device_vector.h>
@@ -131,8 +132,8 @@ class dynamic_map {
    * @param alloc Allocator used to allocate submap device storage
    */
   dynamic_map(std::size_t initial_capacity,
-              Key empty_key_sentinel,
-              Value empty_value_sentinel,
+              sentinel::empty_key<Key> empty_key_sentinel,
+              sentinel::empty_value<Value> empty_value_sentinel,
               Allocator const& alloc = Allocator{});
 
   /**

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}
     {
     }
 
@@ -675,10 +676,11 @@ class static_multimap {
      * @param empty_value_sentinel The reserved value for mapped values to
      * represent empty slots
      */
-    __host__ __device__ device_mutable_view(pair_atomic_type* slots,
-                                            std::size_t capacity,
-                                            Key empty_key_sentinel,
-                                            Value empty_value_sentinel) noexcept
+    __host__ __device__
+    device_mutable_view(pair_atomic_type* slots,
+                        std::size_t capacity,
+                        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 +730,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}
     {
     }