Add atomic test to check copy ctor state and modify atomic/atomic_ref…

… ctor to match specification

.upstream-tests/test/std/atomics/atomics.types.generic/atomic_copyable.pass.cpp

@@ -0,0 +1,71 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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
+//
+//===----------------------------------------------------------------------===//
+//
+// UNSUPPORTED: libcpp-has-no-threads, pre-sm-60
+// UNSUPPORTED: windows && pre-sm-70
+
+// NOTE: atomic<> of a TriviallyCopyable class is wrongly rejected by older
+// clang versions. It was fixed right before the llvm 3.5 release. See PR18097.
+// XFAIL: apple-clang-6.0, clang-3.4, clang-3.3
+
+// <cuda/std/atomic>
+
+#include <cuda/std/atomic>
+#include <cuda/std/cassert>
+// #include <cuda/std/thread> // for thread_id
+// #include <cuda/std/chrono> // for nanoseconds
+
+#include "test_macros.h"
+
+template <class T>
+__host__ __device__
+void test_not_copy_constructible() {
+  static_assert(!cuda::std::is_copy_constructible<T>(), "");
+}
+
+template <class T>
+__host__ __device__
+void test_copy_constructible() {
+  static_assert(cuda::std::is_copy_constructible<T>(), "");
+}
+
+template <class T, class A>
+__host__ __device__
+void test_atomic_ref_copy_ctor() {
+  A val = 0;
+
+  T t0(val);
+  T t1(t0);
+
+  t0++;
+  t1++;
+
+  assert(t1.load() == 2);
+}
+
+int main(int, char**)
+{
+    test_not_copy_constructible<cuda::std::atomic<int>>();
+    test_not_copy_constructible<cuda::atomic<int>>();
+    test_not_copy_constructible<const cuda::std::atomic<int>>();
+    test_not_copy_constructible<const cuda::atomic<int>>();
+
+    test_copy_constructible<cuda::std::atomic_ref<int>>();
+    test_copy_constructible<cuda::atomic_ref<int>>();
+    test_copy_constructible<const cuda::std::atomic_ref<int>>();
+    test_copy_constructible<const cuda::atomic_ref<int>>();
+
+    test_atomic_ref_copy_ctor<cuda::std::atomic_ref<int>, int>();
+    test_atomic_ref_copy_ctor<cuda::atomic_ref<int>, int>();
+    test_atomic_ref_copy_ctor<const cuda::std::atomic_ref<int>, int>();
+    test_atomic_ref_copy_ctor<const cuda::atomic_ref<int>, int>();
+    // test(cuda::std::this_thread::get_id());
+    // test(cuda::std::chrono::nanoseconds(2));
+
+  return 0;
+}

include/cuda/std/detail/libcxx/include/atomic

@@ -1245,6 +1245,7 @@ struct __atomic_base {
     mutable __cxx_atomic_impl<_Tp, _Sco> __a_;
 
     __atomic_base() = default;
+    __atomic_base(const __atomic_base&) = delete;
     _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR
     __atomic_base(const _Tp& __a) _NOEXCEPT : __a_(__a) {}
 
@@ -1364,6 +1365,7 @@ struct __atomic_base_ref {
     mutable __cxx_atomic_ref_impl<_Tp, _Sco> __a_;
 
     __atomic_base_ref() = default;
+    __atomic_base_ref(const __atomic_base_ref&) = default;
     _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR
     __atomic_base_ref(_Tp& __a) _NOEXCEPT : __a_(__a) {}
 
@@ -1483,6 +1485,7 @@ struct __atomic_base_ref {
 template <class _Tp, int _Sco>
 struct __atomic_base<_Tp, _Sco, true> : public __atomic_base<_Tp, _Sco, false> {
     __atomic_base() = default;
+    __atomic_base(const __atomic_base&) = delete;
 
     _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR
     __atomic_base(const _Tp& __a) _NOEXCEPT : __atomic_base<_Tp, _Sco, false>(__a) {}
@@ -1559,6 +1562,7 @@ struct __atomic_base<_Tp, _Sco, true> : public __atomic_base<_Tp, _Sco, false> {
 template <class _Tp, int _Sco>
 struct __atomic_base_ref<_Tp, _Sco, true> : public __atomic_base_ref<_Tp, _Sco, false> {
     __atomic_base_ref() = default;
+    __atomic_base_ref(const __atomic_base_ref&) = default;
 
     _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR
     __atomic_base_ref(_Tp& __a) _NOEXCEPT : __atomic_base_ref<_Tp, _Sco, false>(__a) {}