From 9f9f341561a897790a953126bb2b448162d855c3 Mon Sep 17 00:00:00 2001 From: fbusato Date: Thu, 30 Jan 2025 14:51:42 -0800 Subject: [PATCH] revert from wrong condition removal --- libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h | 4 ++-- libcudacxx/include/cuda/std/__cccl/builtin.h | 5 +++++ 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h b/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h index 7dbeb175e48..2ecd56daf55 100644 --- a/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h +++ b/libcudacxx/include/cuda/std/__atomic/functions/cuda_local.h @@ -39,7 +39,7 @@ _CCCL_DEVICE inline bool __cuda_is_local(const volatile void* __ptr) # if defined(_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE) return false; // Only NVCC+NVRTC define __isLocal, so drop to PTX -# elif _CCCL_CUDA_COMPILER(NVHPC) +# elif _CCCL_CUDACC_BELOW(12, 3) || _CCCL_CUDA_COMPILER(NVHPC) int __tmp = 0; asm("{\n\t" " .reg .pred p;\n\t" @@ -49,7 +49,7 @@ _CCCL_DEVICE inline bool __cuda_is_local(const volatile void* __ptr) : "=r"(__tmp) : "l"(const_cast(__ptr))); return __tmp == 1; -# else // ^^^ _CCCL_CUDA_COMPILER(NVHPC) ^^^ / vvv other compiler vvv +# else // ^^^ _CCCL_CUDACC_BELOW(12, 3) || _CCCL_CUDA_COMPILER(NVHPC) ^^^ / vvv other compiler vvv return __isLocal(const_cast(__ptr)); # endif // _CCCL_CUDACC_AT_LEAST(12, 3) && !_CCCL_CUDA_COMPILER(NVHPC) } diff --git a/libcudacxx/include/cuda/std/__cccl/builtin.h b/libcudacxx/include/cuda/std/__cccl/builtin.h index 0cbd73c9672..99c7f1f9ae1 100644 --- a/libcudacxx/include/cuda/std/__cccl/builtin.h +++ b/libcudacxx/include/cuda/std/__cccl/builtin.h @@ -892,6 +892,11 @@ # define _CCCL_BUILTIN_TYPE_PACK_ELEMENT(...) __type_pack_element<__VA_ARGS__> #endif // _CCCL_HAS_BUILTIN(__type_pack_element) +// NVCC prior to 12.2 have trouble with pack expansion into __type_pack_element in an alias template +#if _CCCL_CUDACC_BELOW(12, 2) +# undef _CCCL_BUILTIN_TYPE_PACK_ELEMENT +#endif // _CCCL_CUDACC_BELOW(12, 2) + #if _CCCL_CHECK_BUILTIN(underlying_type) || _CCCL_COMPILER(GCC, >=, 4, 7) || _CCCL_COMPILER(MSVC) \ || _CCCL_COMPILER(NVRTC) # define _CCCL_BUILTIN_UNDERLYING_TYPE(...) __underlying_type(__VA_ARGS__)