diff --git a/dsv4/kernels/attention/fmha.py b/dsv4/kernels/attention/fmha.py
index 68cf4718..f319a43f 100644
--- a/dsv4/kernels/attention/fmha.py
+++ b/dsv4/kernels/attention/fmha.py
@@ -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).