[FEA]: Refactor cuda::(std::)barrier implementation to isolate inline PTX
- 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
Assessment
This issue has not been assessed yet.