NVIDIA / NVIDIA/cccl

[FEA]: Add `__pipeline_arrive_on_noinc`

Open
#2,927 1 comment 0 reactions 2 assignees Claimed by @bernhardmgruber View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

### Is this a duplicate?

- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Area

libcu++

### Is your feature request related to a problem? Please describe.

Please expose a `noinc` version of `__pipeline_arrive_on`:
It allows to avoid double barrier arrive when synchronizing LDGSTS via a barrier.

### Describe the solution you'd like

```
__device__ __forceinline__
void __pipeline_arrive_on_noinc (uint64_t* barrier)
{
asm volatile (
"cp.async.mbarrier.arrive.noinc.shared.b64 [%0];"
:: "r"(static_cast(__cvta_generic_to_shared(barrier)))
);
}
```

### Describe alternatives you've considered

_No response_

### Additional context

_No response_

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.