From 08b05b159f996c987f94e478ec127fa95eb90213 Mon Sep 17 00:00:00 2001 From: Wesley Maxey Date: Mon, 26 Jul 2021 14:50:59 -0700 Subject: [PATCH 1/2] Add documentation for the barrier::wait_parity extension --- .../synchronization_primitives/barrier.md | 6 ++- .../barrier/wait_parity.md | 45 +++++++++++++++++++ 2 files changed, 49 insertions(+), 2 deletions(-) create mode 100644 docs/extended_api/synchronization_primitives/barrier/wait_parity.md diff --git a/docs/extended_api/synchronization_primitives/barrier.md b/docs/extended_api/synchronization_primitives/barrier.md index 17ad379fc6..c46270363b 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 a `specific` phase 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..475f15e38a --- /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 phase. +`barrier::try_wait_parity` queries the the state of the barrier against the specified phase. + +## 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 + From 219c1d9c5db743c8c255f0fe5e0b8cdfc0294bbd Mon Sep 17 00:00:00 2001 From: Wesley Maxey Date: Thu, 29 Jul 2021 11:20:29 -0700 Subject: [PATCH 2/2] Rephrase barrier docs to use 'parity' more --- docs/extended_api/synchronization_primitives/barrier.md | 2 +- .../synchronization_primitives/barrier/wait_parity.md | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/docs/extended_api/synchronization_primitives/barrier.md b/docs/extended_api/synchronization_primitives/barrier.md index c46270363b..e10d95effb 100644 --- a/docs/extended_api/synchronization_primitives/barrier.md +++ b/docs/extended_api/synchronization_primitives/barrier.md @@ -23,7 +23,7 @@ It has the same interface and semantics as [`cuda::std::barrier`], with the | [`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 a `specific` phase of the barrier | +| [`cuda::barrier::wait_parity/try_wait_parity`] | Wait on or check the parity of the barrier. | ## NVCC `__shared__` Initialization Warnings diff --git a/docs/extended_api/synchronization_primitives/barrier/wait_parity.md b/docs/extended_api/synchronization_primitives/barrier/wait_parity.md index 475f15e38a..58c6bec9ce 100644 --- a/docs/extended_api/synchronization_primitives/barrier/wait_parity.md +++ b/docs/extended_api/synchronization_primitives/barrier/wait_parity.md @@ -12,8 +12,8 @@ __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 phase. -`barrier::try_wait_parity` queries the the state of the barrier against the specified 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