From e7b5a389b823bebe465cefe63ad7b0e95f7fb450 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Mon, 15 Jul 2024 09:06:38 -0700
Subject: [PATCH] Fix a packed_cas build failure with libcudf (#543)

This PR casts data involved in `packed_cas` to `packed_type` to get rid
of build failures.
---
 .../open_addressing_ref_impl.cuh              | 20 +++++++++++--------
 1 file changed, 12 insertions(+), 8 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
index f4169d696..7dccfb353 100644
--- a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
+++ b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh
@@ -1233,20 +1233,24 @@ class open_addressing_ref_impl {
    */
   template <typename Value>
   [[nodiscard]] __device__ constexpr insert_result packed_cas(value_type* address,
-                                                              value_type const& expected,
-                                                              Value const& desired) noexcept
+                                                              value_type expected,
+                                                              Value desired) noexcept
   {
-    cuda::atomic_ref<value_type, Scope> slot_ref(*address);
-    auto expected_slot = expected;
+    using packed_type = cuda::std::conditional_t<sizeof(value_type) == 4, uint32_t, uint64_t>;
 
-    auto const success = slot_ref.compare_exchange_strong(
-      expected_slot, this->native_value(desired), cuda::std::memory_order_relaxed);
+    auto* slot_ptr     = reinterpret_cast<packed_type*>(address);
+    auto* expected_ptr = reinterpret_cast<packed_type*>(&expected);
+    auto* desired_ptr  = reinterpret_cast<packed_type*>(&desired);
+
+    auto slot_ref = cuda::atomic_ref<packed_type, Scope>{*slot_ptr};
+
+    auto const success =
+      slot_ref.compare_exchange_strong(*expected_ptr, *desired_ptr, cuda::memory_order_relaxed);
 
     if (success) {
       return insert_result::SUCCESS;
     } else {
-      return this->predicate_.equal_to(this->extract_key(desired),
-                                       this->extract_key(expected_slot)) ==
+      return this->predicate_.equal_to(this->extract_key(desired), this->extract_key(expected)) ==
                  detail::equal_result::EQUAL
                ? insert_result::DUPLICATE
                : insert_result::CONTINUE;