Commit Graph

75 Commits

Author SHA1 Message Date
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
24a807eae2 debug: test 12w identity softmax with n=256 to verify multi-tile pipeline 2026-05-22 10:24:53 +00:00
572656e79b debug: disable O rescaling to test multi-tile pipeline baseline 2026-05-22 10:23:37 +00:00
8ce257150e fix: revert to scaled row_max, use exp2(old_max - new_max) for O rescaling
row_max is in scaled domain (s_val * scale_log2). The O rescaling
should be exp2(old_max - new_max) without extra scale_log2 because
the max values already include the scaling factor.
2026-05-22 10:22:44 +00:00
e85d50dc3b fix: compute row_max from RAW S values, not scaled
row_max should be the max of the raw QK scores, not pre-scaled.
The scale_log2 is applied during exp2 and rescaling, not stored in row_max.
This fixes the double-scaling bug that broke multi-tile O rescaling.
2026-05-22 10:21:50 +00:00
0bcb5aba2b fix: missing newline after self.s_k = s_k 2026-05-22 10:20:35 +00:00
1982cc4d39 fix: add s_k param to FmhaV3StageC, use self.s_k for V FMHA reconstruction 2026-05-22 10:19:49 +00:00
b80a1ab083 Stage C: add online O rescaling for multi-tile KV + test n=256
- Move O TMEM load/store setup before softmax loop
- After P store: rescale O in TMEM by exp2((old_max - new_max) * scale)
- Only rescale for kt > 0 (first tile has no prior O to rescale)
- Use same TMEM load/modify/store pattern as final normalization
- Test both n=128 (1 tile) and n=256 (2 tiles)
2026-05-22 10:19:08 +00:00
55beaeb2a5 fix: add epilogue warp to tmem_bar, restore wait_for_alloc in epilogue
The epilogue needs tmem_ptr for epilogue_tma_store.
It must be part of the tmem alloc barrier to synchronize.
2026-05-22 10:17:02 +00:00
6514888a5c fix: add softmax_done_bar to synchronize MMA PV with softmax P production
MMA must wait for softmax to produce P in TMEM before starting PV.
Without this, MMA reads stale P data from TMEM, causing deadlock.
softmax_done_bar: softmax warps arrive after P store, MMA waits before PV.
2026-05-22 10:15:26 +00:00
fdea390c71 fix: epilogue warp self-signals acc_pipe producer before consuming 2026-05-22 10:11:55 +00:00
18ab3396b7 fix: remove duplicate tmem free from epilogue (MMA warp handles dealloc) 2026-05-22 10:05:52 +00:00
1994b2ae46 fix: add acc_pipe pipeline for epilogue, matching 12w pattern
- Add acc_bar to SS struct
- Create acc_pipe (full pipeline) before if blocks
- Pass acc_pipe to epilogue_tma_store (needs full pipeline, not participant)
2026-05-22 10:03:08 +00:00
925d85820b fix: epilogue_warp_id must be tuple for epilogue_tma_store, check with [0] 2026-05-22 09:59:20 +00:00
23421bc282 fix: epilogue warp reuse mma_corr_cons pipeline instead of creating new one from st 2026-05-22 09:56:18 +00:00
5b32490b15 fix: define cS and tScS in correction warps (not visible across if blocks) 2026-05-22 09:52:59 +00:00
a5a9413aa5 fix: correct @cute.kernel indentation 2026-05-22 09:49:36 +00:00
cf900a22fe fix: remove duplicate @cute.kernel decorator 2026-05-22 09:46:09 +00:00
bfc1518046 FMHA Stage-C2: production 12-warp pipeline with correction warps
- Softmax warps (0-3): S→softmax→P, vec=[old_max,new_max]→TMEM
- Correction warps (4-7): O rescale in TMEM, final normalize by row_sum
- MMA warp (8): QK→S, PV→O with pipeline chaining
- TMA warp (9): Q/K/V load
- Epilogue warp (10): O TMEM→GMEM via epilogue_tma_store
- Empty warp (11): tmem dealloc mbar init
- Pipeline: mma_s→softmax→s_corr→correction→corr_epi→epilogue + mma_corr→correction
- Supports multi-tile KV with online O rescale
- Follows CUTLASS FMHA correction_rescale pattern exactly
2026-05-22 09:42:39 +00:00
347c107394 test: add multiple seeds to verify softmax consistency 2026-05-22 09:32:08 +00:00
d3682b0c33 fix: use plain range loop for row_max (fmax not allowed in vectorized) 2026-05-22 09:31:07 +00:00
235c7850df fix: add missing old_row_max = row_max before softmax max computation 2026-05-22 09:30:32 +00:00
35056300cb fix vectorize issue: remove vectorize from exp2 pass, add row_sum accumulation
- Remove vectorize=True from exp2 computation loop (carry variable)
- Add row_sum accumulation from P values in exp2 pass
- Compute row_max via fmax in separate pass
2026-05-22 09:29:43 +00:00
c5a504d064 fix: use cute.arch.fmax instead of if-else in vectorized loop 2026-05-22 09:28:32 +00:00
6f4bb0842e softmax: element-wise row_max computation instead of .reduce()
The .reduce() on the C-fragment gives global max across all rows,
not per-row max. Compute row_max element-wise from S values before
the exp2 pass. Also accumulate row_sum in the exp2 pass.
2026-05-22 09:27:36 +00:00
9e145c35f1 fix O normalization: use direct rmem tensor from partition_D shape 2026-05-22 09:23:58 +00:00