Fix CAS dispatching bug (#533)

This PR fixes a macro typo when dispatching the between back-to-back CAS and CAS-dependent write.
NVIDIA · Jul 9, 2024 · 63d313e · 63d313e
1 parent 687948a
commit 63d313e
Show file tree

Hide file tree

Showing 2 changed files with 4 additions and 3 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
@@ -23,6 +23,7 @@
 #include <cuco/probing_scheme.cuh>
 
 #include <cuda/atomic>
+#include <cuda/std/type_traits>
 #include <thrust/distance.h>
 #include <thrust/tuple.h>
 #if defined(CUCO_HAS_CUDA_BARRIER)
@@ -1345,7 +1346,7 @@ class open_addressing_ref_impl {
                                                                     value_type const& expected,
                                                                     Value const& desired) noexcept
   {
-    using mapped_type = decltype(this->empty_value_sentinel());
+    using mapped_type = cuda::std::decay_t<decltype(this->empty_value_sentinel())>;
 
     auto const expected_key     = expected.first;
     auto const expected_payload = expected.second;
@@ -1437,7 +1438,7 @@ class open_addressing_ref_impl {
     if constexpr (sizeof(value_type) <= 8) {
       return packed_cas(address, expected, desired);
     } else {
-#if (_CUDA_ARCH__ < 700)
+#if (__CUDA_ARCH__ < 700)
       return cas_dependent_write(address, expected, desired);
 #else
       return back_to_back_cas(address, expected, desired);

diff --git a/include/cuco/detail/static_map.inl b/include/cuco/detail/static_map.inl
@@ -408,7 +408,7 @@ __device__ bool static_map<Key, Value, Scope, Allocator>::device_mutable_view::i
         }
 
         if constexpr (not cuco::detail::is_packable<value_type>()) {
-#if (_CUDA_ARCH__ < 700)
+#if (__CUDA_ARCH__ < 700)
           return cas_dependent_write(current_slot, insert_pair, key_equal, existing_key);
 #else
           return back_to_back_cas(current_slot, insert_pair, key_equal, existing_key);