fix: use make_tiled_tma_atom instead of _A for P TMA

This commit is contained in:
2026-05-24 03:43:46 +00:00
parent 244aca13ed
commit 3b6f19043f

View File

@@ -115,16 +115,12 @@ class FmhaKernel:
tma_c,mC = cpasync.make_tiled_tma_atom(cpasync.CopyBulkTensorTileS2GOp(),c,epi_s,self.epi_tile)
# SMEM-P: TMA for P (GMEM→SMEM). We reuse Q's GMEM buffer as gP.
# By the time softmax computes P, Q has already been loaded to SMEM,
# so Q's GMEM is free for reuse as a P staging buffer.
if const_expr(self.use_smem_p):
p_s = cute.slice_(self.p_smem_s,(None,None,None,0))
tma_p,gP_tma = cute.nvgpu.make_tiled_tma_atom_A(
utils.sm100.cluster_shape_to_tma_atom_A(self.cluster_shape_mn, pv_mma.thr_id),
mQ, p_s, self.qk_mma_tiler, pv_mma, self.cluster_layout_vmnk.shape
)
tma_p,gP_tma = cpasync.make_tiled_tma_atom(cpasync.CopyBulkTensorTileG2SOp(), mQ, p_s, self.qk_mma_tiler)
else:
tma_p = tma_q # dummy, dead code
tma_p = tma_q # dummy
gP_tma = mQ, dead code
# Always create a valid mLSE tensor for the kernel.
# CuTeDSL doesn't support None parameters in @cute.kernel.
# For normalize=True, mLSE is unused (dead-code-eliminated by compiler).