Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Apply sentinel namespace to other map structures #163

Merged
merged 3 commits into from
May 27, 2022
+28 −24
Merged

Apply sentinel namespace to other map structures #163

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
9d2e198
Apply sentinel namespace to dynamic_map
m3g4d1v3r May 26, 2022
360d8af
Apply sentinel namespace to static_multimap
m3g4d1v3r May 26, 2022
4dc7a77
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] May 26, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 7 additions & 6 deletions include/cuco/detail/dynamic_map.inl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,12 +17,13 @@
namespace cuco {

template <typename Key, typename Value, cuda::thread_scope Scope, typename Allocator>
dynamic_map<Key, Value, Scope, Allocator>::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<Key, Value, Scope, Allocator>::dynamic_map(
std::size_t initial_capacity,
sentinel::empty_key<Key> empty_key_sentinel,
sentinel::empty_value<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),
Expand Down
10 changes: 5 additions & 5 deletions include/cuco/detail/static_multimap/static_multimap.inl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,14 +33,14 @@ template <typename Key,
class ProbeSequence>
static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::static_multimap(
std::size_t capacity,
Key empty_key_sentinel,
Value empty_value_sentinel,
sentinel::empty_key<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel,
cudaStream_t stream,
Allocator const& alloc)
: capacity_{cuco::detail::get_valid_capacity<cg_size(), vector_width(), uses_vector_load()>(
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_},
Expand All @@ -53,7 +53,7 @@ static_multimap<Key, Value, Scope, Allocator, ProbeSequence>::static_multimap(
auto const grid_size = (get_capacity() + stride * block_size - 1) / (stride * block_size);

detail::initialize<atomic_key_type, atomic_mapped_type><<<grid_size, block_size, 0, stream>>>(
slots_.get(), empty_key_sentinel, empty_value_sentinel, get_capacity());
slots_.get(), empty_key_sentinel_, empty_value_sentinel_, get_capacity());
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Good catch

}

template <typename Key,
Expand Down
5 changes: 3 additions & 2 deletions include/cuco/dynamic_map.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,7 @@

#include <cuco/detail/dynamic_map_kernels.cuh>
#include <cuco/detail/error.hpp>
#include <cuco/sentinel.cuh>
#include <cuco/static_map.cuh>

#include <thrust/device_vector.h>
Expand Down Expand Up @@ -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<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel,
Allocator const& alloc = Allocator{});

/**
Expand Down
24 changes: 13 additions & 11 deletions include/cuco/static_multimap.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,7 @@
#include <cuco/detail/prime.hpp>
#include <cuco/detail/static_multimap/kernels.cuh>
#include <cuco/probe_sequences.cuh>
#include <cuco/sentinel.cuh>
#include <cuco/traits.hpp>

#include <thrust/functional.h>
Expand Down Expand Up @@ -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<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel,
cudaStream_t stream = 0,
Allocator const& alloc = Allocator{});

Expand Down Expand Up @@ -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<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel) noexcept
: impl_{slots, capacity, empty_key_sentinel.value, empty_value_sentinel.value}
{
}

Expand Down Expand Up @@ -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<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel) noexcept
: view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel}
{
}
Expand Down Expand Up @@ -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<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel) noexcept
: view_base_type{slots, capacity, empty_key_sentinel, empty_value_sentinel}
{
}
Expand Down