Fix cuda::barrier missing accounting of results in try_wait (#7538) (#7635)

* Fix cuda::barrier missing accounting of results in try_wait

* Cleanup redundant return

* Apply suggestion from @fbusato



---------


(cherry picked from commit ce23268)

Co-authored-by: Wesley Maxey <71408887+wmaxey@users.noreply.github.com>
Co-authored-by: Federico Busato <50413820+fbusato@users.noreply.github.com>
Co-authored-by: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>

Author: 4 people

libcudacxx/include/cuda/__barrier/barrier_block_scope.h

@@ -243,11 +243,11 @@ class barrier<thread_scope_block, ::cuda::std::__empty_completion> : public __bl
       bool __ready = 0;
       ::cuda::std::chrono::high_resolution_clock::time_point const __start =
         ::cuda::std::chrono::high_resolution_clock::now();
-      ::cuda::std::chrono::nanoseconds __elapsed;
+      ::cuda::std::chrono::nanoseconds __elapsed(0);
       do
       {
         const ::cuda::std::uint32_t __wait_nsec = static_cast<::cuda::std::uint32_t>((__nanosec - __elapsed).count());
-        ::cuda::ptx::mbarrier_try_wait(__native_handle(), __token, __wait_nsec);
+        __ready   = ::cuda::ptx::mbarrier_try_wait(__native_handle(), __token, __wait_nsec);
         __elapsed = ::cuda::std::chrono::high_resolution_clock::now() - __start;
       } while (!__ready && (__nanosec > __elapsed));
       return __ready;
@@ -343,11 +343,11 @@ class barrier<thread_scope_block, ::cuda::std::__empty_completion> : public __bl
       int32_t __ready = 0;
       ::cuda::std::chrono::high_resolution_clock::time_point const __start =
         ::cuda::std::chrono::high_resolution_clock::now();
-      ::cuda::std::chrono::nanoseconds __elapsed;
+      ::cuda::std::chrono::nanoseconds __elapsed(0);
       do
       {
         const ::cuda::std::uint32_t __wait_nsec = static_cast<::cuda::std::uint32_t>((__nanosec - __elapsed).count());
-        ::cuda::ptx::mbarrier_try_wait_parity(__native_handle(), __phase_parity, __wait_nsec);
+        __ready   = ::cuda::ptx::mbarrier_try_wait_parity(__native_handle(), __phase_parity, __wait_nsec);
         __elapsed = ::cuda::std::chrono::high_resolution_clock::now() - __start;
       } while (!__ready && (__nanosec > __elapsed));
 
@@ -397,6 +397,7 @@ class barrier<thread_scope_block, ::cuda::std::__empty_completion> : public __bl
       NV_ANY_TARGET,
       (return ::cuda::std::__cccl_thread_poll_with_backoff(
                 ::cuda::std::__barrier_poll_tester_parity<barrier>(this, __phase_parity), __nanosec);))
+    _CCCL_UNREACHABLE();
   }
 
 public:

libcudacxx/test/libcudacxx/force_include.h

@@ -77,9 +77,12 @@ int main(int argc, char** argv)
   int ret = fake_main(argc, argv);
   if (ret != 0)
   {
+    printf("Host testing returned failure\n");
     return ret;
   }
 
+  printf("Testing on device:\n");
+  fflush(stdout);
   list_devices();
   int* cuda_ret = 0;
   CUDA_CALL(err, cudaMalloc(&cuda_ret, sizeof(int)));
@@ -111,7 +114,12 @@ int main(int argc, char** argv)
   CUDA_CALL(err, cudaDeviceSynchronize());
   CUDA_CALL(err, cudaMemcpy(&ret, cuda_ret, sizeof(int), cudaMemcpyDeviceToHost));
   CUDA_CALL(err, cudaFree(cuda_ret));
+  fflush(stdout);
 
+  if (ret != 0)
+  {
+    printf("Device testing returned failure\n");
+  }
   return ret;
 }
 

libcudacxx/test/libcudacxx/std/thread/thread.barrier/try_wait_for.pass.cpp

@@ -18,57 +18,84 @@
 #include "test_macros.h"
 
 template <typename Barrier, template <typename, typename> class Selector, typename Initializer = constructor_initializer>
