diff --git a/include/cuco/detail/dynamic_map.inl b/include/cuco/detail/dynamic_map.inl index 80020232d..27bd3d178 100644 --- a/include/cuco/detail/dynamic_map.inl +++ b/include/cuco/detail/dynamic_map.inl @@ -17,12 +17,13 @@ namespace cuco { template -dynamic_map::dynamic_map(std::size_t initial_capacity, - Key empty_key_sentinel, - Value empty_value_sentinel, - Allocator const& alloc) - : empty_key_sentinel_(empty_key_sentinel), - empty_value_sentinel_(empty_value_sentinel), +dynamic_map::dynamic_map( + std::size_t initial_capacity, + sentinel::empty_key empty_key_sentinel, + sentinel::empty_value empty_value_sentinel, + Allocator const& alloc) + : empty_key_sentinel_(empty_key_sentinel.value), + empty_value_sentinel_(empty_value_sentinel.value), size_(0), capacity_(initial_capacity), min_insert_size_(1E4), diff --git a/include/cuco/detail/static_multimap/static_multimap.inl b/include/cuco/detail/static_multimap/static_multimap.inl index 450c5d871..a3920b181 100644 --- a/include/cuco/detail/static_multimap/static_multimap.inl +++ b/include/cuco/detail/static_multimap/static_multimap.inl @@ -33,14 +33,14 @@ template static_multimap::static_multimap( std::size_t capacity, - Key empty_key_sentinel, - Value empty_value_sentinel, + sentinel::empty_key empty_key_sentinel, + sentinel::empty_value empty_value_sentinel, cudaStream_t stream, Allocator const& alloc) : capacity_{cuco::detail::get_valid_capacity( capacity)}, - empty_key_sentinel_{empty_key_sentinel}, - empty_value_sentinel_{empty_value_sentinel}, + empty_key_sentinel_{empty_key_sentinel.value}, + empty_value_sentinel_{empty_value_sentinel.value}, counter_allocator_{alloc}, slot_allocator_{alloc}, delete_counter_{counter_allocator_}, @@ -53,7 +53,7 @@ static_multimap::static_multimap( auto const grid_size = (get_capacity() + stride * block_size - 1) / (stride * block_size); detail::initialize<<>>( - slots_.get(), empty_key_sentinel, empty_value_sentinel, get_capacity()); + slots_.get(), empty_key_sentinel_, empty_value_sentinel_, get_capacity()); } template #include +#include #include #include @@ -131,8 +132,8 @@ class dynamic_map { * @param alloc Allocator used to allocate submap device storage */ dynamic_map(std::size_t initial_capacity, - Key empty_key_sentinel, - Value empty_value_sentinel, + sentinel::empty_key empty_key_sentinel, + sentinel::empty_value empty_value_sentinel, Allocator const& alloc = Allocator{}); /** diff --git a/include/cuco/static_multimap.cuh b/include/cuco/static_multimap.cuh index 5f7ef35e6..bfe0f8fe7 100644 --- a/include/cuco/static_multimap.cuh +++ b/include/cuco/static_multimap.cuh @@ -21,6 +21,7 @@ #include #include #include +#include #include #include @@ -223,8 +224,8 @@ class static_multimap { * @param alloc Allocator used for allocating device storage */ static_multimap(std::size_t capacity, - Key empty_key_sentinel, - Value empty_value_sentinel, + sentinel::empty_key empty_key_sentinel, + sentinel::empty_value empty_value_sentinel, cudaStream_t stream = 0, Allocator const& alloc = Allocator{}); @@ -576,9 +577,9 @@ class static_multimap { __host__ __device__ device_view_base(pair_atomic_type* slots, std::size_t capacity, - Key empty_key_sentinel, - Value empty_value_sentinel) noexcept - : impl_{slots, capacity, empty_key_sentinel, empty_value_sentinel} + sentinel::empty_key empty_key_sentinel, + sentinel::empty_value empty_value_sentinel) noexcept + : impl_{slots, capacity, empty_key_sentinel.value, empty_value_sentinel.value} { } @@ -675,10 +676,11 @@ class static_multimap { * @param empty_value_sentinel The reserved value for mapped values to * represent empty slots */ - __host__ __device__ device_mutable_view(pair_atomic_type* slots, - std::size_t capacity, - Key empty_key_sentinel, - Value empty_value_sentinel) noexcept + __host__ __device__ + device_mutable_view(pair_atomic_type* slots, + std::size_t capacity, + sentinel::empty_key empty_key_sentinel, + sentinel::empty_value empty_value_sentinel) noexcept : view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel} { } @@ -728,8 +730,8 @@ class static_multimap { */ __host__ __device__ device_view(pair_atomic_type* slots, std::size_t capacity, - Key empty_key_sentinel, - Value empty_value_sentinel) noexcept + sentinel::empty_key empty_key_sentinel, + sentinel::empty_value empty_value_sentinel) noexcept : view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel} { }