Skip to content

Erase Functionality for dynamic_map #149

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

Merged
merged 38 commits into from
Dec 21, 2022
+894 −98
Merged

Erase Functionality for dynamic_map #149

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
38 commits
Select commit Hold shift + click to select a range
976a746
dynamic map erase, still needs work
niskos99 Mar 9, 2022
eead8b8
minor clarity changes
niskos99 Mar 9, 2022
20ac7a3
erase bug fix
niskos99 Apr 5, 2022
ede50d6
dynamic map erase working, only 4 submaps for now
niskos99 Apr 5, 2022
1d8fbd0
type wrappers added
niskos99 Apr 5, 2022
63dd4eb
prevent implicit type conversion of sentinels during construction
niskos99 Apr 5, 2022
52d83f6
erase benchmark added
niskos99 Apr 5, 2022
0878216
num_successes managed pointer updated
niskos99 Apr 6, 2022
7eac9d1
more efficient block reduce
niskos99 Apr 6, 2022
4d10631
doc changes
niskos99 Apr 6, 2022
b59a16b
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Apr 6, 2022
b00fcba
shared mem atomics to keep track of per-submap erases
niskos99 Apr 7, 2022
c146f9d
doc improvements
niskos99 Apr 30, 2022
faf8224
warning fixes
niskos99 Aug 31, 2022
e4b548e
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Aug 31, 2022
93b7983
removed nvtx file
niskos99 Aug 31, 2022
cd21190
num_successes_ removed
niskos99 Oct 7, 2022
4c1952d
doxygen warning fixes
niskos99 Oct 7, 2022
80f4d14
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Oct 7, 2022
6616889
code cleanup
niskos99 Nov 14, 2022
2df247c
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Nov 14, 2022
af7706d
switched typedef to using
niskos99 Nov 14, 2022
54ae254
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Nov 14, 2022
593fe12
responding to PR comments
niskos99 Nov 18, 2022
7598e47
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Nov 18, 2022
ecd3945
Merge remote-tracking branch 'upstream/dev' into erase
PointKernel Dec 20, 2022
1e6ad99
Add more data types for erase tests
PointKernel Dec 20, 2022
9e324fc
Use public murmurhash
PointKernel Dec 20, 2022
bb0e4e9
Update static map benchmark: fix runtime stall bug, remove redundant …
PointKernel Dec 20, 2022
10fd08a
Update dynamic map benchmark: fix conversion warning, add search_none…
PointKernel Dec 20, 2022
788ad29
Cleanups: get rid of host-side counter vector, remove get_ prefixes a…
PointKernel Dec 20, 2022
c71dd60
Get rid of num_successes getter
PointKernel Dec 20, 2022
ab4ef0c
Fix comments
PointKernel Dec 20, 2022
d72e403
Update tests
PointKernel Dec 20, 2022
9478650
Update include/cuco/detail/dynamic_map_kernels.cuh
PointKernel Dec 20, 2022
82e0f2e
Cleanups: relaxed memory atomic, static_assert instead of runtime exp…
PointKernel Dec 21, 2022
412b7f9
Merge branch 'dynamic_map_erase' of github.com:Nicolas-Iskos/cuCollec…
PointKernel Dec 21, 2022
f5ec677
Reorder header groups + remove unused counter allocator
PointKernel Dec 21, 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
44 changes: 25 additions & 19 deletions include/cuco/detail/dynamic_map.inl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,12 +18,12 @@ 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,
sentinel::empty_key<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel,
Allocator const& alloc)
: empty_key_sentinel_(empty_key_sentinel),
empty_value_sentinel_(empty_value_sentinel),
erased_key_sentinel_(empty_key_sentinel),
: empty_key_sentinel_(empty_key_sentinel.value),
empty_value_sentinel_(empty_value_sentinel.value),
erased_key_sentinel_(empty_key_sentinel.value),
size_(0),
capacity_(initial_capacity),
min_insert_size_(1E4),
Expand All @@ -45,13 +45,13 @@ dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(std::size_t initial_capac

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,
Key erased_key_sentinel,
sentinel::empty_key<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel,
sentinel::erased_key<Key> erased_key_sentinel,
Allocator const& alloc)
: empty_key_sentinel_(empty_key_sentinel),
empty_value_sentinel_(empty_value_sentinel),
erased_key_sentinel_(erased_key_sentinel),
: empty_key_sentinel_(empty_key_sentinel.value),
empty_value_sentinel_(empty_value_sentinel.value),
erased_key_sentinel_(erased_key_sentinel.value),
size_(0),
capacity_(initial_capacity),
min_insert_size_(1E4),
Expand All @@ -66,7 +66,6 @@ dynamic_map<Key, Value, Scope, Allocator>::dynamic_map(std::size_t initial_capac
alloc));
submap_views_.push_back(submaps_[0]->get_device_view());
submap_mutable_views_.push_back(submaps_[0]->get_device_mutable_view());

