[FEA][CuteDSL] Expose nvvm warp vote instruction

[tridao]

Is your feature request related to a problem? Please describe. I'm planning using these warp vote instructions to coordinate threads in the same warp.

Describe the solution you'd like Can Cute DSL expose the nvvm warp vote instructions with mode {any, all}? https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#vote Currently only the nvvm.vote.sync with BALLOT mode is exposed via the function cute.arch.vote_ballot_sync.

I looked at cutlass._mlir.dialects._nvvm_ops_gen.py but couldn't find a way to access the nvvm.vote.sync.

[keithzzzzz]

Thanks for pointing this out.

Yes, _nvvm_ops_gen.py sometimes doesn't contain some ops we would like to use. Truly no {any, all} modes exposed in nvvm ir in current version.

Before nvvm exposes those ops, we can use llvm.inline_asm to inline ptx of vote.any.sync & vote.all.sync. You can refer to the cute.arch.exp2 (details here). We will also post an example to show how to use inline ptx of vote.any.sync & vote.all.sync later.

[github-actions[bot]]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[github-actions[bot]]

This issue has been labeled inactive-90d due to no recent activity in the past 90 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.