Fix layernorm_fp8 scale-storage bug and bandwidth miscalculation

[pauliano22]

Summary

• layernorm_quant_kernel only ever stored row 0's dequantization scale. The store was gated on tl.program_id(0) == 0, but each program instance handles exactly one row (row_idx = tl.program_id(0)). Every other row's entry in the returned scales tensor is left as whatever torch.empty((M,), ...) happened to allocate — uninitialized memory. Dequantizing any row other than 0 with that scale would silently produce garbage. Removed the guard so every program stores its own row's scale.
• The FP8 vs. BF16 bandwidth comparison in bench_layernorm_fp8.py used the same byte multiplier (6) for both providers, even though the comment directly above it says FP8's output write is 1 byte instead of 2 (5 bytes/elem total vs. BF16's 6). This overstated the FP8 kernel's reported GB/s and speedup in the README by ~20%. Now uses 5 for the FP8 provider, 6 for BF16.
• Fixed plot_name in bench_layernorm_fp8.py and bench_relu.py to match the actual checked-in results/*.csv/*.png filenames — as written, re-running either script would write to a different filename than the one committed, silently diverging from the results shown in the README.
• Removed a stray self-referential line from requirements.txt: -e git+https://github.com/pauliano22/triton-gpu-kernels.git@<old-sha>#egg=triton_kernels. This looks like a leftover from running pip freeze against a local pip install -e ., and would pull a stale historical commit of the repo itself on a fresh pip install -r requirements.txt.

Test plan

• python3 -m py_compile on all changed .py files — syntax OK
• Verified the layer_norm_fp8.py fix by reading the kernel's execution model (one program per row via tl.program_id(0)) — the fix is a one-line removal of an incorrect conditional
• Not able to run these kernels in this environment — no CUDA GPU, no torch/triton installed here. Would appreciate a sanity run on the H100 box before merging, e.g. layernorm_fp8(x, w, b) on a multi-row input and checking scales has no zero/garbage rows beyond row 0's expected value.

---

Generated by Claude Code