diff --git a/tests/static_map/shared_memory_test.cu b/tests/static_map/shared_memory_test.cu index 046f809ec..6f6700c69 100644 --- a/tests/static_map/shared_memory_test.cu +++ b/tests/static_map/shared_memory_test.cu @@ -173,7 +173,7 @@ template __global__ void shared_memory_hash_table_kernel(bool* key_found) { using slot_type = cuco::pair; - __shared__ cuco::experimental::window windows[NumWindows]; + __shared__ cuco::experimental::window map[NumWindows]; using extent_type = cuco::experimental::extent; using storage_ref_type = cuco::experimental::aow_storage_ref; @@ -188,28 +188,25 @@ __global__ void shared_memory_hash_table_kernel(bool* key_found) cuco::empty_value{-1}, {}, {}, - storage_ref_type{extent_type{}, windows}}; - /* + storage_ref_type{extent_type{}, map}}; + namespace cg = cooperative_groups; + auto const block = cg::this_thread_block(); + // map.initialize(block); - namespace cg = cooperative_groups; - auto map = map_type::make_from_uninitialized_slots( - cg::this_thread_block(), &slots[0], N, cuco::empty_key{-1}, cuco::empty_value{-1}); - - auto g = cg::this_thread_block(); std::size_t index = threadIdx.x + blockIdx.x * blockDim.x; - int rank = g.thread_rank(); + auto const rank = block.thread_rank(); // insert {thread_rank, thread_rank} for each thread in thread-block - map.insert(cuco::pair(rank, rank)); - g.sync(); + auto insert_ref = std::move(raw_ref).with(cuco::experimental::op::insert); + insert_ref.insert(slot_type{rank, rank}); + block.sync(); - auto find_map = find_map_type(map); - auto retrieved_pair = find_map.find(rank); - if (retrieved_pair != find_map.end() && retrieved_pair->second == rank) { + auto find_ref = std::move(insert_ref).with(cuco::experimental::op::find); + auto const retrieved_pair = find_ref.find(rank); + if (retrieved_pair != find_ref.end() && retrieved_pair->second == rank) { key_found[index] = true; } - */ } TEMPLATE_TEST_CASE("Shared memory slots.", "", int32_t)