NVIDIA / NVIDIA/cccl

<cuda::barrier> should add try_wait

Open
#970 0 comments 0 reactions 0 assignees View on GitHub
libcu++
Dominant language
C++
Stars
2.5k
Forks
487
Avg merge
2d 7h
Merged PRs (30d)
296

Description

It can be beneficial to test the state of a barrier prior to needing the barrier to be resolved. This allows one to hide the latency of testing the barrier behind computation. This latency can be quite significant if the barrier is in device or system memory.

For example, instead of:
`arrive(x)`
`compute1`
`compute2`
`wait(x) (usually issues a load which we wait on shortly after, exposing latency)`

One could do:
`arrive(x)`
`compute1`
`x = try_wait(x)`
`compute2`
`wait(x)`

Where the final wait can test the value of x that was previously loaded overlapped with compute-2.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.