Erase Functionality for static_map

include/cuco/detail/error.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -80,3 +80,23 @@ struct cuda_error : public std::runtime_error {
     cudaError_t const status = (expr); \
     assert(cudaSuccess == status);     \
   } while (0)
+
+/**
+ * @brief Macro for checking runtime conditions that throws an exception when
+ * a condition is violated.
+ *
+ * Example usage:
+ *
+ * @code
+ * CUCO_RUNTIME_EXPECTS(key == value, "Key value mismatch");
+ * @endcode
+ *
+ * @param[in] cond Expression that evaluates to true or false
+ * @param[in] reason String literal description of the reason that cond is
+ * expected to be true
+ * @throw std::runtime_error if the condition evaluates to false.
+ */
+#define CUCO_RUNTIME_EXPECTS(cond, reason)                           \
+  (!!(cond)) ? static_cast<void>(0)                                  \
+             : throw std::runtime_error("cuco failure at: " __FILE__ \
+                                       ":" CUCO_STRINGIFY(__LINE__) ": " reason)

include/cuco/detail/static_map.inl

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -15,6 +15,7 @@
  */
 
 #include <cuco/detail/bitwise_compare.cuh>
+#include <cuco/detail/error.hpp>
 
 namespace cuco {
 
@@ -146,15 +147,15 @@ void static_map<Key, Value, Scope, Allocator>::erase(
   InputIt first, InputIt last, Hash hash, KeyEqual key_equal, cudaStream_t stream)
 {
   if (get_empty_key_sentinel() == get_erased_key_sentinel())
-    throw std::runtime_error(
-      "Runtime error: You must provide a unique erased key sentinel value at map construction.\n");
+    CUCO_RUNTIME_EXPECTS(get_empty_key_sentinel() != get_erased_key_sentinel(),
+      "You must provide a unique erased key sentinel value at map construction.");
 
   auto num_keys = std::distance(first, last);
   if (num_keys == 0) { return; }
 
-  auto const block_size = 128;
-  auto const stride     = 1;
-  auto const tile_size  = 4;
+  auto constexpr block_size = 128;
+  auto constexpr stride     = 1;
+  auto constexpr tile_size  = 4;
   auto const grid_size  = (tile_size * num_keys + stride * block_size - 1) / (stride * block_size);
   auto view             = get_device_mutable_view();
 
@@ -438,8 +439,16 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::e
     make_pair<Key, Value>(this->get_erased_key_sentinel(), this->get_empty_value_sentinel());
 
   while (true) {
-    auto existing_key   = current_slot->first.load(cuda::std::memory_order_relaxed);
-    auto existing_value = current_slot->second.load(cuda::std::memory_order_relaxed);
+    //auto existing_key   = current_slot->first.load(cuda::std::memory_order_relaxed);
+    //auto existing_value = current_slot->second.load(cuda::std::memory_order_relaxed);
+
+    static_assert(sizeof(Key) == sizeof(atomic_key_type));
+    static_assert(sizeof(Value) == sizeof(atomic_mapped_type));
+    // TODO: Replace reinterpret_cast with atomic ref when available.
+    value_type slot_contents = *reinterpret_cast<value_type const*>(current_slot);
+    auto existing_key = slot_contents.first;
+    auto existing_value = slot_contents.second;
+
 
     // Key doesn't exist, return false
     if (detail::bitwise_compare(existing_key, this->get_empty_key_sentinel())) { return false; }
@@ -479,8 +488,12 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::e
     make_pair<Key, Value>(this->get_erased_key_sentinel(), this->get_empty_value_sentinel());
 
   while (true) {
-    auto existing_key   = current_slot->first.load(cuda::std::memory_order_relaxed);
-    auto existing_value = current_slot->second.load(cuda::std::memory_order_relaxed);
+    static_assert(sizeof(Key) == sizeof(atomic_key_type));
+    static_assert(sizeof(Value) == sizeof(atomic_mapped_type));
+    // TODO: Replace reinterpret_cast with atomic ref when available.
+    value_type slot_contents = *reinterpret_cast<value_type const*>(current_slot);
+    auto existing_key = slot_contents.first;
+    auto existing_value = slot_contents.second;
 
     auto const slot_is_empty =
       detail::bitwise_compare(existing_key, this->get_empty_key_sentinel());

include/cuco/detail/static_map_kernels.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -165,16 +165,15 @@ template <std::size_t block_size,
 __global__ void erase(
   InputIt first, InputIt last, atomicT* num_successes, viewT view, Hash hash, KeyEqual key_equal)
 {
-  typedef cub::BlockReduce<std::size_t, block_size> BlockReduce;
+  using BlockReduce = cub::BlockReduce<std::size_t, block_size>;
   __shared__ typename BlockReduce::TempStorage temp_storage;
   std::size_t thread_num_successes = 0;
 
   auto tid = block_size * blockIdx.x + threadIdx.x;
   auto it  = first + tid;
 
   while (it < last) {
-    auto k{*it};
-    if (view.erase(k, hash, key_equal)) { thread_num_successes++; }
+    if (view.erase(*it, hash, key_equal)) { thread_num_successes++; }
     it += gridDim.x * block_size;
   }
 
@@ -203,8 +202,7 @@ __global__ void erase(
   auto it   = first + tid / tile_size;
 
   while (it < last) {
-    auto k{*it};
-    if (view.erase(tile, k, hash, key_equal) && tile.thread_rank() == 0) { thread_num_successes++; }
+    if (view.erase(tile, *it, hash, key_equal) and tile.thread_rank() == 0) { thread_num_successes++; }
     it += (gridDim.x * block_size) / tile_size;
   }
 

include/cuco/static_map.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -373,7 +373,7 @@ class static_map {
     using slot_type      = slot_type;
 
     Key empty_key_sentinel_{};  ///< Key value that represents an empty slot
-    Key erased_key_sentinel_{};
+    Key erased_key_sentinel_{}; ///< Key value that represents an erased slot
     Value empty_value_sentinel_{};  ///< Initial Value of empty slot
     pair_atomic_type* slots_{};     ///< Pointer to flat slots storage
     std::size_t capacity_{};        ///< Total number of slots

tests/static_map/erase_test.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,8 +26,9 @@ TEMPLATE_TEST_CASE_SIG("erase key", "", ((typename T), T), (int32_t), (int64_t))
   using Key   = T;
   using Value = T;
 
-  unsigned long num_keys = 1'000'000;
-  cuco::static_map<Key, Value> map{num_keys * 1.1, -1, -1, -2};
+  constexpr std::size_t num_keys = 1'000'000;
+  constexpr std::size_t capacity = 1'100'000;
+  cuco::static_map<Key, Value> map{capacity, -1, -1, -2};
 
   auto m_view = map.get_device_mutable_view();
   auto view   = map.get_device_view();
@@ -81,16 +82,5 @@ TEMPLATE_TEST_CASE_SIG("erase key", "", ((typename T), T), (int32_t), (int64_t))
 
     map.erase(d_keys.begin() + num_keys / 2, d_keys.end());
     REQUIRE(map.get_size() == 0);
-
-    map.insert(pairs_begin, pairs_begin + num_keys / 2);
-    map.insert(pairs_begin + num_keys / 2, pairs_begin + num_keys);
-
-    map.erase(d_keys.begin(), d_keys.begin() + num_keys / 2);
-
-    map.contains(d_keys.begin() + num_keys / 2, d_keys.end(), d_keys_exist.begin());
-
-    REQUIRE(cuco::test::all_of(d_keys_exist.begin(),
-                               d_keys_exist.begin() + num_keys / 2,
-                               [] __device__(const bool key_found) { return key_found; }));
   }
 }