Skip to content

Commit 2dbcfc7

Browse files
PointKernelPointKernel
authored
Merge pull request #260 from PointKernel/avoid-sentinel-namespace
Get rid of the use of sentinel namespace
2 parents 125e331 + 08aedcd commit 2dbcfc7

6 files changed

Lines changed: 86 additions & 92 deletions

File tree

‎include/cuco/detail/dynamic_map.inl‎

Lines changed: 12 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -17,12 +17,11 @@
1717
namespace cuco {
1818

1919
template <typename Key, typename Value, cuda::thread_scope Scope, typename Allocator>
20-
dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(
21-
std::size_t initial_capacity,
22-
sentinel::empty_key<Key> empty_key_sentinel,
23-
sentinel::empty_value<Value> empty_value_sentinel,
24-
Allocator const& alloc,
25-
cudaStream_t stream)
20+
dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(std::size_t initial_capacity,
21+
empty_key<Key> empty_key_sentinel,
22+
empty_value<Value> empty_value_sentinel,
23+
Allocator const& alloc,
24+
cudaStream_t stream)
2625
: empty_key_sentinel_(empty_key_sentinel.value),
2726
empty_value_sentinel_(empty_value_sentinel.value),
2827
erased_key_sentinel_(empty_key_sentinel.value),
@@ -34,8 +33,8 @@ dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(
3433
{
3534
submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
3635
initial_capacity,
37-
sentinel::empty_key<Key>{empty_key_sentinel},
38-
sentinel::empty_value<Value>{empty_value_sentinel},
36+
empty_key<Key>{empty_key_sentinel},
37+
empty_value<Value>{empty_value_sentinel},
3938
alloc,
4039
stream));
4140
submap_views_.push_back(submaps_[0]->get_device_view());
@@ -93,16 +92,16 @@ void dynamic_map<Key, Value, Scope, Allocator>::reserve(std::size_t n, cudaStrea
9392
if (erased_key_sentinel_ != empty_key_sentinel_) {
9493
submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
9594
submap_capacity,
96-
sentinel::empty_key<Key>{empty_key_sentinel_},
97-
sentinel::empty_value<Value>{empty_value_sentinel_},
98-
sentinel::erased_key<Key>{erased_key_sentinel_},
95+
empty_key<Key>{empty_key_sentinel_},
96+
empty_value<Value>{empty_value_sentinel_},
97+
erased_key<Key>{erased_key_sentinel_},
9998
alloc_,
10099
stream));
101100
} else {
102101
submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
103102
submap_capacity,
104-
sentinel::empty_key<Key>{empty_key_sentinel_},
105-
sentinel::empty_value<Value>{empty_value_sentinel_},
103+
empty_key<Key>{empty_key_sentinel_},
104+
empty_value<Value>{empty_value_sentinel_},
106105
alloc_,
107106
stream));
108107
}

‎include/cuco/detail/static_map.inl‎

