Commit Graph

100 Commits

Author SHA1 Message Date
0ced79ab37 Clean up: archive diagnostics and superseded tests
Kept:
- example10 (CUTLASS LLM, O rescale + final normalize)
- example9 (SSA kv_coord version)
- working_softmax_maybe.py (working softmax snapshot from before the nuke)
- test_fmha_v3_stage_c.py (identity softmax baseline, n=128 cos 0.999998)
- test_fmha_v3.py (Stage A+B baseline)
- layertest.py, cudagraph_test.py (required)
- test_cutedsl.py, test_fp4_roundtrip.py (NVFP4 tests)

Archived: diag_tma_*, example8, test_diag_multitile, test_reference_fmha,
test_ref_minimal, test_tma_coord, test_fmha_v3_diag*, test_fmha_v3_12w,
test_dense_router, test_interleave*, test_fused_step1, test_router,
test_cache, test_compile_custom_op, test_custom_op, test_layer_schedule
2026-05-23 00:17:07 +00:00
0dc1fe5f71 FIX: (None,0,None,0) for ALL tma_partition outputs — verified shapes on B200
DIAG OUTPUT (n=256, inside @cute.kernel):
  tAgQ: (((64,128),1), Int32(?), Int32(?), Int32(?))  — 4 modes
  tBgK: (((64,128),1), Int32(?), Int32(?), Int32(?))  — 4 modes
  tVgV: (((64,128),1), 1, 1, 1)                       — 4 modes

After (None,0,None,0) → keeps modes 0 and 2 free → 2D:
  tAgQ: (((64,128),1), Int32(?))
  tBgK: (((64,128),1), Int32(?))
  tVgV: (((64,128),1), 1)

Then [None, kt] indexes the surviving mode 1 (originally mode 2 = KV tiles).
tAgQ[(None, Int32(0))] for Q (1 tile, coordinate is always 0).
Removed diag prints from test_fmha_v3.py.
2026-05-22 23:35:55 +00:00
0534f19aaf auto: pre-test commit 2026-05-22 23:34:03 +00:00
00e103ab91 FIX: (None,0,None,0) pre-slice keeps KV tile axis (mode 2) free
tBgK has 4 modes: (V_grouped, ?, KV_tiles, ?). Mode 2 is the GMEM tile dim.
Old (None,None,0,0) kept modes 0,1 free → mode 2 collapsed to 0 → always tile 0.
8-None no-op slice FAILS — tensor is 4-mode, not 8-mode, at JIT level.

Fix: (None,0,None,0) keeps modes 0,2 free → 2D tensor.
Then tBgK[None, kt] indexes the surviving KV_tiles dim.

Matches CUTLASS reference FMHA pattern:
  tKgK = tKgK_kdl[None, None, 0, batch]
  cute.copy(tma_k, tKgK[None, kv_coord], ...)
2026-05-22 23:25:40 +00:00
30eaba39aa FIX: 8-None no-op pre-slice opens full TMA coordinate space (8 dims)
The tma_partition output has 8 TMA coordinate dimensions, not 4.
The Python-visible shape shows 4 modes, but the TMA descriptor uses
8 coordinates. Without the 8-None no-op pre-slice, modes 4-7 are
collapsed and the GMEM tile axis (mode 4) is pinned to 0.

Pattern that works (confirmed on B200 at n=256 in diag test):
  tBgK = tBgK[(None,None,None,None,None,None,None,None)]  # open 8D
  cute.copy(tma_k, tBgK[None,None,None,None,kt,None,None,None], ...)

The old 4-mode indexing tBgK[(None,None,kt,0)] fails with
'rank mismatch: got 2 and 1' because slicing a 4-mode tensor
produces wrong rank for the TMA coordinate space.

Matches working diag test test_fmha_v3_diag.py exactly.
2026-05-22 23:18:40 +00:00
9c5adcee46 FIX: tma_partition tensors have 4 modes, not 8. Mode 2 is GMEM tile dim.
The 8-mode indexing (tBgK[None,None,None,None,kt,None,None,None]) fails at
JIT compilation with 'coord and shape are weakly congruent' error. The actual
MLIR tensor shape is (((64,128),1),?,?,?) — 4 modes, not 8.

The working fix from commit 845ad98 on the B200 used 4-mode indexing all along:
  tBgK[(None, None, kt, 0)] — mode 2 = GMEM tile dim
  tVgV[(None, 0, kt, 0)] — mode 2 = GMEM tile dim

