Skip to content

Commit

Permalink
[skip-tests] Use same format for standard lib text modification
Browse files Browse the repository at this point in the history
  • Loading branch information
ahendriksen committed Sep 12, 2023
1 parent 2040fa6 commit bb989b1
Showing 1 changed file with 11 additions and 28 deletions.
39 changes: 11 additions & 28 deletions libcudacxx/docs/extended_api/synchronization_primitives/barrier.md
Original file line number Diff line number Diff line change
Expand Up @@ -84,34 +84,17 @@ This may complete the current phase.

### Phase Completion of a `cuda::barrier` with tx-count support

Differences to phase completion of `std::barrier` are highlighted in bold in the
text below.

A barrier is a thread coordination mechanism whose lifetime consists of a
sequence of barrier phases, where each phase allows at most an expected number
of threads to block until the expected number of threads **and the expected number
of transaction-based asynchronous operations** arrive at the barrier.

Each barrier phase consists of the following steps:

1. The `expected count` is decremented by each call to `arrive,arrive_and_drop`, or
**`cuda::device::barrier_arrive_tx`**.
2. **The transaction count is incremented by each call to
`cuda::device::barrier_arrive_tx` and decremented by the completion of
transaction-based asynchronous operations such as `cuda::memcpy_async_tx`**.
3. Exactly once after **both** the expected count and **the transaction count** reach
zero, a thread executes the completion step during its call to `arrive`,
`arrive_and_drop`, or `wait`, except that it is implementation-defined whether
the step executes if no thread calls `wait`.
4. When the completion step finishes, the expected count is reset to what was
specified by the expected argument to the constructor, possibly adjusted by
calls to `arrive_and_drop`, **the transaction count is reset to zero**, and the
next phase starts.

Concurrent invocations of the member functions of barrier **and the non-member
barrier APIs in cuda::device**, other than its destructor, do not introduce data
races. The member functions `arrive` and `arrive_and_drop`, and the non-member
function **cuda::device::barrier_arrive_tx**, execute atomically.
Modify [[thread.barrier.class]](http://eel.is/c++draft/thread.barrier.class) as follows:

> A barrier is a thread coordination mechanism whose lifetime consists of a sequence of barrier phases, where each phase allows at most an expected number of threads to block until the expected number of threads **and the expected number of transaction-based asynchronous operations** arrive at the barrier.
> Each _barrier phase_ consists of the following steps:
>
> 1. The _expected count_ is decremented by each call to `arrive`,`arrive_and_drop`**, or `cuda::device::barrier_arrive_tx`**.
> 2. **The _transaction count_ is incremented by each call to `cuda::device::barrier_arrive_tx` and decremented by the completion of transaction-based asynchronous operations such as `cuda::memcpy_async_tx`.**
> 3. Exactly once after **both** the _expected count_ **and the _transaction count_** reach zero, a thread executes the _completion step_ during its call to `arrive`, `arrive_and_drop`, or `wait`, except that it is implementation-defined whether the step executes if no thread calls `wait`.
> 4. When the completion step finishes, the _expected count_ is reset to what was specified by the `expected` argument to the constructor, possibly adjusted by calls to `arrive_and_drop`, **the _transaction count_ is reset to zero,** and the next phase starts.
>
> Concurrent invocations of the member functions of barrier **and the non-member barrier APIs in `cuda::device`**, other than its destructor, do not introduce data races. The member functions `arrive` and `arrive_and_drop`, **and the non-member function `cuda::device::barrier_arrive_tx`**, execute atomically.
## Implementation-Defined Behavior

Expand Down

0 comments on commit bb989b1

Please sign in to comment.