Add cuda::ptx::* namespace

[ahendriksen]

Description

Add cuda/ptx header and cuda::ptx namespace .

closes #659

• Add the header cuda/ptx
• Add the namespace cuda::ptx
• Add a single PTX wrapper for mbarrier.arrive.expect_tx to demonstrate implementation, testing, documentation, and internal use by higher-level libcu++ APIs.

Summary of intent:

• Provide stable PTX wrapping API that can be used by libcu++ internal code, and included using #include <cuda/ptx> by outside users.
• Expose PTX instructions according to the following priorities: (1) instructions that expose new hardware features, (2) instructions that are not (yet) exposed through other APIs, or (3) instructions that are useful for power users.
• The PTX instructions are tested to assemble correctly (that is: they do not lead to compilation errors)
• The documentation of the PTX instruction is shallow. The PTX ISA documentation is linked for each instruction, and the CTK/CCCL version that the instruction was introduced is noted. Usage or other documentation is optional.
• The exposure uses C++ language features like type-level parameters to ensure the scalable exposure of PTX instructions variants, introduced by .sem, .space, .scope, .op and others.

Checklist

• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[miscco]

I really like this idea. Thanks for bringing it forward.

What I am not 100% sure is if we really want to put this in libcu++ or whether we want to make this its completely own subproject, that just inherits some of the common machinery.

The advantage would be that we would be "separated" from CTK releases which would make it a bit more explicit that we are not providing any guarantees here

[ahendriksen]

Thanks for your comments. I have left responses inline.

[ahendriksen]

Thanks for the comments. Many have been resolved. I replied inline to the open discussions.

[leofang]

This is amazing work. Whatever the team decides, I am happy to see the fruition.

For analogy, this is like CUDA Python providing low-level Python bindings for CUDA C APIs, so that Python users can access CUDA from Python. We are witnessing the advent of low-level C++ binding for PTX APIs (so that C++ users can access CUDA from C++, without writing inline PTX)!

[fduguet-nv]

This is a great initiative.
I think it would be very helpful to get such intrinsic style functions to access ptx instructions, especially on the memory model (e.g. ld.acquire.xxx). I would definitely use this API as I find using inline ptx is error-prone and cumbersome.
Also, I tend to believe that this thin layer allows updates in the code in the future at this header-only library level, without requiring us developers to update code at the pace GPU architecture would evolve.

[ahendriksen]

I think this PR is ready to merge. All comments have been incorporated. In addition, I have made some small improvements after using the API in barrier.h:

• Rename space_shared_cluster to space_cluster. The original name was very close to the PTX naming scheme of shared::cluster, but was unnecessarily long.
• Remove unlikely to be used state spaces, like const, sreg, reg etc.
• Stay even closer to the original PTX exposure by exposing the initial short instructions and also the newer longer variants. The mapping from PTX ISA (CTK 11.0, 11.8, 12.0) to cuda::ptx exposure is documented in the comments of cuda::ptx::mbarrier_arrive.

[jrhemstad]

/backport

[github-actions]

Backport failed for branch/2.3.x, because it was unable to cherry-pick the commit(s).

Please cherry-pick the changes locally.

git fetch origin branch/2.3.x
git worktree add -d .worktree/backport-574-to-branch/2.3.x origin/branch/2.3.x
cd .worktree/backport-574-to-branch/2.3.x
git checkout -b backport-574-to-branch/2.3.x
ancref=$(git merge-base 83b3365bbc12a9db248b67c22052413d41fae97e 9e9fb70b712a6799791997d3c70db6a28a71af72)
git cherry-pick -x $ancref..9e9fb70b712a6799791997d3c70db6a28a71af72