[FEA][CuteDSL] Expose nvvm warp vote instruction
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
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.
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start with cute.arch.vote_ballot_sync and inspect cutlass._mlir.dialects._nvvm_ops_gen.py alongside the NVVM vote specification. Determine how the existing BALLOT exposure is wired, then make the ANY and ALL modes available through Cute DSL and verify that all three vote modes are accessible.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp, python
- Domain
- compilers
- Issue type
- Feature
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 35/100