NVIDIA / NVIDIA/cccl

[FEA]: Refactor cuda::(std::)barrier implementation to isolate inline PTX

Open
#70 0 comments 0 reactions 0 assignees View on GitHub
libcu++
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

The current `cuda::barrier` implementation uses inline PTX directly in the implementation of the member functions [e.g.,](https://github.com/NVIDIA/libcudacxx/blob/9d14ca711b2fda04df9ce44728c9eed2578acdf1/include/cuda/std/detail/libcxx/include/__cuda/barrier.h#L386-L394).

In this approach we are not taking advantage of the full breadth of options available with `mbarrier`. Furthermore, this approach does not scale well to the addition of future architecture improvements.

Ideally, we want to extract the barrier functionality into free functions that hide the architecture specialization. Furthermore, these free functions could be generated by a code generator to better take advantage of the breadth of options.

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.