[EPIC] cuda::ptx platform for exposing PTX instructions

[ahendriksen]

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.

The development of higher-level APIs lags behind the introduction of new PTX instructions in the CUDA Toolkit.

Describe the solution you'd like

Provide a means to add PTX instructions before higher-level exposure has been developed.

That is:

• Add the header cuda/ptx
• Add the namespace cuda::ptx
• Add PTX instructions to expose cluster-level and other features.

Tasks

Preview Give feedback

1. Proof of concept of cuda::ptx #616
   +0
2. [FEA]: First iteration of SM90 PTX instruction exposure #609
   4 of 4
   feature request +1
   

Describe alternatives you've considered

The alternative is to wait for higher-level exposure, which may take a while.

Another alternative is to provide an intermediate exposure like the recent addition of cuda::device::experimental::cp_async_bulk_global_to_shared and friends. These wrapper functions are a hybrid between pure PTX wrapping functions using native types, and higher-level APIs that take higher-level types (like cuda::barrier<..>). These are very useful for documenting features to outside developers, but require design and deviate from the PTX ISA.

Additional context

No response

[ahendriksen]

Closing this epic in favor of a Hopper-specific epic. The goal of setting up a platform for exposing PTX instructions has been achieved..