Lines changed: 11 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -28,12 +28,11 @@
2828
namespace cuco {
2929

3030
template <typename Key, typename Value, cuda::thread_scope Scope, typename Allocator>
31-
static_map<Key, Value, Scope, Allocator>::static_map(
32-
std::size_t capacity,
33-
sentinel::empty_key<Key> empty_key_sentinel,
34-
sentinel::empty_value<Value> empty_value_sentinel,
35-
Allocator const& alloc,
36-
cudaStream_t stream)
31+
static_map<Key, Value, Scope, Allocator>::static_map(std::size_t capacity,
32+
empty_key<Key> empty_key_sentinel,
33+
empty_value<Value> empty_value_sentinel,
34+
Allocator const& alloc,
35+
cudaStream_t stream)
3736
: capacity_{std::max(capacity, std::size_t{1})}, // to avoid dereferencing a nullptr (Issue #72)
3837
empty_key_sentinel_{empty_key_sentinel.value},
3938
empty_value_sentinel_{empty_value_sentinel.value},
@@ -53,13 +52,12 @@ static_map<Key, Value, Scope, Allocator>::static_map(
5352
}
5453

5554
template <typename Key, typename Value, cuda::thread_scope Scope, typename Allocator>
56-
static_map<Key, Value, Scope, Allocator>::static_map(
57-
std::size_t capacity,
58-
sentinel::empty_key<Key> empty_key_sentinel,
59-
sentinel::empty_value<Value> empty_value_sentinel,
60-
sentinel::erased_key<Key> erased_key_sentinel,
61-
Allocator const& alloc,
62-
cudaStream_t stream)
55+
static_map<Key, Value, Scope, Allocator>::static_map(std::size_t capacity,
56+
empty_key<Key> empty_key_sentinel,
57+
empty_value<Value> empty_value_sentinel,
58+
erased_key<Key> erased_key_sentinel,
59+
Allocator const& alloc,
60+
cudaStream_t stream)
6361
: capacity_{std::max(capacity, std::size_t{1})}, // to avoid dereferencing a nullptr (Issue #72)
6462
empty_key_sentinel_{empty_key_sentinel.value},
6563
empty_value_sentinel_{empty_value_sentinel.value},

‎include/cuco/detail/static_multimap/static_multimap.inl‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -33,8 +33,8 @@ template <typename Key,
3333
class ProbeSequence>
3434
static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::static_multimap(
3535
std::size_t capacity,
36-
sentinel::empty_key<Key> empty_key_sentinel,
37-
sentinel::empty_value<Value> empty_value_sentinel,
36+
empty_key<Key> empty_key_sentinel,
37+
empty_value<Value> empty_value_sentinel,
3838
cudaStream_t stream,
3939
Allocator const& alloc)
4040
: capacity_{cuco::detail::get_valid_capacity<cg_size(), vector_width(), uses_vector_load()>(

‎include/cuco/dynamic_map.cuh‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -67,8 +67,8 @@ namespace cuco {
6767
* // within the second insert.
6868
*
6969
* dynamic_map<int, int> m{100'000,
70-
* sentinel::empty_key<int>{empty_key_sentinel},
71-
* sentinel::empty_value<int>{empty_value_sentinel}};
70+
* empty_key<int>{empty_key_sentinel},
71+
* empty_value<int>{empty_value_sentinel}};
7272
*
7373
* // Create a sequence of pairs {{0,0}, {1,1}, ... {i,i}}
7474
* thrust::device_vector<thrust::pair<int,int>> pairs_0(50'000);
@@ -170,9 +170,9 @@ class dynamic_map {
170170
* are the same value
171171
*/
172172
dynamic_map(std::size_t initial_capacity,
173-
sentinel::empty_key<Key> empty_key_sentinel,
174-
sentinel::empty_value<Value> empty_value_sentinel,
175-
sentinel::erased_key<Key> erased_key_sentinel,
173+
empty_key<Key> empty_key_sentinel,
174+
empty_value<Value> empty_value_sentinel,
175+
erased_key<Key> erased_key_sentinel,
176176
Allocator const& alloc = Allocator{},
177177
cudaStream_t stream = nullptr);
178178

‎include/cuco/static_map.cuh‎

Lines changed: 42 additions & 44 deletions
Original file line numberDiff line numberDiff line change
@@ -200,8 +200,8 @@ class static_map {
200200
* @param stream Stream used for executing the kernels
201201
*/
202202
static_map(std::size_t capacity,
203-
sentinel::empty_key<Key> empty_key_sentinel,
204-
sentinel::empty_value<Value> empty_value_sentinel,
203+
empty_key<Key> empty_key_sentinel,
204+
empty_value<Value> empty_value_sentinel,
205205
Allocator const& alloc = Allocator{},
206206
cudaStream_t stream = 0);
207207

@@ -220,9 +220,9 @@ class static_map {
220220
* @param stream Stream used for executing the kernels
221221
*/
222222
static_map(std::size_t capacity,
223-
sentinel::empty_key<Key> empty_key_sentinel,
224-
sentinel::empty_value<Value> empty_value_sentinel,
225-
sentinel::erased_key<Key> erased_key_sentinel,
223+
empty_key<Key> empty_key_sentinel,
224+
empty_value<Value> empty_value_sentinel,
225+
erased_key<Key> erased_key_sentinel,
226226
Allocator const& alloc = Allocator{},
227227
cudaStream_t stream = 0);
228228

@@ -435,8 +435,8 @@ class static_map {
435435

436436
__host__ __device__ device_view_base(pair_atomic_type* slots,
437437
std::size_t capacity,
438-
sentinel::empty_key<Key> empty_key_sentinel,
439-
sentinel::empty_value<Value> empty_value_sentinel) noexcept
438+
empty_key<Key> empty_key_sentinel,
439+
empty_value<Value> empty_value_sentinel) noexcept
440440
: slots_{slots},
441441
capacity_{capacity},
442442
empty_key_sentinel_{empty_key_sentinel.value},
@@ -447,9 +447,9 @@ class static_map {
447447

448448
__host__ __device__ device_view_base(pair_atomic_type* slots,
449449
std::size_t capacity,
450-
sentinel::empty_key<Key> empty_key_sentinel,
451-
sentinel::empty_value<Value> empty_value_sentinel,
452-
sentinel::erased_key<Key> erased_key_sentinel) noexcept
450+
empty_key<Key> empty_key_sentinel,
451+
empty_value<Value> empty_value_sentinel,
452+
erased_key<Key> erased_key_sentinel) noexcept
453453
: slots_{slots},
454454
capacity_{capacity},
455455
empty_key_sentinel_{empty_key_sentinel.value},
@@ -768,11 +768,10 @@ class static_map {
768768
* @param empty_value_sentinel The reserved value for mapped values to
769769
* represent empty slots
770770
*/
771-
__host__ __device__
772-
device_mutable_view(pair_atomic_type* slots,
773-
std::size_t capacity,
774-
sentinel::empty_key<Key> empty_key_sentinel,
775-
sentinel::empty_value<Value> empty_value_sentinel) noexcept
771+
__host__ __device__ device_mutable_view(pair_atomic_type* slots,
772+
std::size_t capacity,
773+
empty_key<Key> empty_key_sentinel,
774+
empty_value<Value> empty_value_sentinel) noexcept
776775
: device_view_base{slots, capacity, empty_key_sentinel, empty_value_sentinel}
777776
{
778777
}
@@ -789,9 +788,9 @@ class static_map {
789788
*/
790789
__host__ __device__ device_mutable_view(pair_atomic_type* slots,
791790
std::size_t capacity,
792-
sentinel::empty_key<Key> empty_key_sentinel,
793-
sentinel::empty_value<Value> empty_value_sentinel,
794-
sentinel::erased_key<Key> erased_key_sentinel) noexcept
791+
empty_key<Key> empty_key_sentinel,
792+
empty_value<Value> empty_value_sentinel,
793+
erased_key<Key> erased_key_sentinel) noexcept
795794
: device_view_base{
796795
slots, capacity, empty_key_sentinel, empty_value_sentinel, erased_key_sentinel}
797796
{
@@ -878,16 +877,16 @@ class static_map {
878877
CG const& g,
879878
pair_atomic_type* slots,
880879
std::size_t capacity,
881-
sentinel::empty_key<Key> empty_key_sentinel,
882-
sentinel::empty_value<Value> empty_value_sentinel) noexcept
880+
empty_key<Key> empty_key_sentinel,
881+
empty_value<Value> empty_value_sentinel) noexcept
883882
{
884883
device_view_base::initialize_slots(
885884
g, slots, capacity, empty_key_sentinel.value, empty_value_sentinel.value);
886885
return device_mutable_view{slots,
887886
capacity,
888887
empty_key_sentinel,
889888
empty_value_sentinel,
890-
sentinel::erased_key<Key>{empty_key_sentinel.value}};
889+
erased_key<Key>{empty_key_sentinel.value}};
891890
}
892891

893892
/**
@@ -909,9 +908,9 @@ class static_map {
909908
CG const& g,
910909
pair_atomic_type* slots,
911910
std::size_t capacity,
912-
sentinel::empty_key<Key> empty_key_sentinel,
913-
sentinel::empty_value<Value> empty_value_sentinel,
914-
sentinel::erased_key<Key> erased_key_sentinel) noexcept
911+
empty_key<Key> empty_key_sentinel,
912+
empty_value<Value> empty_value_sentinel,
913+
erased_key<Key> erased_key_sentinel) noexcept
915914
{
916915
device_view_base::initialize_slots(
917916
g, slots, capacity, empty_key_sentinel, empty_value_sentinel);
@@ -1070,8 +1069,8 @@ class static_map {
10701069
*/
10711070
__host__ __device__ device_view(pair_atomic_type* slots,
10721071
std::size_t capacity,
1073-
sentinel::empty_key<Key> empty_key_sentinel,
1074-
sentinel::empty_value<Value> empty_value_sentinel) noexcept
1072+
empty_key<Key> empty_key_sentinel,
1073+
empty_value<Value> empty_value_sentinel) noexcept
10751074
: device_view_base{slots, capacity, empty_key_sentinel, empty_value_sentinel}
10761075
{
10771076
}
@@ -1088,9 +1087,9 @@ class static_map {
10881087
*/
10891088
__host__ __device__ device_view(pair_atomic_type* slots,
10901089
std::size_t capacity,
1091-
sentinel::empty_key<Key> empty_key_sentinel,
1092-
sentinel::empty_value<Value> empty_value_sentinel,
1093-
sentinel::erased_key<Key> erased_key_sentinel) noexcept
1090+
empty_key<Key> empty_key_sentinel,
1091+
empty_value<Value> empty_value_sentinel,
1092+
erased_key<Key> erased_key_sentinel) noexcept
10941093
: device_view_base{
10951094
slots, capacity, empty_key_sentinel, empty_value_sentinel, erased_key_sentinel}
10961095
{
@@ -1104,9 +1103,9 @@ class static_map {
11041103
__host__ __device__ explicit device_view(device_mutable_view mutable_map)
11051104
: device_view_base{mutable_map.get_slots(),
11061105
mutable_map.get_capacity(),
1107-
sentinel::empty_key<Key>{mutable_map.get_empty_key_sentinel()},
1108-
sentinel::empty_value<Value>{mutable_map.get_empty_value_sentinel()},
1109-
sentinel::erased_key<Key>{mutable_map.get_erased_key_sentinel()}}
1106+
empty_key<Key>{mutable_map.get_empty_key_sentinel()},
1107+
empty_value<Value>{mutable_map.get_empty_value_sentinel()},
1108+
erased_key<Key>{mutable_map.get_erased_key_sentinel()}}
11101109
{
11111110
}
11121111

@@ -1175,12 +1174,11 @@ class static_map {
11751174
g.sync();
11761175
#endif
11771176

1178-
return device_view(
1179-
memory_to_use,
1180-
source_device_view.get_capacity(),
1181-
sentinel::empty_key<Key>{source_device_view.get_empty_key_sentinel()},
1182-
sentinel::empty_value<Value>{source_device_view.get_empty_value_sentinel()},
1183-
sentinel::erased_key<Key>{source_device_view.get_erased_key_sentinel()});
1177+
return device_view(memory_to_use,
1178+
source_device_view.get_capacity(),
1179+
empty_key<Key>{source_device_view.get_empty_key_sentinel()},
1180+
empty_value<Value>{source_device_view.get_empty_value_sentinel()},
1181+
erased_key<Key>{source_device_view.get_erased_key_sentinel()});
11841182
}
11851183

11861184
/**
@@ -1395,9 +1393,9 @@ class static_map {
13951393
{
13961394
return device_view(slots_,
13971395
capacity_,
1398-
sentinel::empty_key<Key>{empty_key_sentinel_},
1399-
sentinel::empty_value<Value>{empty_value_sentinel_},
1400-
sentinel::erased_key<Key>{erased_key_sentinel_});
1396+
empty_key<Key>{empty_key_sentinel_},
1397+
empty_value<Value>{empty_value_sentinel_},
1398+
erased_key<Key>{erased_key_sentinel_});
14011399
}
14021400

14031401
/**
@@ -1409,9 +1407,9 @@ class static_map {
14091407
{
14101408
return device_mutable_view(slots_,
14111409
capacity_,
1412-
sentinel::empty_key<Key>{empty_key_sentinel_},
1413-
sentinel::empty_value<Value>{empty_value_sentinel_},
1414-
sentinel::erased_key<Key>{erased_key_sentinel_});
1410+
empty_key<Key>{empty_key_sentinel_},
1411+
empty_value<Value>{empty_value_sentinel_},
1412+
erased_key<Key>{erased_key_sentinel_});
14151413
}
14161414

14171415
private:

‎include/cuco/static_multimap.cuh‎

Lines changed: 14 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -225,8 +225,8 @@ class static_multimap {
225225
* @param alloc Allocator used for allocating device storage
226226
*/
227227
static_multimap(std::size_t capacity,
228-
sentinel::empty_key<Key> empty_key_sentinel,
229-
sentinel::empty_value<Value> empty_value_sentinel,
228+
empty_key<Key> empty_key_sentinel,
229+
empty_value<Value> empty_value_sentinel,
230230
cudaStream_t stream = 0,
231231
Allocator const& alloc = Allocator{});
232232

@@ -611,8 +611,8 @@ class static_multimap {
611611

612612
__host__ __device__ device_view_base(pair_atomic_type* slots,
613613
std::size_t capacity,
614-
sentinel::empty_key<Key> empty_key_sentinel,
615-
sentinel::empty_value<Value> empty_value_sentinel) noexcept
614+
empty_key<Key> empty_key_sentinel,
615+
empty_value<Value> empty_value_sentinel) noexcept
616616
: impl_{slots, capacity, empty_key_sentinel.value, empty_value_sentinel.value}
617617
{
618618
}
@@ -714,11 +714,10 @@ class static_multimap {
714714
* @param empty_value_sentinel The reserved value for mapped values to
715715
* represent empty slots
716716
*/
717-
__host__ __device__
718-
device_mutable_view(pair_atomic_type* slots,
719-
std::size_t capacity,
720-
sentinel::empty_key<Key> empty_key_sentinel,
721-
sentinel::empty_value<Value> empty_value_sentinel) noexcept
717+
__host__ __device__ device_mutable_view(pair_atomic_type* slots,
718+
std::size_t capacity,
719+
empty_key<Key> empty_key_sentinel,
720+
empty_value<Value> empty_value_sentinel) noexcept
722721
: view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel}
723722
{
724723
}
@@ -770,8 +769,8 @@ class static_multimap {
770769
*/
771770
__host__ __device__ device_view(pair_atomic_type* slots,
772771
std::size_t capacity,
773-
sentinel::empty_key<Key> empty_key_sentinel,
774-
sentinel::empty_value<Value> empty_value_sentinel) noexcept
772+
empty_key<Key> empty_key_sentinel,
773+
empty_value<Value> empty_value_sentinel) noexcept
775774
: view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel}
776775
{
777776
}
@@ -1325,8 +1324,8 @@ class static_multimap {
13251324
{
13261325
return device_view(slots_.get(),
13271326
capacity_,
1328-
sentinel::empty_key<Key>{empty_key_sentinel_},
1329-
sentinel::empty_value<Value>{empty_value_sentinel_});
1327+
empty_key<Key>{empty_key_sentinel_},
1328+
empty_value<Value>{empty_value_sentinel_});
13301329
}
13311330

13321331
/**
@@ -1339,8 +1338,8 @@ class static_multimap {
13391338
{
13401339
return device_mutable_view(slots_.get(),
13411340
capacity_,
1342-
sentinel::empty_key<Key>{empty_key_sentinel_},
1343-
sentinel::empty_value<Value>{empty_value_sentinel_});
1341+
empty_key<Key>{empty_key_sentinel_},
1342+
empty_value<Value>{empty_value_sentinel_});
13441343
}
13451344

13461345
private:

0 commit comments

Comments
 (0)