Skip to content

Commit 31c30d1

Browse files
PointKernelPointKernel
authored
Merge pull request #142 from Nicolas-Iskos/static_map_erase
Erase Functionality for static_map
2 parents ebaba1a + 7b841a8 commit 31c30d1

17 files changed

+695
-113
lines changed

‎benchmarks/hash_table/static_map_bench.cu‎

Lines changed: 58 additions & 38 deletions
Original file line numberDiff line numberDiff line change
@@ -87,16 +87,24 @@ static void BM_static_map_insert(::benchmark::State& state)
8787
}
8888

8989
thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
90+
thrust::device_vector<Key> d_keys(h_keys);
9091

9192
for (auto _ : state) {
92-
state.ResumeTiming();
93-
state.PauseTiming();
94-
map_type map{size, -1, -1};
95-
state.ResumeTiming();
93+
map_type map{size, cuco::sentinel::empty_key<Key>{-1}, cuco::sentinel::empty_value<Value>{-1}};
9694

95+
cudaEvent_t start, stop;
96+
cudaEventCreate(&start);
97+
cudaEventCreate(&stop);
98+
99+
cudaEventRecord(start);
97100
map.insert(d_pairs.begin(), d_pairs.end());
101+
cudaEventRecord(stop);
102+
cudaEventSynchronize(stop);
98103

99-
state.PauseTiming();
104+
float ms;
105+
cudaEventElapsedTime(&ms, start, stop);
106+
107+
state.SetIterationTime(ms / 1000);
100108
}
101109

102110
state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
@@ -112,7 +120,7 @@ static void BM_static_map_search_all(::benchmark::State& state)
112120
float occupancy = state.range(1) / float{100};
113121
std::size_t size = num_keys / occupancy;
114122

115-
map_type map{size, -1, -1};
123+
map_type map{size, cuco::sentinel::empty_key<Key>{-1}, cuco::sentinel::empty_value<Value>{-1}};
116124
auto view = map.get_device_mutable_view();
117125

118126
std::vector<Key> h_keys(num_keys);
@@ -143,50 +151,62 @@ static void BM_static_map_search_all(::benchmark::State& state)
143151
int64_t(state.range(0)));
144152
}
145153

146-
BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIQUE)
147-
->Unit(benchmark::kMillisecond)
148-
->Apply(generate_size_and_occupancy);
154+
template <typename Key, typename Value, dist_type Dist>
155+
static void BM_static_map_erase_all(::benchmark::State& state)
156+
{
157+
using map_type = cuco::static_map<Key, Value>;
149158

150-
BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::UNIQUE)
151-
->Unit(benchmark::kMillisecond)
152-
->Apply(generate_size_and_occupancy);
159+
std::size_t num_keys = state.range(0);
160+
float occupancy = state.range(1) / float{100};
161+
std::size_t size = num_keys / occupancy;
153162

154-
BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIFORM)
155-
->Unit(benchmark::kMillisecond)
156-
->Apply(generate_size_and_occupancy);
163+
// static map with erase support
164+
map_type map{size,
165+
cuco::sentinel::empty_key<Key>{-1},
166+
cuco::sentinel::empty_value<Value>{-1},
167+
cuco::sentinel::erased_key<Key>{-2}};
168+
auto view = map.get_device_mutable_view();
157169

158-
BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::UNIFORM)
159-
->Unit(benchmark::kMillisecond)
160-
->Apply(generate_size_and_occupancy);
170+
std::vector<Key> h_keys(num_keys);
171+
std::vector<Value> h_values(num_keys);
172+
std::vector<cuco::pair_type<Key, Value>> h_pairs(num_keys);
173+
std::vector<Value> h_results(num_keys);
161174

162-
BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::GAUSSIAN)
163-
->Unit(benchmark::kMillisecond)
164-
->Apply(generate_size_and_occupancy);
175+
generate_keys<Dist, Key>(h_keys.begin(), h_keys.end());
165176

166-
BENCHMARK_TEMPLATE(BM_static_map_search_all, int32_t, int32_t, dist_type::GAUSSIAN)
167-
->Unit(benchmark::kMillisecond)
168-
->Apply(generate_size_and_occupancy);
177+
for (auto i = 0; i < num_keys; ++i) {
178+
Key key = h_keys[i];
179+
Value val = h_keys[i];
180+
h_pairs[i].first = key;
181+
h_pairs[i].second = val;
182+
}
169183

170-
BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::UNIQUE)
171-
->Unit(benchmark::kMillisecond)
172-
->Apply(generate_size_and_occupancy);
184+
thrust::device_vector<Key> d_keys(h_keys);
185+
thrust::device_vector<bool> d_results(num_keys);
186+
thrust::device_vector<cuco::pair_type<Key, Value>> d_pairs(h_pairs);
173187

174-
BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::UNIQUE)
175-
->Unit(benchmark::kMillisecond)
176-
->Apply(generate_size_and_occupancy);
188+
for (auto _ : state) {
189+
state.PauseTiming();
190+
map.insert(d_pairs.begin(), d_pairs.end());
191+
state.ResumeTiming();
177192

178-
BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::UNIFORM)
179-
->Unit(benchmark::kMillisecond)
180-
->Apply(generate_size_and_occupancy);
193+
map.erase(d_keys.begin(), d_keys.end());
194+
}
181195

