biondizzle 75288bd12f Wire prefill FMHA into production.py and single_shot ...

- Add dsv4_attention_mixed_fp8_prefill to production.py
- _run_production_fmha_mixed now dispatches to prefill kernel for T>1
- Remove decode-only T==1 restriction
- Update FINAL_STRETCH.md: prefill marked DONE, batched prefill TODO noted

2026-06-03 03:49:57 +00:00

..

__init__.py

B1: mixed FP8/BF16 decode FMHA integration

2026-06-02 22:53:14 +00:00

fmha_6warp_tma_multirow_multitile.cuh

FMHA sink: don't double-scale sink bias

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

fmha_common.cuh

docs: update here-docs with CuTeDSL rationale for NVIDIA

2026-05-28 07:54:01 +00:00

fmha_mixed_fp8_capi.cu

B1: mixed FP8/BF16 decode FMHA integration

2026-06-02 22:53:14 +00:00

fmha_mixed_fp8_decode.cuh

Fix B1 FMHA: swap V matrix canonical layout args (dd, kk) not (kk, dd)

2026-06-03 00:24:20 +00:00

fmha_mixed_fp8_op.py

B1: mixed FP8/BF16 decode FMHA integration

2026-06-02 22:53:14 +00:00

fmha_mixed_fp8_prefill_capi.cu

CRITICAL FIX: Add T-dimension strides to prefill FMHA kernel

2026-06-03 03:48:17 +00:00

fmha_mixed_fp8_prefill_op.py

CRITICAL FIX: Add T-dimension strides to prefill FMHA kernel

2026-06-03 03:48:17 +00:00

fmha_mixed_fp8_prefill.cuh

CRITICAL FIX: Add T-dimension strides to prefill FMHA kernel

2026-06-03 03:48:17 +00:00

fmha_multitile_capi.cu

FMHA sink bias in kernel + single_shot production rewrite

2026-05-31 23:10:13 +00:00

fmha_multitile_op.py

perf: skip MQA GQA expansion in FMHA (stride=0, no 128x K/V copy)

2026-06-02 03:54:03 +00:00

fmha_tma.cuh

P6: TMA store uses mbarrier completion (same as load)

2026-05-30 17:07:24 +00:00

fmha_umma_desc.cuh

B1: mixed FP8/BF16 decode FMHA integration

2026-06-02 22:53:14 +00:00

production.py

Wire prefill FMHA into production.py and single_shot

2026-06-03 03:49:57 +00:00