diff --git a/include/cuco/detail/static_map/kernels.cuh b/include/cuco/detail/static_map/kernels.cuh index c05d0b28b..bf2aced70 100644 --- a/include/cuco/detail/static_map/kernels.cuh +++ b/include/cuco/detail/static_map/kernels.cuh @@ -202,8 +202,8 @@ CUCO_KERNEL __launch_bounds__(BlockSize) void insert_or_apply_shmem( auto shared_map = SharedMapRefType{cuco::empty_key{ref.empty_key_sentinel()}, cuco::empty_value{ref.empty_value_sentinel()}, - {}, - {}, + ref.key_eq(), + ref.probing_scheme(), {}, storage}; auto shared_map_ref = std::move(shared_map).with(cuco::op::insert_or_apply); diff --git a/include/cuco/detail/static_map/static_map.inl b/include/cuco/detail/static_map/static_map.inl index e575114de..08acdacaf 100644 --- a/include/cuco/detail/static_map/static_map.inl +++ b/include/cuco/detail/static_map/static_map.inl @@ -347,19 +347,11 @@ template -template +template void static_map:: insert_or_apply_async( InputIt first, InputIt last, Init init, Op op, cuda::stream_ref stream) noexcept { - using shared_map_type = cuco::static_map>; auto constexpr has_init = true; static_map_ns::detail::dispatch_insert_or_apply( first, last, init, op, ref(op::insert_or_apply), stream); diff --git a/include/cuco/detail/storage/aow_storage.inl b/include/cuco/detail/storage/aow_storage.inl index dfbb90327..5d976e658 100644 --- a/include/cuco/detail/storage/aow_storage.inl +++ b/include/cuco/detail/storage/aow_storage.inl @@ -74,6 +74,8 @@ template void aow_storage::initialize_async( value_type key, cuda::stream_ref stream) noexcept { + if (this->num_windows() == 0) { return; } + auto constexpr cg_size = 1; auto constexpr stride = 4; auto const grid_size = cuco::detail::grid_size(this->num_windows(), cg_size, stride); diff --git a/include/cuco/static_map.cuh b/include/cuco/static_map.cuh index 01a39ad5d..09e980261 100644 --- a/include/cuco/static_map.cuh +++ b/include/cuco/static_map.cuh @@ -37,6 +37,7 @@ #include #include +#include #include namespace cuco { @@ -564,7 +565,10 @@ class static_map { * @param op Callable object to perform apply operation. * @param stream CUDA stream used for insert */ - template + template >> void insert_or_apply_async( InputIt first, InputIt last, Init init, Op op, cuda::stream_ref stream = {}) noexcept; diff --git a/tests/utility/storage_test.cu b/tests/utility/storage_test.cu index 8cbd8c08d..8fb923fb8 100644 --- a/tests/utility/storage_test.cu +++ b/tests/utility/storage_test.cu @@ -37,6 +37,15 @@ TEMPLATE_TEST_CASE_SIG("Storage tests", using allocator_type = cuco::cuda_allocator; auto allocator = allocator_type{}; + SECTION("Initialize empty storage is allowed.") + { + auto s = cuco:: + aow_storage, window_size, cuco::extent, allocator_type>{ + cuco::extent{0}, allocator}; + + s.initialize(cuco::pair{1, 1}); + } + SECTION("Allocate array of pairs with AoS storage.") { auto s = cuco::