Commit Graph
10 Commits
Author SHA1 Message Date
Will LinandClaude Fable 5 95f4f54712 [bugfix] Triton block-sparse backward: scale QK^T after the dot, keep K raw
The backward pre-scaled K by sm_scale*log2(e) in bf16 before recomputing
logits, so the recompute drifted from the forward's saved LSE proportionally
to |logit|; exp2 amplifies that into exponentially wrong probabilities and
garbage dK/dV (plus a mis-scaled dQ) at real activation magnitudes. Forward
was exact and unit-scale tests passed, which is how it survived — the Triton
path is only load-bearing off-Hopper (sm_90 uses the ThunderKittens CUDA
backward).

Fix: pass raw K, apply sm_scale*RCP_LN2 after tl.dot in fp32 in both
_attn_bwd_dkdv and _attn_bwd_dq, and scale dq by sm_scale instead of LN2.
Verified against SDPA to 4 significant digits at sm_scale 1/8/32 and at the
real H3 geometry (37,48,84 tiles, 128-dim heads).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
2026-08-11 04:40:58 +00:00
William Lin 9028953625 [misc] yapf pass under CI's interpreter (3.12) + pin hook language_version (#1702) 2026-08-08 22:24:17 -07:00
9709b7513b [feat] Port NVFP4 QAT/QAD to modular train framework (#1619)
Co-authored-by: Peiyuan Zhang <email>

Co-authored-by: Peiyuan <a>
2026-07-18 15:01:09 -07:00
xsankandmergify[bot] 31b719ae49 [perf] optimize compress & topk kernel (#1421)
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-06-23 18:51:22 +00:00
eed9c4bfbf [kernel] QAD 5090: Add Attn-QAT training Triton kernels (11/12) (#1460)
Co-authored-by: William Lin <SolitaryThinker@users.noreply.github.com>
Co-authored-by: Loay Rashid <42599591+loaydatrain@users.noreply.github.com>
Co-authored-by: Kaiqin Kong <k1kong@ucsd.edu>
2026-06-17 13:35:12 -07:00
Mook 7b872cc41e [Perf] Skip bool-mask round-trip in block-sparse VSA attention (#1243) 2026-04-26 15:14:37 -07:00
alexzms d925ad05f3 [bugfix] fastvideo-kernel: fix VSA Triton padding NaNs and support q/kv length mismatch (#1094) 2026-02-13 12:39:49 -08:00
Loay Rashid dea69cf651 [New Model] Turbodiffusion (#971) 2026-01-02 17:55:56 -06:00
RoyWangandroywang de264fcc99 [fix]: fix STA trition kernel for AMD RDNA archs (#969)
Co-authored-by: roywang <roywang@amd.com>
2025-12-29 14:25:32 -06:00
William LinandShreejithSG 2f6230abcf [kernel] Reorg and fix fastvideo-kernel (#962)
Co-authored-by: ShreejithSG <shreejithsg@gmail.com>
2025-12-26 01:50:24 -06:00