diff --git a/dsv4/kernels/attention/fmha.py b/dsv4/kernels/attention/fmha.py index 57ef90b0..b5ca201b 100644 --- a/dsv4/kernels/attention/fmha.py +++ b/dsv4/kernels/attention/fmha.py @@ -490,15 +490,14 @@ class FmhaKernel: # Step 6: TMA store SMEM → GMEM # The normalized O is now in sC (written by the correction epilog). - # Use the same TMA store pattern as the CUTLASS FMHA reference. - # Partition sC and gC for the bulk TMA copy. - # Get 2D views of sC and gC for TMA. - sC_epi = sC # sC already has c_smem_s layout, TMA can handle it - gC_epi = tCgC # Use the pre-partitioned GMEM tensor + # The tma_c was created with (c, epi_s, epi_tile) where epi_s = select(c_smem_s, mode=[0,1]). + # We need to partition sC and the GMEM output for the TMA copy. + sC_epi_layout = cute.select(self.c_smem_s, mode=[0, 1]) # 2D SMEM layout + sC_epi = cute.make_tensor(sC.iterator, sC_epi_layout) # 2D view of sC + gC_epi = cute.local_tile(mC, epi_tile, (Int32(0), Int32(0))) # 2D GMEM tile bSG_sC, bSG_gC = cpasync.tma_partition( tma_c, 0, cute.make_layout(1), - cute.group_modes(sC, 0, 2), - cute.group_modes(gC_epi, 0, 2), + sC_epi, gC_epi, ) # One TMA store for the full output tile if warp_idx == self.epilogue_warp_id[0]: