@@ -49,6 +49,8 @@ __global__ void shared_memory_test_kernel(Ref* sets,
4949 auto insert_ref = sets[set_id].make_copy (g, sm_buffer, cuco::thread_scope_block);
5050 auto find_ref = insert_ref.rebind_operators (cuco::op::find);
5151
52+ __syncthreads ();
53+
5254 for (int i = g.thread_rank (); i < number_of_elements; i += g.size ()) {
5355 auto found_it = find_ref.find (insterted_keys[offset + i]);
5456
@@ -80,10 +82,11 @@ TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
8082)
8183{
8284 // For uint8_t: sentinel = 0xFF (255), so usable key range is 0..254.
85+ // For uint16_t: sentinel = 0xFFFF (65535), so usable key range is 0..65534.
8386 // thrust::sequence over number_of_sets*elements_in_set keys must not wrap.
84- // Use a smaller set count and element count for 1-byte types.
85- constexpr std::size_t number_of_sets = (sizeof (Key) == 1 ) ? 2 : 1000 ;
86- constexpr std::size_t elements_in_set = (sizeof (Key) == 1 ) ? 100 : 500 ;
87+ // Use smaller set count and element count for smaller types.
88+ constexpr std::size_t number_of_sets = (sizeof (Key) <= 2 ) ? ( sizeof (Key) == 1 ? 2 : 100 ) : 1000 ;
89+ constexpr std::size_t elements_in_set = (sizeof (Key) <= 2 ) ? ( sizeof (Key) == 1 ? 100 : 600 ) : 500 ;
8790 constexpr std::size_t set_capacity = 2 * elements_in_set;
8891
8992 using extent_type = cuco::extent<std::size_t , set_capacity>;
@@ -103,7 +106,8 @@ TEMPLATE_TEST_CASE_SIG("static_set shared memory tests",
103106 // operator yet
104107 std::vector<std::unique_ptr<set_type>> sets;
105108 for (std::size_t set_id = 0 ; set_id < number_of_sets; ++set_id) {
106- sets.push_back (std::make_unique<set_type>(extent_type{}, cuco::empty_key<Key>{-1 }));
109+ sets.push_back (
110+ std::make_unique<set_type>(extent_type{}, cuco::empty_key<Key>{static_cast <Key>(-1 )}));
107111 }
108112
109113 thrust::device_vector<bool > d_keys_exist (number_of_sets * elements_in_set);
0 commit comments