biondizzle

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:25:49 +00:00

11da4daa01 Debug: single PV K-tile

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:25:06 +00:00

8cb32cabc9 Fix asm constraint typo

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:24:55 +00:00

36a50962b3 Full FMHA SMEM-P with scale calibration

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:23:44 +00:00

4a36da9845 Minimal PV SS MMA test: A=128x16, B=16x16

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:22:31 +00:00

77901834a9 Fix P K-tile offset: 2048 BF16 per (128,16) tile, not 1024

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:21:43 +00:00

0bfc943cec FMHA with SMEM-P approach: PV via SS MMA avoids TMEM layout mismatch

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:19:53 +00:00

faeedd3643 Test TS MMA with non-uniform A data

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:18:43 +00:00

570c5b5154 Test softmax→PV with 1 K-tile in isolation

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:17:40 +00:00

a29ef77b64 QK→PV layout test: skip softmax to test TMEM layout compatibility

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:16:45 +00:00

acf17e001e Fix SMEM allocation (was half the needed size) + re-enable full pipeline

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:15:19 +00:00

fa6c124163 Debug: QK only, skip softmax+PV

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:14:36 +00:00

79cee32125 Debug: skip PV step entirely

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:14:00 +00:00

47e9b8a413 Debug: single PV K-tile

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:13:12 +00:00

414b3f4f92 Full FMHA HD=16 with PV GEMM via tcgen05.mma TS

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 14:10:09 +00:00

ed8f48dddf Add systematic SS+TS sequence test to debug MMA coexistence crash

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 13:53:37 +00:00

6a3159dfd9 test: PV then QK to find ordering issue

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 13:52:20 +00:00

640233cb87 test: PV GEMM first (before QK) to test ordering

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 13:50:51 +00:00

d4ed3fa06f test: QK GEMM + PV GEMM combined test

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 13:49:06 +00:00

7cb133c5bf test: exact copy of working TS test (verify it still passes)

biondizzle pushed to master at biondizzle/nvfp4-megamoe-kernel

2026-05-28 13:47:49 +00:00

0dcaa648b3 test: properly aligned V SMEM buffer