PTX: Add cuda::ptx:cp_async_bulk_*

[ahendriksen]

Add:

• cp.async.bulk
• cp.async.bulk.tensor
• cp.reduce.async.bulk.tensor
• cp.async.bulk.wait_group
• cp.async.bulk.commit_group

Description

closes #1398, #1399, #1400, #1401, #1402

Checklist

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

[ahendriksen]

Not sure if this is going to be caught by CI: the .multicast variants of cp.async.bulk{.tensor} are officially part of SM90, but are officiously part of (i.e. fast on) SM90a. Therefore, ptxas emits the following warning when any multicast instruction is used when compiling for sm90:

Advisory: '.multicast::cluster' modifier on instruction 'cp.async.bulk{.tensor}' should be used 
on .target 'sm_90a' instead of .target 'sm_90' as this feature is expected to have substantially 
reduced performance on some future architectures

However, since we compile with --warning-as-error, this advisory gets treated as an error:

#$ ptxas --warning-as-error -arch=sm_90 -m64  "/tmp/tmpxft_00009e7f_00000000-6_ptx.cp.async.bulk.compile.pass.ptx"  -o "/tmp/tmpxft_00009e7f_00000000-8_ptx.cp.async.bulk.compile.pass.cubin" 
ptxas /tmp/tmpxft_00009e7f_00000000-6_ptx.cp.async.bulk.compile.pass.ptx, line 154; error   : Advisory: '.multicast::cluster' modifier on instruction 'cp.async.bulk{.tensor}' should be used on .target 'sm_90a' instead of .target 'sm_90' as this feature is expected to have substantially reduced performance on some future architectures
ptxas fatal   : Ptx assembly aborted due to errors
# --error 0xff --

Can we disable this behavior somehow in lit? @miscco

[ahendriksen]

Thanks for the review! I have implemented most of the comments. One comment is addressed in #1359, and one comment (combining the ptx isa guards in the tests) I hope we can punt to a future PR.