diff --git a/docs/api/synchronization_library/barrier.md b/docs/api/synchronization_library/barrier.md index 3223162391..b3f479278e 100644 --- a/docs/api/synchronization_library/barrier.md +++ b/docs/api/synchronization_library/barrier.md @@ -23,7 +23,7 @@ cuda::std::barrier<> bb; cuda::barrier c; ``` -The class template `barrier` may also be declared without initialization; a +The class template `barrier` may also be declared without initialization in the `cuda::` namespace; a friend function `init` may be used to initialize the object. ```c++ @@ -31,18 +31,27 @@ friend function `init` may be used to initialize the object. __shared__ cuda::barrier b; init(&b, 1); // Use this friend function to initialize the object. +/* +namespace cuda { + template + __host__ __device__ void init(barrier* bar, std::ptrdiff_t expected); + template + __host__ __device__ void init(barrier* bar, std::ptrdiff_t expected, CompletionF completion); +} +*/ ``` +- Expects: `*bar` is trivially initialized. +- Effects: equivalent to initializing `*bar` with a constructor. + In the `device::` namespace, a `__device__` free function is available that provides direct access to the underlying PTX state of a `barrier` object, if its scope is `thread_scope_block` and it is allocated in shared memory. ```c++ namespace cuda { namespace device { - -__device__ std::uint64_t* barrier_native_handle( + __device__ std::uint64_t* barrier_native_handle( barrier& b); - }} ```