Add parity waiting

[ogiroux]

Overview

Add APIs to cuda::barrier that allow for waiting on the parity of the barrier. This is an extension to std::barrier that relies on the internal binary phase of the barrier.

Motivation

• PTX for SM_80 provides parity waiting for barriers rather than using an arrival token that is held in every thread.
  • This is a significant optimization if users can instead derive parity from loop induction variables or other state.
• https://nvbugs/3248969

Requirements

• Add the member function barrier::try_wait_parity
• Add the member function barrier::wait_parity
• Add test coverage for barrier::wait_parity

Design

• Add both front-end wait_parity APIs similarly to wait and try_wait.
  • wait_parity is implemented by polling try_wait_parity with back-off.
  • try_wait_parity queries the parity of the barrier with a parity provided by the user.

Test Plan

• All currently existing tests should continue to function as expected.
• Add a test for wait_parity which also provides coverage for try_wait_parity.

Performance Tests

• There is currently no performance criteria for this feature.

Documentation

• Documentation will be submitted post-release in a subsequent PR.
  • Documentation changes for wait_parity/try_wait_parity #190 contains documentation for this feature.

Complex Internal Systems

• The internal representation of std::barrier is able to queried to obtain the parity of the barrier.

Stakeholders

• Olivier Giroux - Implementer
• Wesley Maxey - Maintainer

[ogiroux]

I'm going to add a bool try_wait(token) facility to this PR before it is merged in 2.0.0. Basically this will just go from having it as an internal name, to having it as an external name (remove the __).

Discussion from internal developer day pointed out that you can get a good benefit from hiding the latency of the load and/or control the polling backoff yourself. It's already in here, might as well make it usable.

[wmaxey]

@ogiroux Does this require any additional tests you would like to see? I was going to permute a few things, but this would be a good opportunity for any edge cases you can think of.

[wmaxey]

Latest clean virtual: https://builds4u.nvidia.com/dvs/#/change/3023208755719989.5?eventType=Virtual

[griwes]

LGTM except for the phrasing in the documentation.

[griwes]

I don't like this phrasing. @ogiroux how would you feel about "Wait on a phase whose parity matches the argument." or some such?

[griwes]

See the previous comment. It's really weird that the description here doesn't mention parity at all.

[wmaxey]

In the interest of time, I'm going to split the documentation changes into another PR and merge the changes.