Always bypass automatic atomic storage checks to prevent potential compiler issues (#4586)

* Always bypass automatic atomic storage checks to prevent potential compiler issues

* Waive atomic.local.pass.cpp if `_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE` is set.

---------

Co-authored-by: Wesley Maxey <[email protected]>
(cherry picked from commit 549dd45)

Author: 2 people

libcudacxx/include/cuda/std/detail/libcxx/include/__config

@@ -257,10 +257,11 @@ using char32_t = unsigned int;
 #    define _LIBCUDACXX_HAS_MSVC_ATOMIC_IMPL
 #  endif
 
+#  define _LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE
 // Enable bypassing automatic storage checks in atomics when using CTK 12.2 and below and if NDEBUG is defined.
-#  if _CCCL_CUDACC_BELOW(12, 2) && !defined(NDEBUG)
-#    define _LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE
-#  endif // _CCCL_CUDACC_BELOW(12, 2)
+// #  if _CCCL_CUDACC_BELOW(12, 2) && !defined(NDEBUG)
+// #    define _LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE
+// #  endif // _CCCL_CUDACC_BELOW(12, 2)
 
 // CUDA Atomics supersede host atomics in order to insert the host/device dispatch layer
 #  if _CCCL_CUDA_COMPILER(NVCC) || _CCCL_COMPILER(NVRTC) || _CCCL_COMPILER(NVHPC) || _CCCL_HAS_CUDA_COMPILER()

libcudacxx/test/libcudacxx/cuda/atomics/atomic.local.pass.cpp

@@ -120,11 +120,13 @@ __device__ inline void tests()
 
 int main(int arg, char** argv)
 {
+#if !defined(_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE)
   NV_IF_ELSE_TARGET(
     NV_IS_HOST,
     (cuda_thread_count = 64;),
     (tests<uint8_t>(); tests<uint16_t>(); tests<uint32_t>(); tests<uint64_t>(); tests<int8_t>(); tests<int16_t>();
      tests<int32_t>();
      tests<int64_t>();))
+#endif
   return 0;
 }