-__host__ __device__ void test(bool add_delay = false)
+__host__ __device__ int test(bool add_delay = false)
 {
+  printf("delay %s\r\n", add_delay ? "enabled" : "disabled");
+
   Selector<Barrier, Initializer> sel;
   SHARED Barrier* b;
-  b          = sel.construct(2);
-  auto delay = cuda::std::chrono::duration<int>(0);
+  b            = sel.construct(2);
+  auto delay   = cuda::std::chrono::nanoseconds(0);
+  auto timeout = cuda::std::chrono::nanoseconds(100000000);
 
   if (add_delay)
   {
-    delay = cuda::std::chrono::duration<int>(1);
+    delay = cuda::std::chrono::nanoseconds(100000);
   }
 
-  typename Barrier::arrival_token* tok = nullptr;
-  execute_on_main_thread([&] {
-    tok = new auto(b->arrive());
-  });
+  auto time = cuda::std::chrono::high_resolution_clock::now();
+  cuda::std::atomic_ref<decltype(time)> time_ref(time);
 
-  auto awaiter = LAMBDA()
+  auto measure = LAMBDA()->cuda::std::chrono::nanoseconds
   {
-    while (b->try_wait_for(cuda::std::move(*tok), delay) == false)
-    {
-    }
+    return cuda::std::chrono::duration_cast<cuda::std::chrono::nanoseconds>(
+      cuda::std::chrono::high_resolution_clock::now() - time_ref.load());
   };
-  auto arriver = LAMBDA()
+
   {
-    (void) b->arrive();
-  };
-  concurrent_agents_launch(awaiter, arriver);
+    typename Barrier::arrival_token* tok = nullptr;
+    execute_on_main_thread([&] {
+      tok = new auto(b->arrive());
+    });
 
-  execute_on_main_thread([&] {
-    auto tok2 = b->arrive(2);
-    while (b->try_wait_for(cuda::std::move(tok2), delay) == false)
+    auto awaiter = LAMBDA()
+    {
+      time_ref = cuda::std::chrono::high_resolution_clock::now();
+      while ((b->try_wait_for(cuda::std::move(*tok), delay) == false) && (measure() < timeout))
+      {
+      }
+      printf("p1 barrier delay: %lluns\r\n", measure().count());
+    };
+    auto arriver = LAMBDA()
     {
+      (void) b->arrive();
+    };
+    concurrent_agents_launch(awaiter, arriver);
+    if (measure() > timeout)
+    {
+      printf("Deadlock detected in p1\r\n");
+      return 1;
     }
-  });
+  }
+  {
+    execute_on_main_thread([&] {
+      auto tok2 = b->arrive(2);
+      time_ref  = ::cuda::std::chrono::high_resolution_clock::now();
+      while ((b->try_wait_for(cuda::std::move(tok2), delay) == false) && (measure() < timeout))
+      {
+      }
+      printf("p2 barrier delay: %lluns\r\n", measure().count());
+    });
+    if (measure() > timeout)
+    {
+      printf("Deadlock detected in p2\r\n");
+      return 1;
+    }
+  }
+  return 0;
 }
 
 int main(int, char**)
 {
-  NV_IF_ELSE_TARGET(
+  int failure = 0;
+  NV_IF_TARGET(
     NV_IS_HOST,
-    (
-      // Required by concurrent_agents_launch to know how many we're launching
-      cuda_thread_count = 2;
-
-      test<cuda::barrier<cuda::thread_scope_block>, local_memory_selector>();
-      test<cuda::barrier<cuda::thread_scope_block>, local_memory_selector>(true);),
-    (test<cuda::barrier<cuda::thread_scope_block>, shared_memory_selector>();
-     test<cuda::barrier<cuda::thread_scope_block>, global_memory_selector>();
-     test<cuda::barrier<cuda::thread_scope_block>, shared_memory_selector>(true);
-     test<cuda::barrier<cuda::thread_scope_block>, global_memory_selector>(true);))
+    (cuda_thread_count = 2; failure |= test<cuda::barrier<cuda::thread_scope_block>, local_memory_selector>();
+     failure |= test<cuda::barrier<cuda::thread_scope_block>, local_memory_selector>(true);),
+    (failure |= test<cuda::barrier<cuda::thread_scope_block>, shared_memory_selector>();
+     failure |= test<cuda::barrier<cuda::thread_scope_block>, global_memory_selector>();
+     failure |= test<cuda::barrier<cuda::thread_scope_block>, shared_memory_selector>(true);
+     failure |= test<cuda::barrier<cuda::thread_scope_block>, global_memory_selector>(true);))
 
-  return 0;
+  return failure;
 }