Atomics backend refactor

libcudacxx/codegen/CMakeLists.txt

@@ -19,8 +19,8 @@ target_compile_features(
 
 add_dependencies(libcudacxx.atomics.codegen codegen)
 
-set(atomic_generated_output "${libcudacxx_BINARY_DIR}/codegen/atomic_cuda_generated.h")
-set(atomic_install_location "${libcudacxx_SOURCE_DIR}/include/cuda/std/detail/libcxx/include/support/atomic")
+set(atomic_generated_output "${libcudacxx_BINARY_DIR}/codegen/cuda_ptx_generated.h")
+set(atomic_install_location "${libcudacxx_SOURCE_DIR}/include/cuda/std/__atomic/functions")
 
 add_custom_target(
     libcudacxx.atomics.codegen.execute
@@ -32,13 +32,13 @@ add_dependencies(libcudacxx.atomics.codegen libcudacxx.atomics.codegen.execute)
 
 add_custom_target(
     libcudacxx.atomics.codegen.install
-    COMMAND ${CMAKE_COMMAND} -E copy "${atomic_generated_output}" "${atomic_install_location}/atomic_cuda_generated.h"
-    BYPRODUCTS "${atomic_install_location}/atomic_cuda_generated.h"
+    COMMAND ${CMAKE_COMMAND} -E copy "${atomic_generated_output}" "${atomic_install_location}/cuda_ptx_generated.h"
+    BYPRODUCTS "${atomic_install_location}/cuda_ptx_generated.h"
 )
 
 add_dependencies(libcudacxx.atomics.codegen.install libcudacxx.atomics.codegen.execute)
 
 add_test(
     NAME libcudacxx.atomics.codegen.diff
-    COMMAND ${CMAKE_COMMAND} -E compare_files "${atomic_install_location}/atomic_cuda_generated.h" "${atomic_generated_output}"
+    COMMAND ${CMAKE_COMMAND} -E compare_files "${atomic_install_location}/cuda_ptx_generated.h" "${atomic_generated_output}"
 )

libcudacxx/codegen/codegen.cpp

@@ -66,7 +66,7 @@ int main()
 
   std::vector<std::string> cv_qualifier{"volatile ", ""};
 
-  std::ofstream out("atomic_cuda_generated.h");
+  std::ofstream out("cuda_ptx_generated.h");
 
   out << R"XXX(//===----------------------------------------------------------------------===//
 //
@@ -78,8 +78,35 @@ int main()
 //
 //===----------------------------------------------------------------------===//
 
-// This is a autogenerated file, we want to ensure that it contains exactly the contentes we want to generate
+// This is an autogenerated file, we want to ensure that it contains exactly the contents we want to generate
 // clang-format off
+
+#ifndef _LIBCUDACXX___ATOMIC_FUNCTIONS_CUDA_PTX_GENERATED_H
+#define _LIBCUDACXX___ATOMIC_FUNCTIONS_CUDA_PTX_GENERATED_H
+
+#include <cuda/std/detail/__config>
+
+#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 <cuda/std/cstdint>
+
+#include <cuda/std/__type_traits/enable_if.h>
+#include <cuda/std/__type_traits/is_signed.h>
+#include <cuda/std/__type_traits/is_unsigned.h>
+
+#include <cuda/std/__atomic/scopes.h>
+#include <cuda/std/__atomic/order.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+#if defined(_CCCL_CUDA_COMPILER)
+
 )XXX";
 
   auto scopenametag = [&](auto scope) {
@@ -302,11 +329,11 @@ int main()
               {
                 out << "template<class _Type, _CUDA_VSTD::__enable_if_t<sizeof(_Type)==" << sz / 8 << ", int> = 0>\n";
                 out << "_CCCL_DEVICE bool __atomic_compare_exchange_cuda(" << cv
-                    << "_Type *__ptr, _Type *__expected, const _Type *__desired, bool, int __success_memorder, int "
+                    << "_Type *__ptr, _Type *__expected, const _Type __desired, bool, int __success_memorder, int "
                        "__failure_memorder, "
                     << scopenametag(s.first) << ") {\n";
                 out << "    uint" << sz << "_t __tmp = 0, __old = 0, __old_tmp;\n";
-                out << "    memcpy(&__tmp, __desired, " << sz / 8 << ");\n";
+                out << "    memcpy(&__tmp, &__desired, " << sz / 8 << ");\n";
                 out << "    memcpy(&__old, __expected, " << sz / 8 << ");\n";
                 out << "    __old_tmp = __old;\n";
                 out << "    NV_DISPATCH_TARGET(\n";
@@ -503,6 +530,9 @@ int main()
     }
   }
 
+  out << "\n#endif // defined(_CCCL_CUDA_COMPILER)\n";
+  out << "\n_LIBCUDACXX_END_NAMESPACE_STD\n";
+  out << "\n#endif // _LIBCUDACXX___ATOMIC_FUNCTIONS_CUDA_PTX_GENERATED_H\n";
   out << "\n// clang-format on\n";
 
   return 0;

libcudacxx/examples/rtc_example.cpp

@@ -50,11 +50,11 @@ template<class T> static constexpr T min(T a, T b) { return a < b ? a : b; }
 
 struct trie {
     struct ref {
-        cuda::std::atomic<trie*> ptr = ATOMIC_VAR_INIT(nullptr);
+        cuda::std::atomic<trie*> ptr = LIBCUDACXX_ATOMIC_VAR_INIT(nullptr);
         // the flag will protect against multiple pointer updates
-        cuda::std::atomic_flag flag = ATOMIC_FLAG_INIT;
+        cuda::std::atomic_flag flag = LIBCUDACXX_ATOMIC_FLAG_INIT;
     } next[26];
-    cuda::std::atomic<int> count = ATOMIC_VAR_INIT(0);
+    cuda::std::atomic<int> count = LIBCUDACXX_ATOMIC_VAR_INIT(0);
 };
 __host__ __device__
 int index_of(char c) {

libcudacxx/examples/trie.cu

@@ -36,11 +36,11 @@ struct trie
 {
   struct ref
   {
-    cuda::atomic<trie*, cuda::thread_scope_device> ptr = ATOMIC_VAR_INIT(nullptr);
+    cuda::atomic<trie*, cuda::thread_scope_device> ptr = LIBCUDACXX_ATOMIC_VAR_INIT(nullptr);
     // the flag will protect against multiple pointer updates
-    cuda::std::atomic_flag flag = ATOMIC_FLAG_INIT;
+    cuda::std::atomic_flag flag = LIBCUDACXX_ATOMIC_FLAG_INIT;
   } next[26];
-  cuda::std::atomic<short> count = ATOMIC_VAR_INIT(0);
+  cuda::std::atomic<short> count = LIBCUDACXX_ATOMIC_VAR_INIT(0);
 };
 __host__ __device__ int index_of(char c)
 {

libcudacxx/examples/trie_mt.cpp

@@ -36,11 +36,11 @@ struct trie
 {
   struct ref
   {
-    std::atomic<trie*> ptr = ATOMIC_VAR_INIT(nullptr);
+    std::atomic<trie*> ptr = LIBCUDACXX_ATOMIC_VAR_INIT(nullptr);
     // the flag will protect against multiple pointer updates
-    std::atomic_flag flag = ATOMIC_VAR_INIT(0);
+    std::atomic_flag flag = LIBCUDACXX_ATOMIC_VAR_INIT(0);
   } next[26];
-  std::atomic<int> count = ATOMIC_VAR_INIT(0);
+  std::atomic<int> count = LIBCUDACXX_ATOMIC_VAR_INIT(0);
 };
 int index_of(char c)
 {

libcudacxx/include/cuda/atomic

@@ -11,6 +11,6 @@
 #ifndef _CUDA_ATOMIC
 #define _CUDA_ATOMIC
 
-#include <cuda/std/atomic>
+#include <cuda/std/__cuda/atomic.h>
 
 #endif // _CUDA_ATOMIC