diff --git a/tests/fmha_v3_stage_c_example9.py b/tests/fmha_v3_stage_c_example9.py
index 23008686..19ac1128 100644
--- a/tests/fmha_v3_stage_c_example9.py
+++ b/tests/fmha_v3_stage_c_example9.py
@@ -179,13 +179,6 @@ class FmhaV3StageCMulti:
         b_lay = cute.make_layout(cute.slice_(cl_vmnk,(0,None,0,0)).shape)
         tBsK,tBgK = cpasync.tma_partition(tma_k,0,b_lay,cute.group_modes(sK,0,3),cute.group_modes(tCgK,0,3))
         tVsV,tVgV = cpasync.tma_partition(tma_v,0,b_lay,cute.group_modes(sV,0,3),cute.group_modes(tCgV,0,3))
-        # DIAG: print shapes before pre-slice
-        cute.printf("tAgQ shape: {}\n", cute.shape(tAgQ))
-        cute.printf("tBgK shape: {}\n", cute.shape(tBgK))
-        cute.printf("tVgV shape: {}\n", cute.shape(tVgV))
-        cute.printf("tAsQ shape: {}\n", cute.shape(tAsQ))
-        cute.printf("tBsK shape: {}\n", cute.shape(tBsK))
-        cute.printf("tVsV shape: {}\n", cute.shape(tVsV))
         tAgQ = tAgQ[(None,0,None,0)]; tBgK = tBgK[(None,None,0,0)]; tVgV = tVgV[(None,0,None,0)]
 
         tCrQ = qk_mma.make_fragment_A(sQ); tCrK = qk_mma.make_fragment_B(sK)
@@ -210,33 +203,24 @@ class FmhaV3StageCMulti:
         pipeline.pipeline_init_wait(cluster_shape_mn=cl_vmnk)
 
         # ===== TMA LOAD warp =====
-        # Multi-tile GMEM indexing — combination that finally makes it work:
-        #
-        # 1. SSA-seeded kv_coord (n_kv_tiles - n_kv_tiles, not Int32(0)). The
-        #    JIT folds literal Int32(0) at trace time; subtracting an SSA
-        #    value from itself produces an SSA register zero that's tracked.
-        #
-        # 2. Drop the try_acquire/pk pattern. We had `pk = kvp.try_acquire()`
-        #    outside the loop and `pk = cutlass.Boolean(1)` at the end of
-        #    each iter — a second loop-carried variable. The JIT's automatic
-        #    iter_args yielding *can* handle one or both, but mixing two
-        #    mutated variables exposed an edge case where `pk` won the
-        #    iter_args slot and `kv_coord` got constant-folded.
-        #
-        #    The reference uses bare `acquire_and_advance()` with no `pk`,
-        #    so just match that. One loop-carried variable (kv_coord),
-        #    pipeline state managed internally.
+        # TMA partition shapes (confirmed by diagnostic):
+        #   tBgK: (((64,128),1),n_kv_tiles,1,1) — mode 1 = GMEM tile dim
+        #   tVgV: (((64,128),1),1,n_kv_tiles,1) — mode 2 = GMEM tile dim
+        # Pre-slice (None,None,0,0) / (None,0,None,0) preserves the tile dim.
+        # After slice, both are 2-mode: (tile_data, n_kv_tiles).
+        # The challenge: getting the GMEM coordinate to propagate at runtime.
+        # kv_coord via SSA-seeding didn't work (constant-folded by JIT).
+        # Using kvh.count (pipeline internal counter) instead — this IS
+        # a proper SSA runtime value tracked through the pipeline state.
         if warp_idx == self.tma_warp_id:
             qp.reset(); qh = qp.acquire_and_advance()
             cute.copy(tma_q, tAgQ[(None, Int32(0))], tAsQ[(None, qh.index)], tma_bar_ptr=qh.barrier)
             qp.tail()
             kvp.reset()
-            kv_coord = n_kv_tiles - n_kv_tiles  # SSA runtime zero
             for kt in cutlass.range(0, n_kv_tiles, 1, unroll=1):
                 kvh = kvp.acquire_and_advance()
-                cute.copy(tma_k, tBgK[(None, kv_coord)], tBsK[(None, kvh.index)], tma_bar_ptr=kvh.barrier)
-                cute.copy(tma_v, tVgV[(None, kv_coord)], tVsV[(None, kvh.index)], tma_bar_ptr=kvh.barrier)
-                kv_coord = kv_coord + 1
+                cute.copy(tma_k, tBgK[(None, kvh.count)], tBsK[(None, kvh.index)], tma_bar_ptr=kvh.barrier)
+                cute.copy(tma_v, tVgV[(None, kvh.count)], tVsV[(None, kvh.index)], tma_bar_ptr=kvh.barrier)
             kvp.tail()
 
         # ===== MMA warp =====