Add documentation for the barrier::wait_parity extension

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
 

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/std/barrier>`:
+
+```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
+
+<!-- TODO: Create an example when trunk is live on godbolt
+## Example
+
+```cuda
+#include <cuda/barrier>
+
+__global__ void example_kernel(cuda::barrier<cuda::thread_scope_block>& bar) {
+  bar.wait_parity(false);
+}
+```
+-->
+
+[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
+