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();
 
+    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()) *
@@ -143,50 +151,59 @@ 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, -1, -1, -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();

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)