From 7608ed3653c6df2b43d4553d428fc0f7faf2c281 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Fri, 17 Jul 2026 20:52:58 +0000 Subject: [PATCH 1/2] Replace cuda::atomic_ref with plain CUDA atomics --- .../detail/hyperloglog/hyperloglog_impl.cuh | 4 +- .../__cuco/detail/open_addressing/kernels.cuh | 5 +- .../open_addressing_ref_impl.cuh | 37 +-- .../__cuco/detail/utility/atomic.cuh | 295 ++++++++++++++++++ cudax/test/CMakeLists.txt | 4 + cudax/test/cuco/utility/test_atomic.cu | 101 ++++++ 6 files changed, 419 insertions(+), 27 deletions(-) create mode 100644 cudax/include/cuda/experimental/__cuco/detail/utility/atomic.cuh create mode 100644 cudax/test/cuco/utility/test_atomic.cu diff --git a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh index a14a698571f..a913fda2496 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh @@ -39,6 +39,7 @@ #include #include +#include #include #include @@ -565,8 +566,7 @@ private: //! @param __value New value _CCCL_DEVICE_API constexpr void __update_max(int __i, __register_type __value) noexcept { - ::cuda::atomic_ref<__register_type, _Scope> __register_ref(__sketch[__i]); - __register_ref.fetch_max(__value, ::cuda::memory_order_relaxed); + (void) ::cuda::experimental::cuco::detail::__atomic_fetch_max<_Scope>(&__sketch[__i], __value); } //! @brief Try expanding the shmem partition for a given kernel beyond 48KB if necessary. diff --git a/cudax/include/cuda/experimental/__cuco/detail/open_addressing/kernels.cuh b/cudax/include/cuda/experimental/__cuco/detail/open_addressing/kernels.cuh index 93ae4b7f7c6..f98b4d51a7a 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/open_addressing/kernels.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/open_addressing/kernels.cuh @@ -23,9 +23,9 @@ #include -#include #include +#include #include #include @@ -128,8 +128,7 @@ _CCCL_KERNEL_ATTRIBUTES _CCCL_LAUNCH_BOUNDS(_BlockSize) void __insert_if_n( const auto __block_num_successes = __block_reduce(__temp_storage).Sum(__thread_num_successes); if (threadIdx.x == 0) { - ::cuda::atomic_ref{*__num_successes}.fetch_add( - __block_num_successes, ::cuda::std::memory_order_relaxed); + ::cuda::experimental::cuco::detail::__atomic_fetch_add<_Ref::thread_scope>(__num_successes, __block_num_successes); } } diff --git a/cudax/include/cuda/experimental/__cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/cudax/include/cuda/experimental/__cuco/detail/open_addressing/open_addressing_ref_impl.cuh index 1001496207b..2af7304518f 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/open_addressing/open_addressing_ref_impl.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/open_addressing/open_addressing_ref_impl.cuh @@ -23,7 +23,6 @@ #include -#include #include #include #include @@ -36,6 +35,7 @@ #include #include +#include #include #include #include @@ -706,10 +706,8 @@ public: auto* __expected_ptr = reinterpret_cast(&__expected); auto* __desired_ptr = reinterpret_cast(&__desired); - auto __slot_ref = ::cuda::atomic_ref{*__slot_ptr}; - - const auto success = - __slot_ref.compare_exchange_strong(*__expected_ptr, *__desired_ptr, ::cuda::memory_order_relaxed); + const auto success = ::cuda::experimental::cuco::detail::__atomic_compare_exchange<_Scope>( + __slot_ptr, *__expected_ptr, *__desired_ptr); if (success) { @@ -745,28 +743,26 @@ public: auto __expected_key = __expected.first; auto __expected_payload = empty_value_sentinel(); - ::cuda::atomic_ref<__key_type, _Scope> key_ref(__address->first); - ::cuda::atomic_ref payload_ref(__address->second); - - const auto key_cas_success = - key_ref.compare_exchange_strong(__expected_key, __key_type{__desired.first}, ::cuda::memory_order_relaxed); - auto payload_cas_success = - payload_ref.compare_exchange_strong(__expected_payload, __desired.second, ::cuda::memory_order_relaxed); + const auto key_cas_success = ::cuda::experimental::cuco::detail::__atomic_compare_exchange<_Scope>( + &__address->first, __expected_key, __key_type{__desired.first}); + auto payload_cas_success = ::cuda::experimental::cuco::detail::__atomic_compare_exchange<_Scope>( + &__address->second, __expected_payload, mapped_type{__desired.second}); // if __key success if (key_cas_success) { while (!payload_cas_success) { - payload_cas_success = payload_ref.compare_exchange_strong( - __expected_payload = empty_value_sentinel(), __desired.second, ::cuda::memory_order_relaxed); + __expected_payload = empty_value_sentinel(); + payload_cas_success = ::cuda::experimental::cuco::detail::__atomic_compare_exchange<_Scope>( + &__address->second, __expected_payload, mapped_type{__desired.second}); } return __insert_result::__success; } else if (payload_cas_success) { // This is insert-specific, cannot for `erase` operations - payload_ref.store(empty_value_sentinel(), ::cuda::memory_order_relaxed); + ::cuda::experimental::cuco::detail::__atomic_store<_Scope>(&__address->second, empty_value_sentinel()); } // Our __key was already present in the slot, so our __key is a duplicate @@ -795,16 +791,14 @@ public: { using mapped_type = ::cuda::std::decay_t; - ::cuda::atomic_ref<__key_type, _Scope> key_ref(__address->first); auto __expected_key = __expected.first; - const auto success = - key_ref.compare_exchange_strong(__expected_key, __key_type{__desired.first}, ::cuda::memory_order_relaxed); + const auto success = ::cuda::experimental::cuco::detail::__atomic_compare_exchange<_Scope>( + &__address->first, __expected_key, __key_type{__desired.first}); // if __key success if (success) { - ::cuda::atomic_ref payload_ref(__address->second); - payload_ref.store(__desired.second, ::cuda::memory_order_relaxed); + ::cuda::experimental::cuco::detail::__atomic_store<_Scope>(&__address->second, mapped_type{__desired.second}); return __insert_result::__success; } @@ -892,12 +886,11 @@ public: template _CCCL_DEVICE_API void __wait_for_payload(_Value& __slot, _Value __sentinel) const noexcept { - auto __ref = ::cuda::atomic_ref<_Value, _Scope>{__slot}; _Value __current; // TODO exponential backoff strategy do { - __current = __ref.load(::cuda::std::memory_order_relaxed); + __current = ::cuda::experimental::cuco::detail::__atomic_load<_Scope>(&__slot); } while (detail::__bitwise_compare(__current, __sentinel)); } #endif // _CCCL_CUDA_COMPILATION() diff --git a/cudax/include/cuda/experimental/__cuco/detail/utility/atomic.cuh b/cudax/include/cuda/experimental/__cuco/detail/utility/atomic.cuh new file mode 100644 index 00000000000..2e335f1b26a --- /dev/null +++ b/cudax/include/cuda/experimental/__cuco/detail/utility/atomic.cuh @@ -0,0 +1,295 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDAX___CUCO_DETAIL_UTILITY_ATOMIC_CUH +#define _CUDAX___CUCO_DETAIL_UTILITY_ATOMIC_CUH + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +namespace cuda::experimental::cuco::detail +{ +#if _CCCL_CUDA_COMPILATION() +template <::cuda::std::size_t _Size> +struct __atomic_word; + +template <> +struct __atomic_word<1> +{ + using __type = ::cuda::std::uint8_t; +}; + +template <> +struct __atomic_word<2> +{ + using __type = ::cuda::std::uint16_t; +}; + +template <> +struct __atomic_word<4> +{ + using __type = unsigned int; +}; + +template <> +struct __atomic_word<8> +{ + using __type = unsigned long long; +}; + +template +using __atomic_word_t = typename __atomic_word::__type; + +template <::cuda::thread_scope _Scope> +inline constexpr bool __is_cuda_atomic_scope = + _Scope == ::cuda::thread_scope_thread || _Scope == ::cuda::thread_scope_block || _Scope == ::cuda::thread_scope_device + || _Scope == ::cuda::thread_scope_system; + +template <::cuda::thread_scope _Scope, class _Word> +[[nodiscard]] _CCCL_DEVICE_API _Word __atomic_cas_word(_Word* __address, _Word __compare, _Word __value) noexcept +{ + static_assert(__is_cuda_atomic_scope<_Scope>, "cuCO atomics support thread, block, device, and system thread scopes"); + static_assert(sizeof(_Word) == 4 || sizeof(_Word) == 8, "CUDA atomicCAS requires a 32-bit or 64-bit word"); + + if constexpr (_Scope == ::cuda::thread_scope_thread) + { + const _Word __old = *__address; + if (__old == __compare) + { + *__address = __value; + } + return __old; + } + else if constexpr (_Scope == ::cuda::thread_scope_block) + { + return ::atomicCAS_block(__address, __compare, __value); + } + else if constexpr (_Scope == ::cuda::thread_scope_device) + { + return ::atomicCAS(__address, __compare, __value); + } + else + { + return ::atomicCAS_system(__address, __compare, __value); + } +} + +template <::cuda::thread_scope _Scope, class _Word> +[[nodiscard]] _CCCL_DEVICE_API _Word __atomic_exchange_word(_Word* __address, _Word __value) noexcept +{ + static_assert(__is_cuda_atomic_scope<_Scope>, "cuCO atomics support thread, block, device, and system thread scopes"); + static_assert(sizeof(_Word) == 4 || sizeof(_Word) == 8, "CUDA atomicExch requires a 32-bit or 64-bit word"); + + if constexpr (_Scope == ::cuda::thread_scope_thread) + { + const _Word __old = *__address; + *__address = __value; + return __old; + } + else if constexpr (_Scope == ::cuda::thread_scope_block) + { + return ::atomicExch_block(__address, __value); + } + else if constexpr (_Scope == ::cuda::thread_scope_device) + { + return ::atomicExch(__address, __value); + } + else + { + return ::atomicExch_system(__address, __value); + } +} + +template +_CCCL_DEVICE_API constexpr void __validate_atomic_type() noexcept +{ + static_assert(::cuda::std::is_trivially_copyable_v<_Tp>, "cuCO atomics require a trivially copyable type"); + static_assert(sizeof(_Tp) == 1 || sizeof(_Tp) == 2 || sizeof(_Tp) == 4 || sizeof(_Tp) == 8, + "cuCO atomics require a 1, 2, 4, or 8 byte type"); +} + +template <::cuda::thread_scope _Scope, class _Tp> +[[nodiscard]] _CCCL_DEVICE_API bool __atomic_compare_exchange(_Tp* __address, _Tp& __expected, _Tp __desired) noexcept +{ + ::cuda::experimental::cuco::detail::__validate_atomic_type<_Tp>(); + + using __word_type = __atomic_word_t<_Tp>; + const __word_type __compare = ::cuda::std::bit_cast<__word_type>(__expected); + const __word_type __value = ::cuda::std::bit_cast<__word_type>(__desired); + + if constexpr (sizeof(_Tp) <= 2) + { + constexpr ::cuda::std::uintptr_t __align_mask = sizeof(unsigned int) - 1; + constexpr unsigned int __size_mask = (1u << (sizeof(_Tp) * 8)) - 1; + const auto __address_value = reinterpret_cast<::cuda::std::uintptr_t>(__address); + auto* const __aligned = reinterpret_cast(__address_value & ~__align_mask); + const auto __offset = static_cast((__address_value & __align_mask) * 8); + const auto __value_mask = __size_mask << __offset; + const auto __window_mask = ~__value_mask; + const auto __compare_bits = static_cast(__compare) << __offset; + const auto __value_bits = static_cast(__value) << __offset; + + auto __old = ::cuda::experimental::cuco::detail::__atomic_cas_word<_Scope>(__aligned, 0u, 0u); + while ((__old & __value_mask) == __compare_bits) + { + const auto __attempt = (__old & __window_mask) | __value_bits; + const auto __result = ::cuda::experimental::cuco::detail::__atomic_cas_word<_Scope>(__aligned, __old, __attempt); + if (__result == __old) + { + return true; + } + __old = __result; + } + __expected = ::cuda::std::bit_cast<_Tp>(static_cast<__word_type>((__old & __value_mask) >> __offset)); + return false; + } + else + { + auto* const __word_address = reinterpret_cast<__word_type*>(__address); + const auto __old = + ::cuda::experimental::cuco::detail::__atomic_cas_word<_Scope>(__word_address, __compare, __value); + const bool __success = __old == __compare; + if (!__success) + { + __expected = ::cuda::std::bit_cast<_Tp>(__old); + } + return __success; + } +} + +template <::cuda::thread_scope _Scope, class _Tp> +[[nodiscard]] _CCCL_DEVICE_API _Tp __atomic_load(const _Tp* __address) noexcept +{ + ::cuda::experimental::cuco::detail::__validate_atomic_type<_Tp>(); + static_assert(sizeof(_Tp) == 4 || sizeof(_Tp) == 8, "cuCO atomic load requires a 4 or 8 byte type"); + + using __word_type = __atomic_word_t<_Tp>; + auto* const __word_address = reinterpret_cast<__word_type*>(const_cast<_Tp*>(__address)); + constexpr __word_type __zero = 0; + const auto __old = ::cuda::experimental::cuco::detail::__atomic_cas_word<_Scope>(__word_address, __zero, __zero); + return ::cuda::std::bit_cast<_Tp>(__old); +} + +template <::cuda::thread_scope _Scope, class _Tp> +_CCCL_DEVICE_API void __atomic_store(_Tp* __address, _Tp __value) noexcept +{ + ::cuda::experimental::cuco::detail::__validate_atomic_type<_Tp>(); + static_assert(sizeof(_Tp) == 4 || sizeof(_Tp) == 8, "cuCO atomic store requires a 4 or 8 byte type"); + + using __word_type = __atomic_word_t<_Tp>; + auto* const __word_address = reinterpret_cast<__word_type*>(__address); + const __word_type __desired = ::cuda::std::bit_cast<__word_type>(__value); + (void) ::cuda::experimental::cuco::detail::__atomic_exchange_word<_Scope>(__word_address, __desired); +} + +template <::cuda::thread_scope _Scope, class _Tp> +[[nodiscard]] _CCCL_DEVICE_API _Tp __atomic_fetch_add(_Tp* __address, _Tp __value) noexcept +{ + ::cuda::experimental::cuco::detail::__validate_atomic_type<_Tp>(); + static_assert(::cuda::std::is_integral_v<_Tp> && !::cuda::std::is_same_v<_Tp, bool>, + "cuCO atomic add requires a non-bool integral type"); + static_assert(sizeof(_Tp) == 4 || sizeof(_Tp) == 8, "cuCO atomic add requires a 4 or 8 byte type"); + + using __word_type = __atomic_word_t<_Tp>; + auto* const __word_address = reinterpret_cast<__word_type*>(__address); + const auto __addend = ::cuda::std::bit_cast<__word_type>(__value); + __word_type __old; + if constexpr (_Scope == ::cuda::thread_scope_thread) + { + __old = *__word_address; + *__word_address = __old + __addend; + } + else if constexpr (_Scope == ::cuda::thread_scope_block) + { + __old = ::atomicAdd_block(__word_address, __addend); + } + else if constexpr (_Scope == ::cuda::thread_scope_device) + { + __old = ::atomicAdd(__word_address, __addend); + } + else if constexpr (_Scope == ::cuda::thread_scope_system) + { + __old = ::atomicAdd_system(__word_address, __addend); + } + else + { + static_assert(__is_cuda_atomic_scope<_Scope>, + "cuCO atomics support thread, block, device, and system thread scopes"); + } + return ::cuda::std::bit_cast<_Tp>(__old); +} + +template <::cuda::thread_scope _Scope, class _Tp> +[[nodiscard]] _CCCL_DEVICE_API _Tp __atomic_fetch_max(_Tp* __address, _Tp __value) noexcept +{ + ::cuda::experimental::cuco::detail::__validate_atomic_type<_Tp>(); + static_assert(::cuda::std::is_integral_v<_Tp> && !::cuda::std::is_same_v<_Tp, bool>, + "cuCO atomic max requires a non-bool integral type"); + static_assert(sizeof(_Tp) == 4 || sizeof(_Tp) == 8, "cuCO atomic max requires a 4 or 8 byte type"); + + using __native_type = + ::cuda::std::conditional_t, int, unsigned int>, + ::cuda::std::conditional_t<::cuda::std::is_signed_v<_Tp>, long long, unsigned long long>>; + auto* const __native_address = reinterpret_cast<__native_type*>(__address); + const auto __native_value = ::cuda::std::bit_cast<__native_type>(__value); + __native_type __old; + if constexpr (_Scope == ::cuda::thread_scope_thread) + { + __old = *__native_address; + if (__old < __native_value) + { + *__native_address = __native_value; + } + } + else if constexpr (_Scope == ::cuda::thread_scope_block) + { + __old = ::atomicMax_block(__native_address, __native_value); + } + else if constexpr (_Scope == ::cuda::thread_scope_device) + { + __old = ::atomicMax(__native_address, __native_value); + } + else if constexpr (_Scope == ::cuda::thread_scope_system) + { + __old = ::atomicMax_system(__native_address, __native_value); + } + else + { + static_assert(__is_cuda_atomic_scope<_Scope>, + "cuCO atomics support thread, block, device, and system thread scopes"); + } + return ::cuda::std::bit_cast<_Tp>(__old); +} +#endif // _CCCL_CUDA_COMPILATION() +} // namespace cuda::experimental::cuco::detail + +#include + +#endif // _CUDAX___CUCO_DETAIL_UTILITY_ATOMIC_CUH diff --git a/cudax/test/CMakeLists.txt b/cudax/test/CMakeLists.txt index f91ca95fdc5..40e15d5df0f 100644 --- a/cudax/test/CMakeLists.txt +++ b/cudax/test/CMakeLists.txt @@ -135,6 +135,10 @@ cudax_add_catch2_test(test_target cuco.fixed_capacity_map.key_sentinel ${cudax_t cuco/fixed_capacity_map/test_key_sentinel.cu ) +cudax_add_catch2_test(test_target cuco.utility.atomic + cuco/utility/test_atomic.cu +) + cudax_add_catch2_test(test_target cuco.utility.capacity ${cudax_target} cuco/utility/test_capacity.cu ) diff --git a/cudax/test/cuco/utility/test_atomic.cu b/cudax/test/cuco/utility/test_atomic.cu new file mode 100644 index 00000000000..79dcae4926f --- /dev/null +++ b/cudax/test/cuco/utility/test_atomic.cu @@ -0,0 +1,101 @@ +//===----------------------------------------------------------------------===// +// +// Part of CUDA Experimental in CUDA C++ Core Libraries, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#include + +#include + +#include +#include + +namespace cudax = cuda::experimental; + +struct atomic_test_values +{ + alignas(4)::cuda::std::uint8_t cas8; + ::cuda::std::uint8_t guard8_0; + ::cuda::std::uint8_t guard8_1; + ::cuda::std::uint8_t guard8_2; + + alignas(4)::cuda::std::uint16_t cas16; + ::cuda::std::uint16_t guard16; + + ::cuda::std::int32_t value32; + ::cuda::std::int64_t value64; + ::cuda::std::int32_t maximum32; + ::cuda::std::int32_t cas32; + ::cuda::std::int32_t stored32; + ::cuda::std::int32_t loaded32; +}; + +__global__ void test_atomic_kernel(atomic_test_values* values) +{ + const auto rank = static_cast<::cuda::std::int32_t>(threadIdx.x); + + (void) cudax::cuco::detail::__atomic_fetch_add<::cuda::thread_scope_device>(&values->value32, ::cuda::std::int32_t{1}); + (void) cudax::cuco::detail::__atomic_fetch_add<::cuda::thread_scope_system>(&values->value64, ::cuda::std::int64_t{1}); + (void) cudax::cuco::detail::__atomic_fetch_max<::cuda::thread_scope_block>(&values->maximum32, rank); + + auto expected8 = ::cuda::std::uint8_t{0}; + (void) cudax::cuco::detail::__atomic_compare_exchange<::cuda::thread_scope_device>( + &values->cas8, expected8, static_cast<::cuda::std::uint8_t>(rank + 1)); + + auto expected16 = ::cuda::std::uint16_t{0}; + (void) cudax::cuco::detail::__atomic_compare_exchange<::cuda::thread_scope_device>( + &values->cas16, expected16, static_cast<::cuda::std::uint16_t>(rank + 1)); + + auto expected32 = ::cuda::std::int32_t{0}; + (void) cudax::cuco::detail::__atomic_compare_exchange<::cuda::thread_scope_device>( + &values->cas32, expected32, rank + 1); + + __syncthreads(); + if (rank == 0) + { + cudax::cuco::detail::__atomic_store<::cuda::thread_scope_device>(&values->stored32, ::cuda::std::int32_t{9}); + values->loaded32 = cudax::cuco::detail::__atomic_load<::cuda::thread_scope_device>(&values->stored32); + } +} + +C2H_TEST("cuco plain CUDA atomics preserve values and packed neighbors", "[atomic]") +{ + atomic_test_values host_values{}; + host_values.guard8_0 = 0xA5; + host_values.guard8_1 = 0x5A; + host_values.guard8_2 = 0xC3; + host_values.guard16 = 0xA55A; + host_values.stored32 = 3; + + atomic_test_values* device_values{}; + REQUIRE_CUDART(::cudaMalloc(&device_values, sizeof(atomic_test_values))); + REQUIRE_CUDART(::cudaMemcpy(device_values, &host_values, sizeof(atomic_test_values), ::cudaMemcpyHostToDevice)); + + constexpr int block_size = 128; + test_atomic_kernel<<<1, block_size>>>(device_values); + REQUIRE_CUDART(::cudaGetLastError()); + REQUIRE_CUDART(::cudaDeviceSynchronize()); + REQUIRE_CUDART(::cudaMemcpy(&host_values, device_values, sizeof(atomic_test_values), ::cudaMemcpyDeviceToHost)); + REQUIRE_CUDART(::cudaFree(device_values)); + + REQUIRE(host_values.cas8 >= 1); + REQUIRE(host_values.cas8 <= block_size); + REQUIRE(host_values.guard8_0 == 0xA5); + REQUIRE(host_values.guard8_1 == 0x5A); + REQUIRE(host_values.guard8_2 == 0xC3); + REQUIRE(host_values.cas16 >= 1); + REQUIRE(host_values.cas16 <= block_size); + REQUIRE(host_values.guard16 == 0xA55A); + REQUIRE(host_values.value32 == block_size); + REQUIRE(host_values.value64 == block_size); + REQUIRE(host_values.maximum32 == block_size - 1); + REQUIRE(host_values.cas32 >= 1); + REQUIRE(host_values.cas32 <= block_size); + REQUIRE(host_values.stored32 == 9); + REQUIRE(host_values.loaded32 == 9); +} From 172e57e5ac8d4a929195e29e512de2c819b7d2a3 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 20 Jul 2026 01:40:29 +0000 Subject: [PATCH 2/2] Fix nodiscard warning in cuco insert kernel --- .../experimental/__cuco/detail/open_addressing/kernels.cuh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cudax/include/cuda/experimental/__cuco/detail/open_addressing/kernels.cuh b/cudax/include/cuda/experimental/__cuco/detail/open_addressing/kernels.cuh index f98b4d51a7a..0d2aae378ce 100644 --- a/cudax/include/cuda/experimental/__cuco/detail/open_addressing/kernels.cuh +++ b/cudax/include/cuda/experimental/__cuco/detail/open_addressing/kernels.cuh @@ -128,7 +128,8 @@ _CCCL_KERNEL_ATTRIBUTES _CCCL_LAUNCH_BOUNDS(_BlockSize) void __insert_if_n( const auto __block_num_successes = __block_reduce(__temp_storage).Sum(__thread_num_successes); if (threadIdx.x == 0) { - ::cuda::experimental::cuco::detail::__atomic_fetch_add<_Ref::thread_scope>(__num_successes, __block_num_successes); + (void) ::cuda::experimental::cuco::detail::__atomic_fetch_add<_Ref::thread_scope>( + __num_successes, __block_num_successes); } }