[FEA]: First iteration of SM90 PTX instruction exposure

[jrhemstad]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

libcu++

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

As a follow on to the proof of concept for exposing C++ APIs for PTX instructions.

I would like C++ wrapper APIs for the following PTX instructions:

Tasks

Preview Give feedback

1. st.async
2. red.async
3. mbarrier.test_wait/mbarrier.try_wait .cluster scope
4. PTX wrappers for mbarrier.arrive variants #659
   2 of 2
   feature request +1
   

Describe the solution you'd like

C++ APIs consistent with what was decided from the proof of concept here: #573

Describe alternatives you've considered

No response

Additional context

No response

[ahendriksen]

The mbarrier.arrive variants are added in #574. So once that is merged, they can be checked off in the task list. It appears that I do not have access to the task list. @jrhemstad : can you pull the tasks out into separate issues that I can close? Or make me co-owner of the task list somehow?

[jrhemstad]

Odd, you should already have permissions to be able to edit the task list. In the meantime, I just went ahead and made a sub-issue (#659) and updated your PR to close that issue when merged.

[ahendriksen]

I submitted the following PRs:

• Add cuda::ptx::* namespace #574 (mbarrier.arrive)
• Add cuda::ptx::red_async #1080
• Add cuda::ptx::st_async #1078
• Add cuda::ptx:mbarrier_{try/test}_wait{_parity} #674

[jrhemstad]

Closing as completed! Thanks @ahendriksen !