Updated all files: example10, test_fmha_v3_stage_c, README, docstrings.
2026-05-22 23:08:27 +00:00
d7cdf63c58 Fix test_fmha_v3_stage_c.py: 8-mode TMA indexing (mode 4 = GMEM tile dim) 2026-05-22 22:58:10 +00:00
79e35f10ad Add diag test with 8-mode TMA indexing from commit 2711611 2026-05-22 22:40:09 +00:00
dbd77f2bc4 DOCUMENT: TMA 8-mode indexing — the bug that cost us a full day. README + inline comments. 2026-05-22 21:28:58 +00:00
3a4524c318 Fix identity diag: same 8D TMA indexing fix 2026-05-22 21:21:52 +00:00
2fd6f02e8e FIX: Use full 8D indexing for tBgK/tVgV — mode 4 is the GMEM tile dim 2026-05-22 21:21:23 +00:00
51173b4805 REVERT to working baseline (n=128 cos 0.999998). Multi-tile TMA is a CuTeDSL JIT limitation. 2026-05-22 20:37:21 +00:00
8d5733bb4c Test: use kvh.index (pipeline state) as TMA GMEM coordinate 2026-05-22 20:36:21 +00:00
f2b1cc39d3 SMEM counter: separate allocate_tensor instead of struct field 2026-05-22 20:35:42 +00:00
4955fb5204 Fix SMEM counter type: cutlass.Int32 for MemRange 2026-05-22 20:35:17 +00:00
5f6e33c054 SMEM-backed kv_coord counter — JIT can't constant-fold SMEM reads 2026-05-22 20:34:52 +00:00
18a5c60e26 DEBUG: hardcoded Int32(1) to test if TMA can read tile 1 2026-05-22 20:34:21 +00:00
b7a1deed52 DEBUG: use Int32(kt) directly to test if coordinate matters 2026-05-22 20:34:03 +00:00
c01291f16a Test: kv_coord = warp_idx() * 0 — force SSA from runtime value 2026-05-22 20:33:40 +00:00
da4eebc9a4 DEBUG: add cute.printf for kv_coord runtime value 2026-05-22 20:33:03 +00:00
b8d4ac0ecb Test: Python range() instead of cutlass.range() for TMA loop 2026-05-22 20:32:44 +00:00
ae6197fc78 Test example9: drop try_acquire/pk, single loop-carried kv_coord 2026-05-22 20:32:25 +00:00
20290d7439 REVERT to working example7 (n=128 cos 0.999998). Example8 TMA fix didn't work. 2026-05-22 20:28:15 +00:00
72ad0eca9e Update stage_c test to example8: SSA kv_coord + per-tile O rescale 2026-05-22 20:27:58 +00:00
b1ee693998 Clean up tests: archive superseded files, keep only essential unit tests
Kept in tests/unit/:
- test_fmha_v3.py (stages A+B)
- test_fmha_v3_diag.py (identity softmax, n=128+256)
- test_fmha_v3_stage_c.py (real softmax, n=128 cos 0.999998)
- layertest.py + cudagraph_test.py (required for every change)
- infrastructure: cache, custom_op, cutedsl, router, fp4, fused, interleave

Archived: 19 superseded unit tests + 10 root-level scratch files
Root level: only fmha_v3_stage_c_example7.py remains (now in unit/)
2026-05-22 20:25:27 +00:00
1e0805ad60 Diag: test n=384 (3 tiles) to find crash boundary 2026-05-22 18:07:07 +00:00
1aa4a91d01 Diag: test all sizes 128-1024 2026-05-22 18:06:28 +00:00
8586603280 DEBUG: disable O rescale to isolate NaN cause 2026-05-22 18:05:46 +00:00
45a0fb9971 Add NaN/inf checking to stage C test 2026-05-22 18:01:11 +00:00
147d85a617 CRITICAL FIX: K GMEM slice (None,None,0,0) not (None,0,None,0)
K from QK MMA B-partition has GMEM iter at mode 1, NOT mode 2.
(None,0,None,0) hardcodes mode 1 to 0 → TMA always loads tile 0.
(None,None,0,0) keeps mode 1 free → correct multi-tile loading.

