fix: layernorm_fp8 kernel produces garbage scales for all but row 0

[pauliano22]

Summary

kernels/layer_norm_fp8.py had two compounding bugs that made the kernel completely non-functional:

1. .to(tl.float8e4m3fn) referenced a dtype that doesn't exist in triton.language — checked both the pinned triton==3.6.0 and current 3.7.1, neither has a float8e4m3fn attribute (triton/language/core.py only defines float8e4nv, float8e5, float8e5b16, float8e4b8, float8e4b15). Every call raised AttributeError immediately, so this kernel could never actually run. Fixed to tl.float8e4nv, triton's name for the NVIDIA E4M3 format (matches the torch.float8_e4m3fn dtype already used for the output tensor on the host side).
2. if tl.program_id(0) == 0: guarded the scale store, so only row 0's scale was ever written to Scale_ptr; rows 1..M-1 kept whatever uninitialized memory torch.empty happened to allocate. Scale_ptr is sized (M,) and each row already computes its own max_val/scale a few lines above, so a per-row scale was clearly the intent — the guard looks like leftover confusion (the removed comment literally said "once for the whole tensor or per row").

Verification

No test harness or CI exists for this repo's kernels, so I verified manually via TRITON_INTERPRET=1 (CPU interpreter mode) against a reference LayerNorm + per-row quantization computation:

• Before the fix: kernel raised AttributeError and could not run at all.
• With only the dtype fix (bug 2 still present): 7/8 rows returned garbage — including subnormal floats where uninitialized memory happened to look like a valid-ish number, which is what actually surfaced this on inspection.
• With both fixes: all rows match the reference computation to rtol=1e-3, atol=1e-3.

scales: tensor([0.0071, 0.0057, 0.0057, 0.0048, 0.0056, 0.0075, 0.0075, 0.0067])
expected: tensor([0.0071, 0.0057, 0.0057, 0.0048, 0.0056, 0.0075, 0.0075, 0.0067])
PASS

Test plan

• TRITON_INTERPRET=1 run against a hand-written reference implementation — passes for all rows
• Negative control: reverting the scale-store fix reproduces the original garbage-output bug under the same test, confirming the test actually discriminates
• Not run on real GPU hardware (none available in this environment) — recommend a quick real-H100 smoke test before merging, though the interpreter-mode numerics should carry over directly since both bugs are dtype-name/control-flow issues, not precision-sensitive math

---

Generated by Claude Code