[QST][CUTEDSL] Address Misalignment in FP8 Gemm
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
I am seeking help from Cutlass community. I encountered misaligned address (16 bytes requested) when load matrix from smem to register. I found out the problem started at "thr_copy_ldmatrix_A.partition_S(sA)".
Does anyone know how to debug and solve this issue?
--------------------------------
DEBUG: sA:
tensor<ptr<i8, smem, align<1024>> o ((64,1),(8,8),(1,4)):((1,0),(64,512),(0,4096))>
DEBUG: thr_copy_ldmatrix_A:
Tiled Copy
Tiler MN: (32:1,32:1)
TV Layout tiled: ((4,8,2,2),((4,2,2),(1,1))):((128,1,16,0),((32,8,512),(0,0)))
Copy Atom
ThrID: 32:1
TV Layout Src: ((2,2,4,2),16):((16,128,32,0),1)
TV Layout Dst: ((4,8),(1,2,2,2)):((32,1),(1,16,8,128))
Value type: i8
--------------------------------
tCsA_copy_view = thr_copy_ldmatrix_A.partition_S(sA)
DEBUG: tCsA_copy_view:
tensor<ptr<i8, smem, align<8>> o (((8,2),2),2,2,(1,4)):(((1,128),1024),32,2048,(0,4096))>
--------------------------------
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 at thr_copy_ldmatrix_A.partition_S(sA) and inspect the reported shared-memory tensor and copy-view layouts, especially their alignment values. Reproduce the FP8 GEMM load and determine what change prevents the 16-byte misaligned-address error when loading from shared memory to registers.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc, performance
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Needs clarification
- Newbie friendliness
- 35/100