submap_num_successes_.push_back(submaps_[0]->get_num_successes());

CUCO_CUDA_TRY(cudaMallocManaged(&num_successes_, sizeof(atomic_ctr_type)));
Expand Down Expand Up @@ -94,15 +93,22 @@ void dynamic_map<Key, Value, Scope, Allocator>::reserve(std::size_t n)
// if the submap does not exist yet, create it
else {
submap_capacity = capacity_;
submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
submap_capacity,
sentinel::empty_key<Key>{empty_key_sentinel_},
sentinel::empty_value<Value>{empty_value_sentinel_},
sentinel::erased_key<Key>{erased_key_sentinel_},
alloc_));
if(erased_key_sentinel_ != empty_key_sentinel_) {
submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
submap_capacity,
sentinel::empty_key<Key>{empty_key_sentinel_},
sentinel::empty_value<Value>{empty_value_sentinel_},
sentinel::erased_key<Key>{erased_key_sentinel_},
alloc_));
} else {
submaps_.push_back(std::make_unique<static_map<Key, Value, Scope, Allocator>>(
submap_capacity,
sentinel::empty_key<Key>{empty_key_sentinel_},
sentinel::empty_value<Value>{empty_value_sentinel_},
alloc_));
}
submap_views_.push_back(submaps_[submap_idx]->get_device_view());
submap_mutable_views_.push_back(submaps_[submap_idx]->get_device_mutable_view());

submap_num_successes_.push_back(submaps_[submap_idx]->get_num_successes());

capacity_ *= 2;
Expand Down
6 changes: 3 additions & 3 deletions include/cuco/dynamic_map.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -138,9 +138,9 @@ class dynamic_map {
Allocator const& alloc = Allocator{});

dynamic_map(std::size_t initial_capacity,
Key empty_key_sentinel,
Value empty_value_sentinel,
Key erased_key_sentinel,
sentinel::empty_key<Key> empty_key_sentinel,
sentinel::empty_value<Value> empty_value_sentinel,
sentinel::erased_key<Key> erased_key_sentinel,
Allocator const& alloc = Allocator{});

/**
Expand Down
2 changes: 1 addition & 1 deletion tests/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,7 @@ ConfigureTest(STATIC_MAP_TEST
###################################################################################################
# - dynamic_map tests -----------------------------------------------------------------------------
ConfigureTest(DYNAMIC_MAP_TEST
#dynamic_map/unique_sequence_test.cu
dynamic_map/unique_sequence_test.cu
dynamic_map/erase_test.cu)

###################################################################################################
Expand Down
5 changes: 4 additions & 1 deletion tests/dynamic_map/erase_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,10 @@ TEMPLATE_TEST_CASE_SIG(
using Value = T;

unsigned long num_keys = 1'000'000;
cuco::dynamic_map<Key, Value> map{num_keys * 2, -1, -1, -2};
cuco::dynamic_map<Key, Value> map{num_keys * 2,
cuco::sentinel::empty_key<Key>{-1},
cuco::sentinel::empty_value<Value>{-1},
cuco::sentinel::erased_key<Key>{-2}};

thrust::device_vector<Key> d_keys(num_keys);
thrust::device_vector<Value> d_values(num_keys);
Expand Down
6 changes: 4 additions & 2 deletions tests/dynamic_map/unique_sequence_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,8 +38,10 @@ TEMPLATE_TEST_CASE_SIG("Unique sequence of keys",
(int64_t, int64_t))
{
constexpr std::size_t num_keys{50'000'000};
cuco::dynamic_map<Key, Value> map{
30'000'000, cuco::sentinel::empty_key<Key>{-1}, cuco::sentinel::empty_value<Value>{-1}};

cuco::dynamic_map<Key, Value> map{30'000'000,
cuco::sentinel::empty_key<Key>{-1},
cuco::sentinel::empty_value<Value>{-1}};

thrust::device_vector<Key> d_keys(num_keys);
thrust::device_vector<Value> d_values(num_keys);
Expand Down