[05yuki]

Motivation

DeepSeek-V4.1-Flash-1017's reference code (inference/model.py in the checkpoint) rounds the attention output o to FP8 and back before the wo_a einsum. The rounding uses act_quant(o, 32, "ue8m0", ..., inplace=True): blocks of 32 values along each group's row, with power-of-two scales. The reference does this in the attention and in the DSpark draft's attention:

# wo_a is block-diagonal over groups (each projects only its own heads), hence einsum not
# Linear. convert.py dequantizes it to bf16; quantize/dequantize o before einsum.

The checkpoint's weights expect the rounded o. SGLang's bf16 wo_a path skips the step. That is the path 1017 takes: its weight_block_size is [32, 32], so wo_a_fp8_gemm_enabled is false and wo_a is dequantized at load. On 1017 the step changes o by 1.8 to 2.6% (relative).

Modifications

• python/sglang/srt/models/deepseek_common/o_fp8_qdq.py (new): an in-place Triton FP8 (e4m3) quantize-dequantize.
  • Matches the reference: amax floored at 1e-4, scale 2^ceil(log2(amax / 448)), clamp to ±448, round to e4m3, multiply back.
  • Uses the hardware e4m3 conversion on sm89 and newer.
  • Older GPUs get an fp32 emulation of the same rounding: round half to even onto the e4m3 grid, signed zero kept.
• deepseek_v4.py (MQALayer) and deepseek_v4_dspark.py (DSparkAttention): apply the step on the bf16 wo_a path when SGLANG_DSV4_O_FP8_QDQ_BLOCK is set.
• environ.py: two switches.
  • SGLANG_DSV4_O_FP8_QDQ_BLOCK (default 0, off).
  • SGLANG_DSV4_O_FP8_QDQ_EMULATE forces the emulated path, for checking it against the hardware conversion.

The step is opt-in, so other DeepSeek-V4 checkpoints are untouched. The earlier V4.1 reference does not have it. The checkpoint's config carries no flag for it; 1017's [32, 32] block size could switch it on automatically if you prefer that. I left it explicit.

Accuracy Tests

The kernel was compared with the checkpoint's own tilelang act_quant on bf16 tensors of the attention output's shape at T = 1, 3, 37 and 4096 tokens, including tiny and huge blocks and an all-zero block. 0 bf16 values differ, on the hardware conversion and on the emulation. 1017 runs with the step on in our setup: 2x RTX 5070 Ti, KTransformers CPU experts, TP2.

Speed Tests and Profiling

One small elementwise kernel per attention layer on the bf16 wo_a path, only when enabled.

Checklist

• Format your code according to the Format code with pre-commit.
• Add unit tests according to the Run and add unit tests.
• Update documentation according to Write documentations.
• Provide accuracy and speed benchmark results according to Test the accuracy and Benchmark the speed.
• Follow the SGLang code style guidance.

🤖 Generated with Claude Code

https://claude.ai/code/session_01C2UBZdnrDZSBh8yQz8rc9g

---

CI States

Latest PR Test (Base): ❌ Run #39262925866
Latest PR Test (Extra): ❌ Run #39262925552
Latest PR Test (AMD ROCm 10): ❌ Run #39262925672