Change tests to use validate_pinned API.

libcudacxx/test/libcudacxx/heterogeneous/atomic.pass.cpp

@@ -10,6 +10,7 @@
 // UNSUPPORTED: windows && pre-sm-70
 
 #include <cuda/std/atomic>
+#include <cuda/std/cassert>
 
 #include "helpers.h"
 
@@ -144,18 +145,18 @@ using basic_testers =
               exchange_tester<-12, 17>>;
 
 using arithmetic_atomic_testers =
-  extend_tester_list<basic_testers,
-                     fetch_add_tester<17, 13, 30>,
-                     fetch_sub_tester<30, 21, 9>,
-                     fetch_sub_tester<9, 17, -8>>;
+  append<basic_testers,
+        fetch_add_tester<17, 13, 30>,
+        fetch_sub_tester<30, 21, 9>,
+        fetch_sub_tester<9, 17, -8>>;
 
 using bitwise_atomic_testers =
-  extend_tester_list<arithmetic_atomic_testers,
-                     fetch_add_tester<-8, 10, 2>,
-                     fetch_or_tester<2, 13, 15>,
-                     fetch_and_tester<15, 8, 8>,
-                     fetch_and_tester<8, 13, 8>,
-                     fetch_xor_tester<8, 12, 4>>;
+  append<arithmetic_atomic_testers,
+        fetch_add_tester<-8, 10, 2>,
+        fetch_or_tester<2, 13, 15>,
+        fetch_and_tester<15, 8, 8>,
+        fetch_and_tester<8, 13, 8>,
+        fetch_xor_tester<8, 12, 4>>;
 
 class big_not_lockfree_type
 {
@@ -197,39 +198,39 @@ __host__ __device__ void validate_not_lock_free()
 
 void kernel_invoker()
 {
-  validate_not_movable<cuda::std::atomic<signed char>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<signed short>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<signed int>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<signed long>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<signed long long>, arithmetic_atomic_testers>();
-
-  validate_not_movable<cuda::std::atomic<unsigned char>, bitwise_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<unsigned short>, bitwise_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<unsigned int>, bitwise_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<unsigned long>, bitwise_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<unsigned long long>, bitwise_atomic_testers>();
-
-  validate_not_movable<cuda::std::atomic<float>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::std::atomic<double>, arithmetic_atomic_testers>();
-
-  validate_not_movable<cuda::std::atomic<big_not_lockfree_type>, basic_testers>();
-
-  validate_not_movable<cuda::atomic<signed char, cuda::thread_scope_system>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::atomic<signed short, cuda::thread_scope_system>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::atomic<signed int, cuda::thread_scope_system>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::atomic<signed long, cuda::thread_scope_system>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::atomic<signed long long, cuda::thread_scope_system>, arithmetic_atomic_testers>();
-
-  validate_not_movable<cuda::atomic<unsigned char, cuda::thread_scope_system>, bitwise_atomic_testers>();
-  validate_not_movable<cuda::atomic<unsigned short, cuda::thread_scope_system>, bitwise_atomic_testers>();
-  validate_not_movable<cuda::atomic<unsigned int, cuda::thread_scope_system>, bitwise_atomic_testers>();
-  validate_not_movable<cuda::atomic<unsigned long, cuda::thread_scope_system>, bitwise_atomic_testers>();
-  validate_not_movable<cuda::atomic<unsigned long long, cuda::thread_scope_system>, bitwise_atomic_testers>();
-
-  validate_not_movable<cuda::atomic<float>, arithmetic_atomic_testers>();
-  validate_not_movable<cuda::atomic<double>, arithmetic_atomic_testers>();
-
-  validate_not_movable<cuda::atomic<big_not_lockfree_type, cuda::thread_scope_system>, basic_testers>();
+  validate_pinned<cuda::std::atomic<signed char>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::std::atomic<signed short>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::std::atomic<signed int>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::std::atomic<signed long>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::std::atomic<signed long long>, arithmetic_atomic_testers>();
+
+  validate_pinned<cuda::std::atomic<unsigned char>, bitwise_atomic_testers>();
+  validate_pinned<cuda::std::atomic<unsigned short>, bitwise_atomic_testers>();
+  validate_pinned<cuda::std::atomic<unsigned int>, bitwise_atomic_testers>();
+  validate_pinned<cuda::std::atomic<unsigned long>, bitwise_atomic_testers>();
+  validate_pinned<cuda::std::atomic<unsigned long long>, bitwise_atomic_testers>();
+
+  validate_pinned<cuda::std::atomic<float>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::std::atomic<double>, arithmetic_atomic_testers>();
+
+  validate_pinned<cuda::std::atomic<big_not_lockfree_type>, basic_testers>();
+
+  validate_pinned<cuda::atomic<signed char, cuda::thread_scope_system>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::atomic<signed short, cuda::thread_scope_system>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::atomic<signed int, cuda::thread_scope_system>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::atomic<signed long, cuda::thread_scope_system>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::atomic<signed long long, cuda::thread_scope_system>, arithmetic_atomic_testers>();
+
+  validate_pinned<cuda::atomic<unsigned char, cuda::thread_scope_system>, bitwise_atomic_testers>();
+  validate_pinned<cuda::atomic<unsigned short, cuda::thread_scope_system>, bitwise_atomic_testers>();
+  validate_pinned<cuda::atomic<unsigned int, cuda::thread_scope_system>, bitwise_atomic_testers>();
+  validate_pinned<cuda::atomic<unsigned long, cuda::thread_scope_system>, bitwise_atomic_testers>();
+  validate_pinned<cuda::atomic<unsigned long long, cuda::thread_scope_system>, bitwise_atomic_testers>();
+
+  validate_pinned<cuda::atomic<float>, arithmetic_atomic_testers>();
+  validate_pinned<cuda::atomic<double>, arithmetic_atomic_testers>();
+
+  validate_pinned<cuda::atomic<big_not_lockfree_type, cuda::thread_scope_system>, basic_testers>();
 }
 
 int main(int arg, char** argv)

libcudacxx/test/libcudacxx/heterogeneous/atomic_flag.pass.cpp

@@ -10,6 +10,7 @@
 // UNSUPPORTED: windows && pre-sm-70
 
 #include <cuda/std/atomic>
+#include <cuda/std/cassert>
 
 #include "helpers.h"
 
@@ -51,7 +52,7 @@ using atomic_flag_testers = tester_list<clear_tester, clear, test_and_set_tester
 
 void kernel_invoker()
 {
-  validate_not_movable<cuda::std::atomic_flag, atomic_flag_testers>();
+  validate_pinned<cuda::std::atomic_flag, atomic_flag_testers>();
 }
 
 int main(int argc, char** argv)

libcudacxx/test/libcudacxx/heterogeneous/atomic_ref.pass.cpp

@@ -10,6 +10,7 @@
 // UNSUPPORTED: windows && pre-sm-70
 
 #include <cuda/std/atomic>
+#include <cuda/std/cassert>
 
 #include "helpers.h"
 
@@ -154,53 +155,53 @@ using basic_testers =
               exchange_tester<-12, 17>>;
 
 using arithmetic_atomic_testers =
-  extend_tester_list<basic_testers,
-                     fetch_add_tester<17, 13, 30>,
-                     fetch_sub_tester<30, 21, 9>,
-                     fetch_sub_tester<9, 17, -8>>;
+  append<basic_testers,
+         fetch_add_tester<17, 13, 30>,
+         fetch_sub_tester<30, 21, 9>,
+         fetch_sub_tester<9, 17, -8>>;
 
 using bitwise_atomic_testers =
-  extend_tester_list<arithmetic_atomic_testers,
-                     fetch_add_tester<-8, 10, 2>,
-                     fetch_or_tester<2, 13, 15>,
-                     fetch_and_tester<15, 8, 8>,
-                     fetch_and_tester<8, 13, 8>,
-                     fetch_xor_tester<8, 12, 4>>;
+  append<arithmetic_atomic_testers,
+         fetch_add_tester<-8, 10, 2>,
+         fetch_or_tester<2, 13, 15>,
+         fetch_and_tester<15, 8, 8>,
+         fetch_and_tester<8, 13, 8>,
+         fetch_xor_tester<8, 12, 4>>;
 
 void kernel_invoker()
 {
 // todo
 #ifdef _LIBCUDACXX_ATOMIC_REF_SUPPORTS_SMALL_INTEGRAL
-  validate_not_movable<signed char, arithmetic_atomic_testers>();
-  validate_not_movable<signed short, arithmetic_atomic_testers>();
+  validate_pinned<signed char, arithmetic_atomic_testers>();
+  validate_pinned<signed short, arithmetic_atomic_testers>();
 #endif
-  validate_not_movable<signed int, arithmetic_atomic_testers>();
-  validate_not_movable<signed long, arithmetic_atomic_testers>();
-  validate_not_movable<signed long long, arithmetic_atomic_testers>();
+  validate_pinned<signed int, arithmetic_atomic_testers>();
+  validate_pinned<signed long, arithmetic_atomic_testers>();
+  validate_pinned<signed long long, arithmetic_atomic_testers>();
 
 #ifdef _LIBCUDACXX_ATOMIC_REF_SUPPORTS_SMALL_INTEGRAL
-  validate_not_movable<unsigned char, bitwise_atomic_testers>();
-  validate_not_movable<unsigned short, bitwise_atomic_testers>();
+  validate_pinned<unsigned char, bitwise_atomic_testers>();
+  validate_pinned<unsigned short, bitwise_atomic_testers>();
 #endif
-  validate_not_movable<unsigned int, bitwise_atomic_testers>();
-  validate_not_movable<unsigned long, bitwise_atomic_testers>();
-  validate_not_movable<unsigned long long, bitwise_atomic_testers>();
+  validate_pinned<unsigned int, bitwise_atomic_testers>();
+  validate_pinned<unsigned long, bitwise_atomic_testers>();
+  validate_pinned<unsigned long long, bitwise_atomic_testers>();
 
 #ifdef _LIBCUDACXX_ATOMIC_REF_SUPPORTS_SMALL_INTEGRAL
-  validate_not_movable<signed char, arithmetic_atomic_testers>();
-  validate_not_movable<signed short, arithmetic_atomic_testers>();
+  validate_pinned<signed char, arithmetic_atomic_testers>();
+  validate_pinned<signed short, arithmetic_atomic_testers>();
 #endif
-  validate_not_movable<signed int, arithmetic_atomic_testers>();
-  validate_not_movable<signed long, arithmetic_atomic_testers>();
-  validate_not_movable<signed long long, arithmetic_atomic_testers>();
+  validate_pinned<signed int, arithmetic_atomic_testers>();
+  validate_pinned<signed long, arithmetic_atomic_testers>();
+  validate_pinned<signed long long, arithmetic_atomic_testers>();
 
 #ifdef _LIBCUDACXX_ATOMIC_REF_SUPPORTS_SMALL_INTEGRAL
-  validate_not_movable<unsigned char, bitwise_atomic_testers>();
-  validate_not_movable<unsigned short, bitwise_atomic_testers>();
+  validate_pinned<unsigned char, bitwise_atomic_testers>();
+  validate_pinned<unsigned short, bitwise_atomic_testers>();
 #endif
-  validate_not_movable<unsigned int, bitwise_atomic_testers>();
-  validate_not_movable<unsigned long, bitwise_atomic_testers>();
-  validate_not_movable<unsigned long long, bitwise_atomic_testers>();
+  validate_pinned<unsigned int, bitwise_atomic_testers>();
+  validate_pinned<unsigned long, bitwise_atomic_testers>();
+  validate_pinned<unsigned long long, bitwise_atomic_testers>();
 }
 
 int main(int arg, char** argv)

libcudacxx/test/libcudacxx/heterogeneous/barrier.pass.cpp

@@ -12,6 +12,7 @@
 // uncomment for a really verbose output detailing what test steps are being launched
 // #define DEBUG_TESTERS
 
+#include <cuda/std/cassert>
 #include <cuda/barrier>
 
 #include "helpers.h"
@@ -148,23 +149,23 @@ using cuda_barrier_system = cuda::barrier<cuda::thread_scope_system, Completion>
 
 void kernel_invoker()
 {
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, a_aw_w>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_aw_w>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, a_aw_w>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_aw_w>(2);
 
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, aw_aw>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, aw_aw>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, aw_aw>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, aw_aw>(2);
 
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, a_w_aw>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_w_aw>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, a_w_aw>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_w_aw>(2);
 
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, a_w_a_w>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_w_a_w>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, a_w_a_w>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_w_a_w>(2);
 
-  validate_not_movable<barrier_and_token_with_completion<cuda::std::barrier>, completion_performers_a>(2);
-  validate_not_movable<barrier_and_token_with_completion<cuda_barrier_system>, completion_performers_a>(2);
+  validate_pinned<barrier_and_token_with_completion<cuda::std::barrier>, completion_performers_a>(2);
+  validate_pinned<barrier_and_token_with_completion<cuda_barrier_system>, completion_performers_a>(2);
 
-  validate_not_movable<barrier_and_token_with_completion<cuda::std::barrier>, completion_performers_b>(2);
-  validate_not_movable<barrier_and_token_with_completion<cuda_barrier_system>, completion_performers_b>(2);
+  validate_pinned<barrier_and_token_with_completion<cuda::std::barrier>, completion_performers_b>(2);
+  validate_pinned<barrier_and_token_with_completion<cuda_barrier_system>, completion_performers_b>(2);
 }
 
 int main(int arg, char** argv)

libcudacxx/test/libcudacxx/heterogeneous/barrier_abi_v2.pass.cpp

@@ -15,6 +15,7 @@
 #define _LIBCUDACXX_CUDA_ABI_VERSION 2
 
 #include <cuda/barrier>
+#include <cuda/std/cassert>
 
 #include "helpers.h"
 
@@ -153,23 +154,23 @@ using cuda_barrier_system = cuda::barrier<cuda::thread_scope_system, Completion>
 
 void kernel_invoker()
 {
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, a_aw_w>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_aw_w>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, a_aw_w>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_aw_w>(2);
 
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, aw_aw>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, aw_aw>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, aw_aw>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, aw_aw>(2);
 
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, a_w_aw>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_w_aw>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, a_w_aw>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_w_aw>(2);
 
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, a_w_a_w>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_w_a_w>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, a_w_a_w>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, a_w_a_w>(2);
 
-  validate_not_movable<barrier_and_token_with_completion<cuda::std::barrier>, completion_performers_a>(2);
-  validate_not_movable<barrier_and_token_with_completion<cuda_barrier_system>, completion_performers_a>(2);
+  validate_pinned<barrier_and_token_with_completion<cuda::std::barrier>, completion_performers_a>(2);
+  validate_pinned<barrier_and_token_with_completion<cuda_barrier_system>, completion_performers_a>(2);
 
-  validate_not_movable<barrier_and_token_with_completion<cuda::std::barrier>, completion_performers_b>(2);
-  validate_not_movable<barrier_and_token_with_completion<cuda_barrier_system>, completion_performers_b>(2);
+  validate_pinned<barrier_and_token_with_completion<cuda::std::barrier>, completion_performers_b>(2);
+  validate_pinned<barrier_and_token_with_completion<cuda_barrier_system>, completion_performers_b>(2);
 }
 
 int main(int arg, char** argv)

libcudacxx/test/libcudacxx/heterogeneous/barrier_parity.cuda.pass.cpp

@@ -14,6 +14,7 @@
 // #define DEBUG_TESTERS
 
 #include <cuda/barrier>
+#include <cuda/std/cassert>
 
 #include <atomic>
 
@@ -89,8 +90,8 @@ using aw_aw_pw2 =
 
 void kernel_invoker()
 {
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, aw_aw_pw1>(2);
-  validate_not_movable<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, aw_aw_pw2>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, aw_aw_pw1>(2);
+  validate_pinned<barrier_and_token<cuda::barrier<cuda::thread_scope_system>>, aw_aw_pw2>(2);
 }
 
 int main(int arg, char** argv)

libcudacxx/test/libcudacxx/heterogeneous/barrier_parity.std.pass.cpp

@@ -14,6 +14,7 @@
 // #define DEBUG_TESTERS
 
 #include <cuda/barrier>
+#include <cuda/std/cassert>
 
 #include <atomic>
 
@@ -88,8 +89,8 @@ using aw_aw_pw2 =
 
 void kernel_invoker()
 {
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, aw_aw_pw1>(2);
-  validate_not_movable<barrier_and_token<cuda::std::barrier<>>, aw_aw_pw2>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, aw_aw_pw1>(2);
+  validate_pinned<barrier_and_token<cuda::std::barrier<>>, aw_aw_pw2>(2);
 }
 
 int main(int arg, char** argv)