[gitover22]

Motivation

Extend the CUDA TileLang raw FP8 KV path added in #42957. Kernel dispatch currently accepts packed rows and CUDA-incompatible E4M3FNUZ tensors until a lower-level failure; attention CP is not rejected by the raw-layout validation. CUDA raw KV writes also use separate casts before scattering despite an existing fused writer.

Modifications

• Reject FP8 KV rows whose width differs from the query and reject E4M3FNUZ on CUDA before kernel dispatch.
• Validate resolved attn_cp_size alongside the existing DCP constraint, retaining the existing mixed-backend and HiSparse checks.
• Select the fused cast/paged writer by raw versus scaled layout, including NoPE-only inputs with no RoPE tensor.
• Add numerical, masking, negative-control and page-boundary coverage, plus a model-accuracy report and sampled IDs/hashes. The raw layout remains unscaled; this does not add support for group-scaled packed KV rows.

The earlier raw-layout CUDA approach is documented in #36904. Already merged allocator/backend routing and the existing NoPE kernel logic are retained rather than duplicated.

Accuracy Tests

Real sampled evaluation: GLM-5.3-Flash-FP8, raw fp8_e4m3 KV, TileLang prefill/decode, TP8/EP4 + MTP, 8×H20-96GB. EvalScope 1.12.0; temperature 0, top_p 1, seed 42, native chat template/default thinking, one candidate per sample, concurrency 8.

Dataset / metric	Correct / scored	Score	Input TPS	Input TPS/GPU
IFEval strict prompt	35/40	87.50%	N/A	N/A
C-Eval validation	50/52	96.15%	N/A	N/A
GSM8K	48/50	96.00%	N/A	N/A
LiveCodeBench release_v6 pass@1	16/18	88.89%	N/A	N/A

Accuracy-only: input/output/total TPS are unavailable. IFEval loose prompt was 90.00%, strict instruction 56/61=91.80%, and loose instruction 57/61=93.44%.

Samples were selected without replacement using seed 42: IFEval by instruction-count strata (0-shot), C-Eval one per all 52 subjects (5-shot), GSM8K uniform (4-shot), and LiveCodeBench six per difficulty (0-shot; easy 6/6, medium 6/6, hard 4/6). LiveCodeBench's equal-difficulty sample is not a population-weighted full-dataset score.

All 160 selected samples were scored. One C-Eval and one LiveCodeBench response exhausted the respective 32768/65536-token output budgets without a final answer; both count as incorrect. The other failed code sample raised IndexError. No API/transport failures occurred. A sampled evaluator cache-ID mismatch was repaired and request/prediction/review identities audited; two historical duplicate requests remain in the 162 actual requests, with no correctness-based response selection.

Provenance/scope: the model evaluation was completed on runtime base 7760d5264ac16b85392aefb1bf5c13ce38d9632e with validation patch f9c6b44d964504c6ab9e0088c090aa8a10a5ff5e, not rerun on this PR's latest upstream base. There is no paired BF16 KV accuracy control or full-dataset accuracy claim. Full settings, failure treatment, source IDs and hashes are in the validation report and its machine-readable summary.

Rebased component validation: on H20 SM90, PyTorch 2.14.1+cu130 and TileLang 0.1.12, the exact committed Python files passed 15 tests and 9 subtests. Tests cover NoPE/RoPE, 8/16/32 heads, quantized-input FP32 references, one-hot and all-masked rows, wrong-index negative control, invalid layouts/formats, page boundaries and exported buffer lengths. The two CP regression assertions fail against the pre-fix upstream override code and pass with this patch.

python -m pytest -v \
  test/registered/unit/test_dsa_tilelang_fp8_validation.py \
  test/registered/kernels/ops/attention/test_tilelang_dsa_fp8_cuda.py \
  test/registered/mem_cache/test_tilelang_raw_fp8_pool.py

All changed-file pre-commit checks passed. Full-tree non-Rust checks passed; unrelated Rust hooks could not run because the local Rust toolchain is unavailable.

Speed Tests and Profiling

No new end-to-end serving benchmark or profile was run for this upstream port, and no throughput gain is claimed. The raw writer uses the existing fused cast/write kernel instead of separate casts and scatter. For a 512-dimensional NoPE row, raw FP8 stores 512 bytes versus BF16's 1024 bytes; this describes row storage, not a measured serving speedup.

Checklist

• Format code with pre-commit.
• Add focused registered numerical and runtime tests.
• Update validation documentation.
• Provide real sampled accuracy results with runtime and sampling limitations.
• Provide a fresh paired end-to-end speed benchmark/profile.
• Follow the SGLang code style guidance.

---

CI States

Latest PR Test (Base): ❌ Run #38366193208
Latest PR Test (Extra): ❌ Run #38366192837
Latest PR Test (AMD ROCm 10): ❌ Run #38366193272