[SM80] Backport mHC fused post + int8 all-reduce; fixes #28, #29

[sjug]

What this fixes

Two flags documented in the README as opt-in performance knobs are broken on the shipped v0.6.6-sm80 image (e8953f6e5):

• [Bug]: VLLM_MHC_POST_FUSE_SQRSUM imports missing mhc_post_sqrsum_tilelang in v0.6.6-sm80 #28 — VLLM_MHC_POST_FUSE_SQRSUM=1 crashes at worker profiling: tilelang.py imports mhc_post_sqrsum_tilelang, which was never ported to tilelang_kernels.py. Because VLLM_MHC_PRENORM_SHARD is deliberately coupled to the fused sqrsum, it is silently inactive as well.
• [Bug]: VLLM_MHC_AR_INT8=1 drops the attention all-reduce without re-adding it in v0.6.6-sm80 (rank-local garbage output, no crash) #29 — VLLM_MHC_AR_INT8=1 suppresses the attention TP all-reduce (attention.py: reduce_results=not ar_hoisted) but nothing re-adds it: nvidia/model.py at e8953f6e5 has no _hoisted_all_reduce, no int8 TileLang consumers exist, and the native all_reduce_int8 op is absent — every output is rank-local garbage with no crash.

This PR ports the missing pieces from the A100 campaign tree: the decoder-side hoisted all-reduce (int8 above the 2048-token threshold, bf16 below), the int8/sqrsum TileLang post kernels, the native all_reduce_int8 transport, the PP x_scales handoff, and a flag-enabled test.

Test commands and results

All on 8×A100-SXM4-40GB (NVSwitch, driver 610.57.04), model deepseek-ai/DeepSeek-V4-Flash-0731, TP=8, 512K max-model-len, DSpark k=5. Image built from this branch on the target host.

• python3 -m pytest tests/models/test_deepseek_v4_input_gemm_fusion.py -v → 16 passed in 13.54 s (run on the 8×A100 target, in a container from this branch's image; includes test_hoisted_all_reduce_readds_the_suppressed_collective and test_decoder_carries_int8_scales_to_each_mhc_post)
• End-to-end A/B on the built image (each arm: greedy sanity, 32k c8/c16/c64 decode sweep, gsm8k-100 @16-shot nc=8):
  • control: sanity ✓, gsm8k 0.96, c8/c16/c64 = 165/300/334 steps/s
  • +AR_INT8: sanity ✓ (was gibberish on e8953f6e5), gsm8k 0.96, decode-neutral
  • +SQRSUM/PRENORM: boots (was ImportError), sanity ✓, gsm8k 0.95, decode-neutral
  • all three: sanity ✓, gsm8k 0.95, decode-neutral
• Cold-prefill A/B (standalone, 8k/64k): control 11,072/12,556 tok/s → all three flags 11,575/13,259 tok/s (+4.5%/+5.6%); AR_INT8 alone −20 ms TTFT@8k (campaign measured −18.75 ms on A100-80GB).
• Full serving certification of the production candidate (this branch, max-num-seqs 16, capture cap 96, 8448/8192, util 0.875): 342k-token needle retrieval ✓, gsm8k-200 @16-shot nc=8 = 0.975 ± 0.011, tool-call parser ✓, zero allocator OOM events, archive matrix c1–c16 clean.

Model evaluation

gsm8k scores above; all certified configs sit within each other's confidence intervals and match the shipped-image baseline (0.965 ± 0.013) — the recovered flags are quality-neutral.

Build note (pre-existing, NOT addressed by this PR)

Fresh cache-less CUDA-13 builds of any branch (including the current default branch) fail in the DeepEP stage before reaching this PR's changes: tools/ep_kernels/install_python_libraries.sh installs unpinned torch, which resolves cu12.x wheels under the CUDA 13.0.3 base and aborts on DeepEP's CUDA check; additionally torch 2.13.0+cu130 pins nvidia-nccl-cu13==2.28.3.post1, which no longer exists on any index. CI is shielded by its registry build-cache. This is an independent defect tracked separately; local workaround used for the verification builds here: pin torch==2.13.0 with the cu130 extra index under UV_OVERRIDE mapping nvidia-nccl-cu13 to 2.30.7.

AI assistance statement

AI assistance was used in developing and verifying this change. Every changed line was reviewed by the human submitter, and all test and benchmark results were measured on the target hardware described.

[wtdcode]

I'm rebasing upstream vllm changes, please wait a bit. Thanks in advance!