Fix flakey heterogeneous tests

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,15 @@ 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 +195,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,50 @@ 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

@@ -13,12 +13,10 @@
 // #define DEBUG_TESTERS
 
 #include <cuda/barrier>
+#include <cuda/std/cassert>
 
 #include "helpers.h"
 
-__managed__ bool completed_from_host   = false;
-__managed__ bool completed_from_device = false;
-
 template <typename Barrier>
 struct barrier_and_token
 {
@@ -46,8 +44,6 @@ struct barrier_and_token_with_completion
     {
       assert(completed.load() == false);
       completed.store(true);
-
-      NV_IF_ELSE_TARGET(NV_IS_HOST, completed_from_host = true;, completed_from_device = true;)
     }
   };
 
@@ -148,32 +144,28 @@ 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)
 {
-  NV_IF_TARGET(
-    NV_IS_HOST,
-    (kernel_invoker();
-
-     if (check_managed_memory_support(true)) { assert(completed_from_host); } assert(completed_from_device);))
+  NV_IF_TARGET(NV_IS_HOST, (kernel_invoker();))
 
   return 0;
 }

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

@@ -15,15 +15,13 @@
 #define _LIBCUDACXX_CUDA_ABI_VERSION 2
 
 #include <cuda/barrier>
+#include <cuda/std/cassert>
 
 #include "helpers.h"
 
 static_assert(sizeof(cuda::barrier<cuda::thread_scope_system>) == 64, "");
 static_assert(sizeof(cuda::barrier<cuda::thread_scope_system, void (*)()>) == 4 * 64, "");
 
-__managed__ bool completed_from_host   = false;
-__managed__ bool completed_from_device = false;
-
 template <typename Barrier>
 struct barrier_and_token
 {
@@ -51,8 +49,6 @@ struct barrier_and_token_with_completion
     {
       assert(completed.load() == false);
       completed.store(true);
-
-      NV_IF_ELSE_TARGET(NV_IS_HOST, completed_from_host = true;, completed_from_device = true;)
     }
   };
 
@@ -153,32 +149,28 @@ 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)
 {
-  NV_IF_TARGET(
-    NV_IS_HOST,
-    (kernel_invoker();
-
-     if (check_managed_memory_support(true)) { assert(completed_from_host); } assert(completed_from_device);))
+  NV_IF_TARGET(NV_IS_HOST, (kernel_invoker();))
 
   return 0;
 }

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)