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>
