[QST]some question about blackwell fmha
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
when ElementAccumulatorQK=float, acc_qk = q@k, will exist in tensor memory with float type, when next gemm acc_qk @ v, type of acc_qk is float, type of v is cutlass::half_t, acc_qk should trans to half_t firstly, but I don't see the transform code in sm100_fmha_fwd_kernel_tma_warpspecialized.hpp, where is it ? And question2, for the code: tOrP0.data() = tOrP0.data().get() + uint32_t(TmemAllocation::P0); why should tOrP0 offset 32 before matmul with v ? shouldn't be tmem_base_addr? I see the tOrP0 layout is tmem_16b o ((_128,_16),_1,(_4,_2)):((_131072,_1),_0,(_16,_64))
,why is the row coordinate stride _131072 , not the original _65536 ?
I am very confused and would appreciate it if you could answer my question. Thank you.
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 by reading sm100_fmha_fwd_kernel_tma_warpspecialized.hpp and tracing the accumulator and tOrP0 layout around the two operations named in the issue. Compare the declared layouts and address offsets with the relevant tensor-memory and matmul documentation; done means providing a clear explanation of the type conversion, the 32-element offset, and the row stride.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- performance
- Issue type
- Documentation
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Quiet
- Clarity
- Needs clarification
- Newbie friendliness
- 38/100