Updates

NVIDIA · Nov 10, 2023 · 38be8cf · 38be8cf
1 parent 9326c33
commit 38be8cf
Show file tree

Hide file tree

Showing 2 changed files with 3 additions and 13 deletions.
diff --git a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
@@ -402,12 +402,10 @@ class open_addressing_ref_impl {
             using mapped_type           = decltype(this->empty_slot_sentinel_.second);
             auto const expected_payload = this->empty_slot_sentinel_.second;
             auto ref =
-              cuda::atomic<mapped_type, cuda::thread_scope_device>{(window_ptr + i)->second};
+              cuda::atomic_ref<mapped_type, cuda::thread_scope_device>{(window_ptr + i)->second};
             mapped_type old;
-            int n = 0;
             do {
               old = ref.load(cuda::std::memory_order_relaxed);
-              // printf("old: %d expected_payload: %d n: %d\n", int(old), int(expected_payload), n);
             } while (cuco::detail::bitwise_compare(old, expected_payload));
           }
           return {iterator{&window_ptr[i]}, false};
@@ -423,13 +421,10 @@ class open_addressing_ref_impl {
               auto const res    = cas_dependent_write(window_ptr + i, window_slots[i], value);
               using mapped_type = decltype(this->empty_slot_sentinel_.second);
               auto ref =
-                cuda::atomic<mapped_type, cuda::thread_scope_device>{(window_ptr + i)->second};
+                cuda::atomic_ref<mapped_type, cuda::thread_scope_device>{(window_ptr + i)->second};
               mapped_type old;
-              int n = 0;
               do {
                 old = ref.load(cuda::std::memory_order_relaxed);
-                // printf("old: %d expected_payload: %d n: %d\n", int(old), int(expected_payload),
-                // n);
               } while (cuco::detail::bitwise_compare(old, expected_payload));
               return res;
             }

diff --git a/tests/static_map/insert_and_find_test.cu b/tests/static_map/insert_and_find_test.cu
@@ -37,7 +37,6 @@ __global__ void parallel_sum(Ref v)
 #endif
     {
       auto [iter, inserted] = v.insert_and_find(cuco::pair{i, 1});
-      if (inserted) { printf("key: %d payload:%d \n", int(iter->first), int(iter->second)); }
       // for debugging...
       // if (iter->second < 0) {
       //   asm("trap;");
@@ -73,7 +72,7 @@ TEMPLATE_TEST_CASE_SIG("Parallel insert-or-update",
     thrust::equal_to<Key>{},
     cuco::experimental::linear_probing<1, cuco::murmurhash3_32<Key>>{}};
 
-  static constexpr int Blocks  = 64;
+  static constexpr int Blocks  = 1024;
   static constexpr int Threads = 128;
 
   parallel_sum<<<Blocks, Threads>>>(
@@ -86,10 +85,6 @@ TEMPLATE_TEST_CASE_SIG("Parallel insert-or-update",
   thrust::sequence(thrust::device, d_keys.begin(), d_keys.end());
   m.find(d_keys.begin(), d_keys.end(), d_values.begin());
 
-  for (auto const t : d_values) {
-    printf("### %d\n", int(t));
-  }
-
   REQUIRE(cuco::test::all_of(
     d_values.begin(), d_values.end(), [] __device__(Value v) { return v == Blocks * Threads; }));
 }