revert from wrong condition removal

NVIDIA · Jan 30, 2025 · 9f9f341 · 9f9f341
1 parent 427d7c8
commit 9f9f341
Show file tree

Hide file tree

Showing 2 changed files with 7 additions and 2 deletions.
diff --git 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<const void*>(__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<const void*>(__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
@@ -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__)