Erase Functionality for static_map

benchmarks/hash_table/static_map_bench.cu

@@ -87,16 +87,24 @@ static void BM_static_map_insert(::benchmark::State& state)
   }
 
   thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
+  thrust::device_vector<Key> d_keys(h_keys);
 
   for (auto _ : state) {
-    state.ResumeTiming();
-    state.PauseTiming();
-    map_type map{size, -1, -1};
-    state.ResumeTiming();
+    map_type map{size, cuco::sentinel::empty_key<Key>{-1}, cuco::sentinel::empty_value<Value>{-1}};
 
+    cudaEvent_t start, stop;
+    cudaEventCreate(&start);
+    cudaEventCreate(&stop);
+
+    cudaEventRecord(start);
     map.insert(d_pairs.begin(), d_pairs.end());
+    cudaEventRecord(stop);
+    cudaEventSynchronize(stop);
 
-    state.PauseTiming();
+    float ms;
+    cudaEventElapsedTime(&ms, start, stop);
+
+    state.SetIterationTime(ms / 1000);
   }
 
   state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
@@ -112,7 +120,7 @@ static void BM_static_map_search_all(::benchmark::State& state)
   float occupancy      = state.range(1) / float{100};
   std::size_t size     = num_keys / occupancy;
 
-  map_type map{size, -1, -1};
+  map_type map{size, cuco::sentinel::empty_key<Key>{-1}, cuco::sentinel::empty_value<Value>{-1}};
   auto view = map.get_device_mutable_view();
 
   std::vector<Key> h_keys(num_keys);
@@ -143,50 +151,62 @@ static void BM_static_map_search_all(::benchmark::State& state)
                           int64_t(state.range(0)));
 }
 
-BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIQUE)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+template <typename Key, typename Value, dist_type Dist>
+static void BM_static_map_erase_all(::benchmark::State& state)
+{
+  using map_type = cuco::static_map<Key, Value>;
 
-BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::UNIQUE)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  std::size_t num_keys = state.range(0);
+  float occupancy      = state.range(1) / float{100};
+  std::size_t size     = num_keys / occupancy;
 
-BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIFORM)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  // static map with erase support
+  map_type map{size,
+               cuco::sentinel::empty_key<Key>{-1},
+               cuco::sentinel::empty_value<Value>{-1},
+               cuco::sentinel::erased_key<Key>{-2}};
+  auto view = map.get_device_mutable_view();
 
-BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::UNIFORM)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  std::vector<Key> h_keys(num_keys);
+  std::vector<Value> h_values(num_keys);
+  std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);
+  std::vector<Value> h_results(num_keys);
 
-BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::GAUSSIAN)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
 
-BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::GAUSSIAN)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  for (auto i = 0; i < num_keys; ++i) {
+    Key key           = h_keys[i];
+    Value val         = h_keys[i];
+    h_pairs[i].first  = key;
+    h_pairs[i].second = val;
+  }
 
-BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::UNIQUE)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  thrust::device_vector<Key> d_keys(h_keys);
+  thrust::device_vector<bool> d_results(num_keys);
+  thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
 
-BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::UNIQUE)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  for (auto _ : state) {
+    state.PauseTiming();
+    map.insert(d_pairs.begin(), d_pairs.end());
+    state.ResumeTiming();
 
-BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::UNIFORM)
-  ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+    map.erase(d_keys.begin(), d_keys.end());
+  }
 
-BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::UNIFORM)
+  state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
+                          int64_t(state.range(0)));
+}
+
+BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIQUE)
   ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  ->Apply(generate_size_and_occupancy)
+  ->UseManualTime();
 
-BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::GAUSSIAN)
+BENCHMARK_TEMPLATE(BM_static_map_erase_all, int32_t, int32_t, dist_type::UNIQUE)
   ->Unit(benchmark::kMillisecond)
   ->Apply(generate_size_and_occupancy);
 
-BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::GAUSSIAN)
+BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIQUE)
   ->Unit(benchmark::kMillisecond)
-  ->Apply(generate_size_and_occupancy);
+  ->Apply(generate_size_and_occupancy)
+  ->UseManualTime();

examples/static_map/custom_type_example.cu

@@ -91,7 +91,9 @@ int main(void)
   // Construct a map with 100,000 slots using the given empty key/value sentinels. Note the
   // capacity is chosen knowing we will insert 80,000 keys, for an load factor of 80%.
   cuco::static_map<custom_key_type, custom_value_type> map{
-    100'000, empty_key_sentinel, empty_value_sentinel};
+    100'000,
+    cuco::sentinel::empty_key{empty_key_sentinel},
+    cuco::sentinel::empty_value{empty_value_sentinel}};
 
   // Inserts 80,000 pairs into the map by using the custom hasher and custom equality callable
   map.insert(pairs_begin, pairs_begin + num_pairs, custom_hash{}, custom_key_equals{});

examples/static_map/static_map_example.cu

@@ -32,8 +32,11 @@ int main(void)
   // for an load factor of 50%.
   cudaStream_t str;
   cudaStreamCreate(&str);
-  cuco::static_map<int, int> map{
-    100'000, empty_key_sentinel, empty_value_sentinel, cuco::cuda_allocator<char>{}, str};
+  cuco::static_map<int, int> map{100'000,
+                                 cuco::sentinel::empty_key{empty_key_sentinel},
+                                 cuco::sentinel::empty_value{empty_value_sentinel},
+                                 cuco::cuda_allocator<char>{},
+                                 str};
 
   thrust::device_vector<thrust::pair<int, int>> pairs(50'000);
 

include/cuco/detail/dynamic_map.inl

@@ -30,7 +30,10 @@ dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(std::size_t initial_capac
     alloc_{alloc}
 {
   submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
-    initial_capacity, empty_key_sentinel, empty_value_sentinel, alloc));
+    initial_capacity,
+    sentinel::empty_key<Key>{empty_key_sentinel},
+    sentinel::empty_value<Value>{empty_value_sentinel},
+    alloc));
   submap_views_.push_back(submaps_[0]->get_device_view());
   submap_mutable_views_.push_back(submaps_[0]->get_device_mutable_view());
 
@@ -59,7 +62,10 @@ void dynamic_map<Key, Value, Scope, Allocator>::reserve(std::size_t n)
     else {
       submap_capacity = capacity_;
       submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
-        submap_capacity, empty_key_sentinel_, empty_value_sentinel_, alloc_));
+        submap_capacity,
+        sentinel::empty_key<Key>{empty_key_sentinel_},
+        sentinel::empty_value<Value>{empty_value_sentinel_},
+        alloc_));
       submap_views_.push_back(submaps_[submap_idx]->get_device_view());
       submap_mutable_views_.push_back(submaps_[submap_idx]->get_device_mutable_view());
 

include/cuco/detail/error.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-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)