Proof: diag n=256 went from cos 0.711 → 0.999999 with this one change.
2026-05-22 17:59:57 +00:00
bc3e94ff45 Diag: try K slice (None,None,0,0) keeping mode 1 (CUTLASS ref style) 2026-05-22 17:59:01 +00:00
601e662dd4 Diag: try runtime Int32(0+0) for kv_coord with cutlass.range 2026-05-22 17:57:58 +00:00
e5030cbea5 Diag: use Python range() unrolling like stage C test 2026-05-22 17:56:59 +00:00
5ee34d925b Fix diagnostic test: same Int32(kt) + n_kv_tiles fixes 2026-05-22 17:56:15 +00:00
d2bbdd59f6 Try cutlass.range with Int32(kt) — now n_kv_tiles is Python int 2026-05-22 17:51:25 +00:00
bf80fbee99 FIX: n_kv_tiles as Python int (s_k//128) for range() unrolling
cute.size() returns a CuTeDSL symbol, not a Python int.
range() on a symbol can't iterate — the loop never unrolls.
Now n_kv_tiles is computed in __init__ as s_k // 128 (Python int).
2026-05-22 17:50:07 +00:00
0b3bc3a16d Option 2: Python range() with Int32(kt) for TMA GMEM coord
cutlass.range traces once - kv_coord/kt are trace-time values,
not runtime loop-carried state. Python range() fully unrolls at
trace time, emitting distinct Int32(k) constants per iteration.
Int32(1) hardcoded already proved TMA CAN load from tile 1.
2026-05-22 17:47:43 +00:00
f80f8eb38f Clean up debug prints, set kv_coord as Int32(0)
Key findings to relay to CUTLASS LLM:
- kv_coord=Int32(1) hardcode CHANGES the output (TMA CAN load from different tiles)
- kv_coord=Int32(0) + kv_coord += 1 does NOT increment at runtime
  (all multi-tile outputs identical to kv_coord=0)
- kv_coord=0 (plain Python int) also doesn't work
- Pipeline handle .count doesn't work either
- The TMA GMEM tile coordinate must be dynamic at kernel runtime,
  but CuTeDSL appears to constant-fold or not propagate the increment
2026-05-22 17:39:27 +00:00
36cf1a363b DEBUG: try plain Python int kv_coord (like CUTLASS ref) 2026-05-22 17:34:30 +00:00
d95e2221c2 DEBUG: hardcode kv_coord=1 to test if TMA uses it 2026-05-22 17:32:53 +00:00
59746b46fc DEBUG: try K slice (None,0,None,0) keeping mode 2 free 2026-05-22 17:30:06 +00:00
b6cefba31c DEBUG: print tBgK/tVgV shapes before/after slice 2026-05-22 17:28:45 +00:00
ba2cefb668 Stage C: manual kv_coord + correct K GMEM slice + O rescale fence
Key fixes:
1. GMEM tile coord: manual Int32 kv_coord (not kvh.count)
2. K GMEM slice: (None,None,0,0) keeps mode 1 free (GMEM iter)
3. V GMEM slice: (None,0,None,0) keeps mode 2 free (GMEM iter)
4. Add fence_view_async_tmem_load before O rescale for visibility
2026-05-22 17:26:56 +00:00
4f6853e1ae FIX: only slice GMEM tensors (SMEM already 2D from tma_partition) 2026-05-22 16:57:31 +00:00
c61590ac6d FIX: consistent GMEM/SMEM slicing for K and V TMA partitions
Both GMEM and SMEM sides must be sliced to the same rank for cute.copy.
K (QK MMA B-partition): slice [(None,None,0,0)] keeps modes 0,1
  - mode 1 = GMEM iteration, indexed by kvh.count
V (PV MMA B-partition): slice [(None,0,None,0)] keeps modes 0,2
  - mode 2 = GMEM iteration, indexed by kvh.count
Q: only 1 tile, (None,0,None,0) hardcode is fine.
2026-05-22 16:56:38 +00:00
7aaf9ccbda FIX: keep GMEM iteration dimension FREE in TMA K/V partition slices
Root cause of multi-tile failure: (None,0,None,0) slice hardcodes the
GMEM tile dimension to 0, so TMA always loads from tile 0 regardless
of kvh.count. K from QK MMA has GMEM iter at mode 1, V from PV MMA
has it at mode 2 (different layouts: K,D,L vs D,K,L).

Fix follows CUTLASS FMHA reference:
- K: tBgK[(None,None,None,0)] + tBgK[(None, kvh.count, None)]
- V: tVgV[(None,0,None,0)] + tVgV[(None, kvh.count)]
2026-05-22 16:51:57 +00:00
04da36e18c Add diagnostic test for multi-tile TMA pipeline (identity softmax) 2026-05-22 16:47:08 +00:00
b50968dfaf FIX: acc_scale was double-multiplying by scale_log2
row_max is already in log2(scaled) space (S * scale_log2), so
old_row_max - row_max_safe is the correct exponent for exp2.
The old code computed exp2(scale_log2 * (old_row_max - row_max_safe))
which is exp2(scale_log2^2 * (old_max_S - new_max_S)) — wrong.
2026-05-22 16:42:45 +00:00
c9fe26a5fc Stage C: integrate example3 multi-tile fixes into unit test
- Combined K+V barrier (one acquire per kt, kvh.count == kt)
- O rescale for kt > 0 (online softmax O correction)
- final_o_bar sync (MMA signals before producer_tail)
- s_k as constructor param (compile-time for V layout)
- kv_tx_bytes covers both K and V transfers
- Test covers n=128, 256, 512, 1024
2026-05-22 16:39:45 +00:00
ad2a494968 Revert "debug: test 12w identity softmax with n=256 to verify multi-tile pipeline"
This reverts commit 6cf8702e3c.
2026-05-22 10:25:48 +00:00