NVIDIA / NVIDIA/cutlass

[QST]some question about blackwell fmha

Open
#2,909 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage inactive-30d inactive-90d question
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

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.