arrive_tx: Trap on older architectures

NVIDIA · Aug 18, 2023 · 904ab27 · 904ab27
1 parent 701175b
commit 904ab27
Showing 1 changed file with 2 additions and 4 deletions.
diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h b/libcudacxx/include/cuda/std/detail/libcxx/include/__cuda/barrier.h
@@ -161,7 +161,7 @@ friend class _CUDA_VSTD::__barrier_poll_tester_parity;
         )
     }
 
-    _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_INLINE_VISIBILITY
+    _LIBCUDACXX_NODISCARD_ATTRIBUTE _LIBCUDACXX_DEVICE
     arrival_token arrive_tx(_CUDA_VSTD::ptrdiff_t __arrive_count_update,
                             _CUDA_VSTD::ptrdiff_t __transaction_count_update) {
 #if (_LIBCUDACXX_DEBUG_LEVEL >= 2)
@@ -186,9 +186,7 @@ friend class _CUDA_VSTD::__barrier_poll_tester_parity;
                     , "r"(static_cast<_CUDA_VSTD::uint32_t>(__transaction_count_update))
                     : "memory");
             ), NV_IS_DEVICE, (
-                static_assert("barrier::arrive_tx is only supported for sm_90 and up.");
-            ), NV_IS_HOST, (
-                static_assert("barrier::arrive_tx is only supported on device code.");
+                __trap(); // mbarrier.arrive.expect_tx is not supported for architectures before SM90.
             )
         )
         return __token;