[QST] Difference between `cute.arch.block_idx_in_cluster` vs `cute.arch.block_in_cluster_idx`
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
What is your question?
The example hopper dense gemm kernel retrieved first:
bidx_in_cluster = cute.arch.block_in_cluster_idx()
and then
cta_rank_in_cluster = cute.arch.make_warp_uniform(
cute.arch.block_idx_in_cluster()
)
cluster_coord_mnk = cta_layout_mnk.get_flat_coord(cta_rank_in_cluster)
My understanding was that cute.arch.block_idx_in_cluster() was nothing but a colexicographical traversal of the 3D cute.arch.block_in_cluster_idx() coordinate, meaning cluster_coord_mnk and bidx_in_cluster were the same thing.
I wonder if my understanding was correct? If so, is there any particular reason why the example kernel tries to retrieve both?
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 examples/python/CuTeDSL/hopper/dense_gemm.py at the linked lines and trace the uses of block_in_cluster_idx() and block_idx_in_cluster(). Compare the two returned values in the surrounding kernel context; done means documenting their distinction and whether cluster_coord_mnk and bidx_in_cluster are equivalent.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python
- Domain
- performance
- Issue type
- Documentation
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Mostly clear
- Newbie friendliness
- 42/100