diff --git a/docs/extended_api/synchronization_primitives/barrier.md b/docs/extended_api/synchronization_primitives/barrier.md index 17ad379fc6..e10d95effb 100644 --- a/docs/extended_api/synchronization_primitives/barrier.md +++ b/docs/extended_api/synchronization_primitives/barrier.md @@ -21,8 +21,9 @@ It has the same interface and semantics as [`cuda::std::barrier`], with the ## Barrier Operations -| [`cuda::barrier::init`] | Initialize a `cuda::barrier`. `(friend function)` | -| [`cuda::device::barrier_native_handle`] | Get the native handle to a `cuda::barrier`. `(function template)` | +| [`cuda::barrier::init`] | Initialize a `cuda::barrier`. `(friend function)` | +| [`cuda::device::barrier_native_handle`] | Get the native handle to a `cuda::barrier`. `(function template)` | +| [`cuda::barrier::wait_parity/try_wait_parity`] | Wait on or check the parity of the barrier. | ## NVCC `__shared__` Initialization Warnings @@ -98,6 +99,7 @@ __global__ void example_kernel() { [`cuda::barrier::init`]: ./barrier/init.md [`cuda::device::barrier_native_handle`]: ./barrier/barrier_native_handle.md +[`cuda::barrier::wait_parity/try_wait_parity`]: ./barrier/wait_parity.md [`cuda::std::barrier`]: https://en.cppreference.com/w/cpp/thread/barrier diff --git a/docs/extended_api/synchronization_primitives/barrier/wait_parity.md b/docs/extended_api/synchronization_primitives/barrier/wait_parity.md new file mode 100644 index 0000000000..58c6bec9ce --- /dev/null +++ b/docs/extended_api/synchronization_primitives/barrier/wait_parity.md @@ -0,0 +1,45 @@ +--- +grand_parent: Extended API +parent: Barriers +--- + +# `cuda::barrier::wait_parity` and `cuda::barrier::try_wait_parity` + +Defined in header ``: + +```cuda +__host__ __device__ void cuda::std::barrier::wait_parity(bool phase); +__host__ __device__ bool cuda::std::barrier::try_wait_parity(bool phase); +``` + +`barrier::wait_parity` stalls execution while the barrier is not at the specified parity. +`barrier::try_wait_parity` returns true if the parity of the barrier matches the given parity. + +## Return Value + +`barrier::try_wait_parity` returns a boolean representing whether the barrier is at the given phase + + + +[See it on Godbolt](https://godbolt.org/z/dr4798Y76){: .btn } + + +[`cuda::thread_scope`]: ./thread_scopes.md + +[thread.barrier.class paragraph 12]: https://eel.is/c++draft/thread.barrier.class#12 + +[coalesced threads]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#coalesced-group-cg + +[`concurrentManagedAccess` property]: https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html#structcudaDeviceProp_116f9619ccc85e93bc456b8c69c80e78b +[`hostNativeAtomicSupported` property]: https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaDeviceProp.html#structcudaDeviceProp_1ef82fd7d1d0413c7d6f33287e5b6306f +