biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
1,094 Commits 1 Branch 17 Tags
2b76b691cbaad7fae6ab7daa4e655ba1b2da4d8f
Commit Graph

1094 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
biondizzle
2b76b691cb fix: block_idx() returns tuple, use [1] for y 2026-05-24 23:29:59 +00:00
biondizzle
4c79e5533e D2: add multi-CTA grid with block_idx_y for Q/O head indexing 2026-05-24 23:27:38 +00:00
biondizzle
335e310c79 Update D2 status in README 2026-05-24 22:58:23 +00:00
biondizzle
e3e67c3992 NVFP4-3: enable 2-CTA UMMA when MMA tile M >= 256 (1.7-1.9x throughput) 2026-05-24 22:57:49 +00:00
biondizzle
e0339a92fc D2: revert multi-CTA grid params (using per-head launch approach instead) 2026-05-24 22:52:21 +00:00
biondizzle
a5271821a8 D2: add scale test (more heads, larger hd) 2026-05-24 22:49:44 +00:00
biondizzle
d563c93fc5 D2: add per-head launch test 2026-05-24 22:48:22 +00:00
biondizzle
9b476d87f9 fix: compare un-normalized O against un-normalized reference 2026-05-24 22:44:11 +00:00
biondizzle
0ca7b58a6a D1: fully revert LSE change back to original sfw_idx==0 guard 2026-05-24 22:41:32 +00:00
biondizzle
db353ec35a D2: add simple n_h=1 regression test 2026-05-24 22:39:25 +00:00
biondizzle
4418e04a28 D1: revert per-row LSE to sfw_idx=0 for now (debugging D2 regression) 2026-05-24 22:28:11 +00:00
biondizzle
2cc66bff68 D2: add initial multi-head test file 2026-05-24 22:26:10 +00:00
biondizzle
49e66fb6e4 D1: corrected KV merge test with proper normalized output formula 2026-05-24 22:24:27 +00:00
biondizzle
c47f648617 fix lse verify 2026-05-24 22:23:08 +00:00
biondizzle
3577e09603 D1: add LSE verification test 2026-05-24 22:22:31 +00:00
biondizzle
674c5b9c18 D1: fix per-row LSE output + add KV merge test v2 with per-row LSE 2026-05-24 22:21:51 +00:00
biondizzle
18f3274c0b D1: DEBUG - NO-OP O rescale (multiply by 1.0) to test TMEM round-trip 2026-05-24 22:19:16 +00:00
biondizzle
c33185ca0a D1: add rescale diagnostic 2026-05-24 22:18:12 +00:00
biondizzle
02edff5ac7 D1: add KV merge test using log-sum-exp (avoids TMEM round-trip) 2026-05-24 22:17:24 +00:00
biondizzle
0f30319e06 Revert "D1: move O rescale atoms outside const_expr guard (match CUTLASS pattern)" 
This reverts commit aaf21d8ac1.
 2026-05-24 22:15:38 +00:00
biondizzle
aaf21d8ac1 D1: move O rescale atoms outside const_expr guard (match CUTLASS pattern) 2026-05-24 22:07:18 +00:00
biondizzle
35a3c04e8e fix debug test 2026-05-24 22:04:51 +00:00
biondizzle
a391aa1fd3 D1: add rescale debug test 2026-05-24 22:04:20 +00:00
biondizzle
55c6903980 D1: fix O rescale identity tensor - use PV MMA shape not QK shape 2026-05-24 22:02:55 +00:00
biondizzle
f1aab1bfc1 D1: add multi-KV-tile O rescale test (s_k=256,384,512) 2026-05-24 22:00:42 +00:00
biondizzle
77b366d44b Scrub B200 password from markdown files 2026-05-24 21:52:54 +00:00
biondizzle
83506e6ad2 Add MAY_24_26_PLAN.md: next session startup plan 2026-05-24 21:50:32 +00:00
biondizzle
9435bf9653 Restore NVFP4 Precision Roadmap + add O rescale gap to D1.5 2026-05-24 21:48:58 +00:00
biondizzle
03cbd8ffa6 Add STAGE_D2.md: Multi-query grid + head packing plan 2026-05-24 21:43:04 +00:00
biondizzle
f4e0315af9 Remove obsolete STAGE_D1.3.md and SMEM_P_GUIDANCE_REQUEST.md 2026-05-24 21:41:17 +00:00
biondizzle
dadfad8f89 Docs: Update STAGE_D.md, README.md with hd=512 compilation blocker, lessons learned 2026-05-24 21:35:25 +00:00
biondizzle
a5fef69363 D1.4: Use cutlass.range(unroll=1) for k_sub loops in both TMA and MMA warps 2026-05-24 17:55:33 +00:00
biondizzle
c11ac38ceb D1.4: Remove --opt-level 0 from hd512 test (use default opt level) 2026-05-24 16:42:01 +00:00
biondizzle
b14d88f37f D1.4: Fix merge test - use use_smem_p=False for hd=256 kernel (SMEM budget) 2026-05-24 16:36:48 +00:00
biondizzle
e6c9e6c0d0 D1.4: Add external k_sub merge test for hd=512 (avoids slow in-kernel k_sub compilation) 2026-05-24 16:31:06 +00:00
biondizzle
13fcf16b14 D1.4: Use --opt-level 0 only (ptxas -j not supported, MLIR is the bottleneck) 2026-05-24 15:43:17 +00:00
biondizzle
b4da412b30 D1.4: Use options string for compile flags (--ptxas-options -j64 --opt-level 0) 2026-05-24 15:40:39 +00:00
biondizzle
4f69dffc93 D1.4: Add PtxasOptions -j64 + OptLevel(0) for faster hd=512 compilation 2026-05-24 15:36:35 +00:00
biondizzle
331ddb29b7 D1.4: Fix regression test for un-normalized O output (D5a) 2026-05-24 15:13:16 +00:00
biondizzle
25201d0c3d D1.4: Guard LSE computation with const_expr(not normalize) - fixes BF16 type mismatch in regression test 2026-05-24 15:11:39 +00:00
biondizzle
7f64a11eea D1.4: Switch k_sub from cutlass.range to Python range (unrolled at trace time) 2026-05-24 15:10:28 +00:00
biondizzle
6d7b8fed3e D1.4: Fix tTMrO placeholder - define only inside const_expr block 2026-05-24 14:23:22 +00:00
biondizzle
7a4ff959bf D1.4: Use cutlass.range loop for k_sub (reduce IR), guard O rescale with const_expr(n_kv_tiles>1) 2026-05-24 14:22:45 +00:00
biondizzle
449a6e7ede Fix: add cutlass import to test_d1_qk512 2026-05-24 14:20:32 +00:00
biondizzle
ce267909ad Fix: add cpasync import to test_d1_qk512 2026-05-24 14:20:01 +00:00
biondizzle
625837fd44 D1.4: Add hd=512 QK-only and standalone test for compilation debugging 2026-05-24 14:19:26 +00:00
biondizzle
592873b560 D1.4: Reduce pv_n_tile to 128 for hd=512 to fit SMEM budget (192KB) 2026-05-24 08:07:32 +00:00
biondizzle
e7c146dbfd D1: Unrolled k_sub path (hardcoded k_sub=0,1) to avoid cutlass.range IR explosion 2026-05-24 07:03:14 +00:00
biondizzle
dd39c2ebdf D1: Use cutlass.range for k_sub loops (CuTeDSL immutable handle) 2026-05-24 06:43:30 +00:00
biondizzle
2bf3ee40aa D1: Fix kvh scoping - define before loops, consume V via pipeline 2026-05-24 06:42:26 +00:00