biondizzle a30ebfb197 FMHA SM100: Full kernel with TMET PTX, UMMA descriptors, softmax loop ...

- TMEM alloc/dealloc/load/store via inline PTX (tcgen05.*)
- UMMA SMEM descriptor construction (make_umma_desc)
- QK GEMM via tcgen05.mma.kind::f16 inline asm
- Online softmax with D3/D4/D5c masks
- O rescale in REGISTERS (D1.5 fix — no TMEM round-trip!)
- FP4 quantize helpers (hs2e2m1, fp8_e4m3_encode)
- Still needs: PV GEMM, proper P staging, TMEM O load/store

2026-05-28 05:19:34 +00:00

..

__init__.py

Stage E: head-packed MQA/GQA, batch dim, custom_op, integration API

2026-05-27 15:15:03 +00:00

fmha_backup_pre_epilog.py

D1.5: Replace TMEM round-trip normalize with correction epilog (one-way: TMEM→reg→SMEM→GMEM)

2026-05-24 00:24:24 +00:00

fmha_backup_v2.py

D1.5: Prepare for SMEM accumulator implementation

2026-05-26 21:00:41 +00:00

fmha_sm100_launch.cu

fix: rename .cpp to .cu for CUDA compilation

2026-05-28 05:16:41 +00:00

fmha_sm100.cuh

FMHA SM100: Full kernel with TMET PTX, UMMA descriptors, softmax loop

2026-05-28 05:19:34 +00:00

fmha_smem_acc.py

Stage E: production attention wrapper + Python KV merge, clean fmha_smem_acc

2026-05-27 06:34:10 +00:00

fmha.py

Revert "D1.5: WIP SMEM accumulator — framework in place, accumulation logic TODO"

2026-05-27 02:17:26 +00:00

fmha.py.backup

auto: pre-test commit

2026-05-23 20:08:31 +00:00

production.py

Stage E: head-packed MQA/GQA, batch dim, custom_op, integration API

2026-05-27 15:15:03 +00:00