From dfc37039eb2c7c17ce9139fdf7569d74cd78a577 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Fri, 17 Jul 2026 15:04:49 +0200 Subject: [PATCH 1/4] [libcu++] Implement P3793R2 Better Shifting --- libcudacxx/include/cuda/std/__bit/shl.h | 87 +++++++++ libcudacxx/include/cuda/std/__bit/shr.h | 87 +++++++++ libcudacxx/include/cuda/std/bit | 4 +- .../std/numerics/bit/bit.shift/shl.pass.cpp | 164 ++++++++++++++++ .../std/numerics/bit/bit.shift/shr.pass.cpp | 176 ++++++++++++++++++ libcudacxx/test/support/test_macros.h | 11 +- 6 files changed, 524 insertions(+), 5 deletions(-) create mode 100644 libcudacxx/include/cuda/std/__bit/shl.h create mode 100644 libcudacxx/include/cuda/std/__bit/shr.h create mode 100644 libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp diff --git a/libcudacxx/include/cuda/std/__bit/shl.h b/libcudacxx/include/cuda/std/__bit/shl.h new file mode 100644 index 00000000000..7cc1c69d669 --- /dev/null +++ b/libcudacxx/include/cuda/std/__bit/shl.h @@ -0,0 +1,87 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// 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 _CUDA_STD___BIT_SHL_H +#define _CUDA_STD___BIT_SHL_H + +#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 + +#if _CCCL_CUDA_COMPILATION() +# include +# include +#endif // _CCCL_CUDA_COMPILATION() + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA_STD + +_CCCL_TEMPLATE(class _Tp, class _Shift) +_CCCL_REQUIRES(__cccl_is_integer_v<_Tp> _CCCL_AND __cccl_is_integer_v<_Shift>) +[[nodiscard]] _CCCL_API constexpr _Tp shl(_Tp __v, _Shift __shift) noexcept +{ + constexpr auto __width = static_cast(__num_bits_v<_Tp>); + const auto __ushift = ::cuda::uabs(__shift); + + if constexpr (is_signed_v<_Shift>) + { + if (__shift < 0) + { + _CCCL_IF_NOT_CONSTEVAL_DEFAULT + { + // On device, shr PTX instruction clamps the shift to width, however only 32-bit shifts are supported. + NV_IF_TARGET(NV_IS_DEVICE, ({ + if constexpr (sizeof(_Shift) <= sizeof(uint32_t) && sizeof(_Tp) <= sizeof(int64_t)) + { + using _Up = __make_nbit_int_t>; + return static_cast<_Tp>(::cuda::ptx::shr(_Up{__v}, static_cast(__ushift))); + } + })) + } + return (__ushift < __width) ? (__v >> __ushift) : static_cast<_Tp>(::cuda::std::cmp_less(__v, 0) ? -1 : 0); + } + } + + _CCCL_IF_NOT_CONSTEVAL_DEFAULT + { + // On device, shl PTX instruction clamps the shift to width, however only 32-bit shifts are supported. + NV_IF_TARGET(NV_IS_DEVICE, ({ + if constexpr (sizeof(_Shift) <= sizeof(uint32_t) && sizeof(_Tp) <= sizeof(int64_t)) + { + using _Up = __make_nbit_int_t>; + return static_cast<_Tp>(::cuda::ptx::shl(_Up{__v}, static_cast(__ushift))); + } + })) + } + return (__ushift < __width) ? static_cast<_Tp>(::cuda::std::__to_unsigned_like(__v) << __ushift) : _Tp{0}; +} + +_CCCL_END_NAMESPACE_CUDA_STD + +#include + +#endif // _CUDA_STD___BIT_SHL_H diff --git a/libcudacxx/include/cuda/std/__bit/shr.h b/libcudacxx/include/cuda/std/__bit/shr.h new file mode 100644 index 00000000000..1a8c085616d --- /dev/null +++ b/libcudacxx/include/cuda/std/__bit/shr.h @@ -0,0 +1,87 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// 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 _CUDA_STD___BIT_SHR_H +#define _CUDA_STD___BIT_SHR_H + +#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 + +#if _CCCL_CUDA_COMPILATION() +# include +# include +#endif // _CCCL_CUDA_COMPILATION() + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA_STD + +_CCCL_TEMPLATE(class _Tp, class _Shift) +_CCCL_REQUIRES(__cccl_is_integer_v<_Tp> _CCCL_AND __cccl_is_integer_v<_Shift>) +[[nodiscard]] _CCCL_API constexpr _Tp shr(_Tp __v, _Shift __shift) noexcept +{ + constexpr auto __width = static_cast(__num_bits_v<_Tp>); + const auto __ushift = ::cuda::uabs(__shift); + + if constexpr (is_signed_v<_Shift>) + { + if (__shift < 0) + { + _CCCL_IF_NOT_CONSTEVAL_DEFAULT + { + // On device, shl PTX instruction clamps the shift to width, however only 32-bit shifts are supported. + NV_IF_TARGET(NV_IS_DEVICE, ({ + if constexpr (sizeof(_Shift) <= sizeof(uint32_t) && sizeof(_Tp) <= sizeof(int64_t)) + { + using _Up = __make_nbit_int_t>; + return static_cast<_Tp>(::cuda::ptx::shl(_Up{__v}, static_cast(__ushift))); + } + })) + } + return (__ushift < __width) ? static_cast<_Tp>(::cuda::std::__to_unsigned_like(__v) << __ushift) : _Tp{0}; + } + } + + _CCCL_IF_NOT_CONSTEVAL_DEFAULT + { + // On device, shr PTX instruction clamps the shift to width, however only 32-bit shifts are supported. + NV_IF_TARGET(NV_IS_DEVICE, ({ + if constexpr (sizeof(_Shift) <= sizeof(uint32_t) && sizeof(_Tp) <= sizeof(int64_t)) + { + using _Up = __make_nbit_int_t>; + return static_cast<_Tp>(::cuda::ptx::shr(_Up{__v}, static_cast(__ushift))); + } + })) + } + return (__ushift < __width) ? (__v >> __ushift) : static_cast<_Tp>(::cuda::std::cmp_less(__v, 0) ? -1 : 0); +} + +_CCCL_END_NAMESPACE_CUDA_STD + +#include + +#endif // _CUDA_STD___BIT_SHR_H diff --git a/libcudacxx/include/cuda/std/bit b/libcudacxx/include/cuda/std/bit index b0f33a48c22..0140697f317 100644 --- a/libcudacxx/include/cuda/std/bit +++ b/libcudacxx/include/cuda/std/bit @@ -4,7 +4,7 @@ // 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) 2024 NVIDIA CORPORATION & AFFILIATES. +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. // //===----------------------------------------------------------------------===// @@ -30,6 +30,8 @@ #include #include #include +#include +#include #include #endif // _CUDA_STD_BIT diff --git a/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp new file mode 100644 index 00000000000..c3525387de3 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp @@ -0,0 +1,164 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// 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. +// +//===----------------------------------------------------------------------===// + +// template +// constexpr T shl(T x, S s) noexcept; + +#include +#include +#include +#include +#include +#include +#include + +#include "test_macros.h" + +template +TEST_FUNC constexpr T invoke_shl(T v, S shift) +{ + if (cuda::std::is_constant_evaluated()) + { + return cuda::std::shl(v, shift); + } + else + { + DoNotOptimize(v); + DoNotOptimize(shift); + return cuda::std::shl(v, shift); + } +} + +template +TEST_FUNC constexpr void test_pos(T v, S shift) +{ + assert(cuda::std::cmp_greater_equal(shift, 0)); + + auto result = invoke_shl(v, shift); + if (cuda::std::cmp_less(shift, sizeof(T) * CHAR_BIT)) + { + using U = cuda::std::make_unsigned_t; + assert(result == static_cast(static_cast(v) << shift)); + } + else + { + assert(result == T{0}); + } +} + +template +TEST_FUNC constexpr void test_neg(T v, S shift) +{ + assert(shift < 0); + + auto result = invoke_shl(v, shift); + if (cuda::std::cmp_less(cuda::uabs(shift), (sizeof(T) * CHAR_BIT))) + { + assert(result == static_cast(v >> cuda::uabs(shift))); + } + else + { + assert(result == T(cuda::std::cmp_less(v, 0) ? -1 : 0)); + } +} + +template +TEST_FUNC constexpr void test() +{ + constexpr auto tmin = cuda::std::numeric_limits::min(); + constexpr auto tmax = cuda::std::numeric_limits::max(); + constexpr auto smin = cuda::std::numeric_limits::min(); + constexpr auto smax = cuda::std::numeric_limits::max(); + + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::std::shl(T{}, S{}))); + + const T vs[] = {tmin, T(-24), T(-1), T{0}, T{1}, T{20}, T{99}, T(1225), tmax}; + const S pos_shifts[] = {S{0}, S{1}, S{7}, S{17}, S{23}, S{32}, S{33}, S{65}, smax}; + + // Disable loop unrolling to reduce ptxas compile times. + _CCCL_PRAGMA_NOUNROLL() + for (auto v : vs) + { + _CCCL_PRAGMA_NOUNROLL() + for (auto shift : pos_shifts) + { + test_pos(v, shift); + } + } + + if constexpr (cuda::std::is_signed_v) + { + const S neg_shifts[] = {smin, S{-65}, S{-33}, S{-32}, S{-15}, S{-7}, S{-1}}; + + _CCCL_PRAGMA_NOUNROLL() + for (auto v : vs) + { + _CCCL_PRAGMA_NOUNROLL() + for (auto shift : neg_shifts) + { + test_neg(v, shift); + } + } + } +} + +template +TEST_FUNC constexpr void test() +{ + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test(); +#endif // _CCCL_HAS_INT128() + + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test(); +#endif // _CCCL_HAS_INT128() +} + +TEST_FUNC constexpr bool test() +{ + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test<__int128_t>(); +#endif // _CCCL_HAS_INT128() + + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test<__uint128_t>(); +#endif // _CCCL_HAS_INT128() + + return true; +} + +int main(int, char**) +{ + test(); + static_assert(test()); + + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp new file mode 100644 index 00000000000..83887266103 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp @@ -0,0 +1,176 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// 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. +// +//===----------------------------------------------------------------------===// + +// template +// constexpr T shr(T x, S s) noexcept; + +#include +#include +#include +#include +#include +#include +#include + +#include "test_macros.h" + +template +TEST_FUNC constexpr T invoke_shr(T v, S shift) +{ + if (cuda::std::is_constant_evaluated()) + { + return cuda::std::shr(v, shift); + } + else + { + DoNotOptimize(v); + DoNotOptimize(shift); + return cuda::std::shr(v, shift); + } +} + +template +TEST_FUNC constexpr void test_pos(T v, S shift) +{ + assert(cuda::std::cmp_greater_equal(shift, 0)); + + auto result = invoke_shr(v, shift); + if (cuda::std::cmp_less(shift, (sizeof(T) * CHAR_BIT))) + { + assert(result == static_cast(v >> shift)); + } + else + { + assert(result == T(cuda::std::cmp_less(v, 0) ? -1 : 0)); + } +} + +template +TEST_FUNC constexpr void test_neg(T v, S shift) +{ + assert(shift < 0); + + // There is a bug on sm100+ when cuda::std::shr returns invalid result when the shift is int64_t min. Re-enable + // once nvbug 6471375 is resolved. + NV_IF_TARGET(NV_PROVIDES_SM_100, ({ + if constexpr (sizeof(S) == sizeof(cuda::std::int64_t)) + { + if (shift == cuda::std::numeric_limits::min()) + { + return; + } + } + })) + + auto result = invoke_shr(v, shift); + if (cuda::std::cmp_less(cuda::uabs(shift), sizeof(T) * CHAR_BIT)) + { + using U = cuda::std::make_unsigned_t; + assert(result == static_cast(static_cast(v) << cuda::uabs(shift))); + } + else + { + assert(result == T{0}); + } +} + +template +TEST_FUNC constexpr void test() +{ + constexpr auto tmin = cuda::std::numeric_limits::min(); + constexpr auto tmax = cuda::std::numeric_limits::max(); + constexpr auto smin = cuda::std::numeric_limits::min(); + constexpr auto smax = cuda::std::numeric_limits::max(); + + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::std::shr(T{}, S{}))); + + const T vs[] = {tmin, T(-24), T(-1), T{0}, T{1}, T{20}, T{99}, T(1225), tmax}; + const S pos_shifts[] = {S{0}, S{1}, S{7}, S{17}, S{23}, S{32}, S{33}, S{65}, smax}; + + // Disable loop unrolling to reduce ptxas compile times. + _CCCL_PRAGMA_NOUNROLL() + for (auto v : vs) + { + _CCCL_PRAGMA_NOUNROLL() + for (auto shift : pos_shifts) + { + test_pos(v, shift); + } + } + + if constexpr (cuda::std::is_signed_v) + { + const S neg_shifts[] = {smin, S{-65}, S{-33}, S{-32}, S{-15}, S{-7}, S{-1}}; + + _CCCL_PRAGMA_NOUNROLL() + for (auto v : vs) + { + _CCCL_PRAGMA_NOUNROLL() + for (auto shift : neg_shifts) + { + test_neg(v, shift); + } + } + } +} + +template +TEST_FUNC constexpr void test() +{ + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test(); +#endif // _CCCL_HAS_INT128() + + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test(); +#endif // _CCCL_HAS_INT128() +} + +TEST_FUNC constexpr bool test() +{ + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test<__int128_t>(); +#endif // _CCCL_HAS_INT128() + + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test<__uint128_t>(); +#endif // _CCCL_HAS_INT128() + + return true; +} + +int main(int, char**) +{ + test(); + static_assert(test()); + + return 0; +} diff --git a/libcudacxx/test/support/test_macros.h b/libcudacxx/test/support/test_macros.h index 6c7fbaaf4aa..257ee9195b1 100644 --- a/libcudacxx/test/support/test_macros.h +++ b/libcudacxx/test/support/test_macros.h @@ -122,12 +122,15 @@ TEST_FUNC inline void DoNotOptimize(Tp const& value) } template -TEST_FUNC inline void DoNotOptimize(Tp& value) -{ +TEST_FUNC inline void DoNotOptimize(Tp& value){ + NV_IF_TARGET(NV_IS_DEVICE, ({ asm volatile("" ::"l"(&value) + : "memory"); })) # if TEST_COMPILER(CLANG) - asm volatile("" : "+r,m"(value) : : "memory"); + NV_IF_TARGET(NV_IS_HOST, ({ + asm volatile("" : "+r,m"(value) : : "memory"); + })) # else - asm volatile("" : "+m,r"(value) : : "memory"); + NV_IF_TARGET(NV_IS_HOST, ({ asm volatile("" : "+m,r"(value) : : "memory"); })) # endif } #endif // !TEST_COMPILER(MSVC) From d62708c8d0b7522165e9355808ce73f95da27d1d Mon Sep 17 00:00:00 2001 From: David Bayer Date: Fri, 17 Jul 2026 19:04:23 +0200 Subject: [PATCH 2/4] fixes --- .../test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp | 5 +++-- .../test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp | 5 +++-- 2 files changed, 6 insertions(+), 4 deletions(-) diff --git a/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp index c3525387de3..663c1169d62 100644 --- a/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shl.pass.cpp @@ -24,7 +24,7 @@ template TEST_FUNC constexpr T invoke_shl(T v, S shift) { - if (cuda::std::is_constant_evaluated()) + if (cuda::std::__cccl_default_is_constant_evaluated()) { return cuda::std::shl(v, shift); } @@ -74,7 +74,6 @@ TEST_FUNC constexpr void test() { constexpr auto tmin = cuda::std::numeric_limits::min(); constexpr auto tmax = cuda::std::numeric_limits::max(); - constexpr auto smin = cuda::std::numeric_limits::min(); constexpr auto smax = cuda::std::numeric_limits::max(); static_assert(cuda::std::is_same_v); @@ -96,6 +95,8 @@ TEST_FUNC constexpr void test() if constexpr (cuda::std::is_signed_v) { + constexpr auto smin = cuda::std::numeric_limits::min(); + const S neg_shifts[] = {smin, S{-65}, S{-33}, S{-32}, S{-15}, S{-7}, S{-1}}; _CCCL_PRAGMA_NOUNROLL() diff --git a/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp index 83887266103..67b576b56c1 100644 --- a/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp +++ b/libcudacxx/test/libcudacxx/std/numerics/bit/bit.shift/shr.pass.cpp @@ -24,7 +24,7 @@ template TEST_FUNC constexpr T invoke_shr(T v, S shift) { - if (cuda::std::is_constant_evaluated()) + if (cuda::std::__cccl_default_is_constant_evaluated()) { return cuda::std::shr(v, shift); } @@ -86,7 +86,6 @@ TEST_FUNC constexpr void test() { constexpr auto tmin = cuda::std::numeric_limits::min(); constexpr auto tmax = cuda::std::numeric_limits::max(); - constexpr auto smin = cuda::std::numeric_limits::min(); constexpr auto smax = cuda::std::numeric_limits::max(); static_assert(cuda::std::is_same_v); @@ -108,6 +107,8 @@ TEST_FUNC constexpr void test() if constexpr (cuda::std::is_signed_v) { + constexpr auto smin = cuda::std::numeric_limits::min(); + const S neg_shifts[] = {smin, S{-65}, S{-33}, S{-32}, S{-15}, S{-7}, S{-1}}; _CCCL_PRAGMA_NOUNROLL() From 0aed510f29fb3c7944d1e55acf604538b6f0f218 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Fri, 17 Jul 2026 19:45:08 +0200 Subject: [PATCH 3/4] fix donotoptimize --- libcudacxx/test/support/test_macros.h | 35 +++++++++------------------ 1 file changed, 12 insertions(+), 23 deletions(-) diff --git a/libcudacxx/test/support/test_macros.h b/libcudacxx/test/support/test_macros.h index 257ee9195b1..05f5c89402c 100644 --- a/libcudacxx/test/support/test_macros.h +++ b/libcudacxx/test/support/test_macros.h @@ -106,34 +106,23 @@ # define TEST_NVRTC_VIRTUAL_DEFAULT_DTOR_ANNOTATION #endif -#if TEST_COMPILER(MSVC) -# include template -inline void DoNotOptimize(Tp const& value) +TEST_FUNC inline void DoNotOptimize(Tp& value) { - [[maybe_unused]] const volatile void* volatile unused = __builtin_addressof(value); - _ReadWriteBarrier(); -} -#else // ^^^ TEST_COMPILER(MSVC) ^^^ / vvv !TEST_COMPILER(MSVC) vvv -template -TEST_FUNC inline void DoNotOptimize(Tp const& value) -{ - asm volatile("" : : "r,m"(value) : "memory"); -} + [[maybe_unused]] const volatile void* volatile ptr = &reinterpret_cast(__x); -template -TEST_FUNC inline void DoNotOptimize(Tp& value){ - NV_IF_TARGET(NV_IS_DEVICE, ({ asm volatile("" ::"l"(&value) - : "memory"); })) -# if TEST_COMPILER(CLANG) - NV_IF_TARGET(NV_IS_HOST, ({ - asm volatile("" : "+r,m"(value) : : "memory"); - })) -# else + // Device path. + NV_IF_TARGET(NV_IS_DEVICE, ({ asm volatile("" ::"l"(ptr) : "memory"); })) + + // Host path. +#if TEST_COMPILER(CLANG) + NV_IF_TARGET(NV_IS_HOST, ({ asm volatile("" : "+r,m"(value) : : "memory"); })) +#elif TEST_COMPILER(MSVC) + NV_IF_TARGET(NV_IS_HOST, ({ _ReadWriteBarrier(); })) +#else NV_IF_TARGET(NV_IS_HOST, ({ asm volatile("" : "+m,r"(value) : : "memory"); })) -# endif +#endif } -#endif // !TEST_COMPILER(MSVC) // NVCC can't handle static member variables, so with a little care // a function returning a reference will result in the same thing From bd57e62030efb7b34637874a5a2c27914a6cb8c0 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Fri, 17 Jul 2026 19:48:14 +0200 Subject: [PATCH 4/4] fix typo --- libcudacxx/test/support/test_macros.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/test/support/test_macros.h b/libcudacxx/test/support/test_macros.h index 05f5c89402c..90e651d91f5 100644 --- a/libcudacxx/test/support/test_macros.h +++ b/libcudacxx/test/support/test_macros.h @@ -109,7 +109,7 @@ template TEST_FUNC inline void DoNotOptimize(Tp& value) { - [[maybe_unused]] const volatile void* volatile ptr = &reinterpret_cast(__x); + [[maybe_unused]] const volatile void* volatile ptr = &reinterpret_cast(value); // Device path. NV_IF_TARGET(NV_IS_DEVICE, ({ asm volatile("" ::"l"(ptr) : "memory"); }))