biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Use kvh.count for GMEM tile coordinate (pipeline-tracked SSA value)

Browse Source
This commit is contained in:
biondizzle
2026-05-22 21:58:43 +00:00
parent 254f7be884
commit 7e832aa527
1 changed files with 11 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files

38
tests/fmha_v3_stage_c_example9.py
View File

@@ -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 =====