182-
BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::UNIFORM)
196+
state.SetBytesProcessed((sizeof(Key) + sizeof(Value)) * int64_t(state.iterations()) *
197+
int64_t(state.range(0)));
198+
}
199+
200+
BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIQUE)
183201
->Unit(benchmark::kMillisecond)
184-
->Apply(generate_size_and_occupancy);
202+
->Apply(generate_size_and_occupancy)
203+
->UseManualTime();
185204

186-
BENCHMARK_TEMPLATE(BM_static_map_insert, int64_t, int64_t, dist_type::GAUSSIAN)
205+
BENCHMARK_TEMPLATE(BM_static_map_erase_all, int32_t, int32_t, dist_type::UNIQUE)
187206
->Unit(benchmark::kMillisecond)
188207
->Apply(generate_size_and_occupancy);
189208

190-
BENCHMARK_TEMPLATE(BM_static_map_search_all, int64_t, int64_t, dist_type::GAUSSIAN)
209+
BENCHMARK_TEMPLATE(BM_static_map_insert, int32_t, int32_t, dist_type::UNIQUE)
191210
->Unit(benchmark::kMillisecond)
192-
->Apply(generate_size_and_occupancy);
211+
->Apply(generate_size_and_occupancy)
212+
->UseManualTime();

‎examples/static_map/custom_type_example.cu‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -91,7 +91,9 @@ int main(void)
9191
// Construct a map with 100,000 slots using the given empty key/value sentinels. Note the
9292
// capacity is chosen knowing we will insert 80,000 keys, for an load factor of 80%.
9393
cuco::static_map<custom_key_type, custom_value_type> map{
94-
100'000, empty_key_sentinel, empty_value_sentinel};
94+
100'000,
95+
cuco::sentinel::empty_key{empty_key_sentinel},
96+
cuco::sentinel::empty_value{empty_value_sentinel}};
9597

9698
// Inserts 80,000 pairs into the map by using the custom hasher and custom equality callable
9799
map.insert(pairs_begin, pairs_begin + num_pairs, custom_hash{}, custom_key_equals{});

‎examples/static_map/static_map_example.cu‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -32,8 +32,11 @@ int main(void)
3232
// for an load factor of 50%.
3333
cudaStream_t str;
3434
cudaStreamCreate(&str);
35-
cuco::static_map<int, int> map{
36-
100'000, empty_key_sentinel, empty_value_sentinel, cuco::cuda_allocator<char>{}, str};
35+
cuco::static_map<int, int> map{100'000,
36+
cuco::sentinel::empty_key{empty_key_sentinel},
37+
cuco::sentinel::empty_value{empty_value_sentinel},
38+
cuco::cuda_allocator<char>{},
39+
str};
3740

3841
thrust::device_vector<thrust::pair<int, int>> pairs(50'000);
3942

‎include/cuco/detail/dynamic_map.inl‎

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,10 @@ dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(std::size_t initial_capac
3030
alloc_{alloc}
3131
{
3232
submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
33-
initial_capacity, empty_key_sentinel, empty_value_sentinel, alloc));
33+
initial_capacity,
34+
sentinel::empty_key<Key>{empty_key_sentinel},
35+
sentinel::empty_value<Value>{empty_value_sentinel},
36+
alloc));
3437
submap_views_.push_back(submaps_[0]->get_device_view());
3538
submap_mutable_views_.push_back(submaps_[0]->get_device_mutable_view());
3639

@@ -59,7 +62,10 @@ void dynamic_map<Key, Value, Scope, Allocator>::reserve(std::size_t n)
5962
else {
6063
submap_capacity = capacity_;
6164
submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
62-
submap_capacity, empty_key_sentinel_, empty_value_sentinel_, alloc_));
65+
submap_capacity,
66+
sentinel::empty_key<Key>{empty_key_sentinel_},
67+
sentinel::empty_value<Value>{empty_value_sentinel_},
68+
alloc_));
6369
submap_views_.push_back(submaps_[submap_idx]->get_device_view());
6470
submap_mutable_views_.push_back(submaps_[submap_idx]->get_device_mutable_view());
6571

‎include/cuco/detail/error.hpp‎

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2020, NVIDIA CORPORATION.
2+
* Copyright (c) 2020-2022, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.
@@ -80,3 +80,23 @@ struct cuda_error : public std::runtime_error {
8080
cudaError_t const status = (expr); \
8181
assert(cudaSuccess == status); \
8282
} while (0)
83+
84+
/**
85+
* @brief Macro for checking runtime conditions that throws an exception when
86+
* a condition is violated.
87+
*
88+
* Example usage:
89+
*
90+
* @code
91+
* CUCO_RUNTIME_EXPECTS(key == value, "Key value mismatch");
92+
* @endcode
93+
*
94+
* @param[in] cond Expression that evaluates to true or false
95+
* @param[in] reason String literal description of the reason that cond is
96+
* expected to be true
97+
* @throw std::runtime_error if the condition evaluates to false.
98+
*/
99+
#define CUCO_RUNTIME_EXPECTS(cond, reason) \
100+
(!!(cond)) ? static_cast<void>(0) \
101+
: throw std::runtime_error("cuco failure at: " __FILE__ \
102+
":" CUCO_STRINGIFY(__LINE__) ": " reason)

0 commit comments

Comments
 (0)