[Archive #39210 2026-10-01] NVFP4 W4A16 MegaMoE before rebase

[zianglih]

Motivation

@HumansAnd

Add opt-in NVFP4 W4A16 to the existing FlashInfer MegaMoE integration. Reuse
SGLANG_FLASHINFER_CUTEDSL_NVFP4_W4A16=1 for serialized ModelOpt NVFP4 and
--quantization nvfp4_online, keeping expert activations and EP transport in BF16.

• Dependency: FlashInfer #5019
  merged as a03f2205263d4e691d68e485bff287e37a19b6c3 on 2026-09-28 and is
  not yet in a release. Use a source build containing that merge; no package
  version bump is made here. The adapter API was checked against the merged
  source. Historical GPU results below retain their original FlashInfer pins.
• Structure: extend the existing W4A4 MegaMoE adapter, following the W4A16
  opt-in pattern from #35120.

Modifications

• Preparation: use FlashInfer's W4A16 preprocessing and
  Sm100_Bf16_Nvfp4_Bf16_Cutedsl_MegaMoeConfig. Preserve packed weights, block
  scales, and separate per-expert FP32 weight scales. Reuse the existing layer
  construction and fused runner. Use standard (weight, block_scale) pairs and
  stage per-expert FP32 weight scales each forward; FlashInfer owns reference-counted
  workspace pooling. Remove SGLang's private workspace assignments/cache.
• Activations: neutralize checkpoint activation scales and supply BF16
  inputs without activation-scale tensors. Online weight quantization retains
  its source-precision/exclusion semantics; the selected runner bypasses
  activation quantization.
• W4A16 adapter contract: SM100/SM103, BF16 model dtype, gated SiLU/SwiGLU or SiTU, post-FC2 routing,
  BF16 combine and DP=TP=EP with DP attention.
  SGLANG_FLASHINFER_MEGAMOE_IN_KERNEL_FC2_REDUCE=1 now permits BF16 atomic
  FC2 reduction; tuning may still select external reduction. Default remains
  disabled (external FP32 accumulation). Reject non-BF16 W4A16 combine.
  Nemotron ReLU-squared experts remain unsupported.
• Reload: restore canonical loading views over the prepared allocations,
  then copy newly prepared weights back into the same buffers. Apply this
  W4A16 hook and retain its live MegaMoE instance so captured weight/scale
  addresses and the shared workspace remain stable across /update_weights_from_disk.
• GLM-5.2 + MTP: admit GlmMoeDsaForCausalLM in the existing MegaMoE
  architecture guard. Preserve full DP token counts and synchronized graph
  buckets across ranks for profile selection and for a target using a
  standard-communication MTP draft. Expert inputs stay local.
• Kimi K3: wire SiTU for NVFP4 W4A4/W4A16 using existing GEMM1 beta fields,
  separate from SwiGLU clamps. Admit both K3 architectures; use standard TopK and
  local-token EP dispatch/shared-expert handling. Keep DeepGEMM separate.
• Autotuning: tune once per geometry/capacity at startup with pooled public
  workspaces. Cover captured/eager decode bounds and adaptive MTP verify width,
  plus one DP-adjusted prefill chunk (max_prefill_tokens without chunking).
  Larger decode shapes use that profile; reject overflow. Prepare before capture.
  Reuse FlashInfer's namespaced cache: load/publish once per startup context,
  stage once across layers, cache geometry keys and binary-search capacities.
  No separate sidecar; SGLANG_FLASHINFER_AUTOTUNE_EXTEND is unchanged.
• Dummy loading: set floating-point weight_scale_2 to one alongside inverse
  scales, keeping gate/up compatible and ordinary weights random. Checkpoint
  loading/updates and quantization interfaces are unchanged.
• Tests: retain focused adapter/storage, cache/profile and adaptive decode
  bound regressions, plus one dummy-scale test covering FC1/FC2 neutral scales
  and random ordinary weights. Qwen3 online GSM8K and serialized NVFP4 reload
  use external FP32 combine (SGLANG_FLASHINFER_MEGAMOE_IN_KERNEL_FC2_REDUCE=0);
  the online-quantized RL reload fixture is excluded.
• Scope: no new environment variable, kernel or package pin. W4A4/MXFP8 use
  enable_in_kernel_fc2_reduce and native workspace pooling. No drafter MegaMoE
  or deterministic-mode special handling. Prior W4A4 GPU coverage is dummy
  startup/generation; full-model/native reuse coverage is W4A16. SiTU checks are
  CPU-only. Native tuning requires EP=WORLD for the upstream collectives.

Example on four Blackwell GPUs (requires FlashInfer sgl-project#5019):

SGLANG_FLASHINFER_CUTEDSL_NVFP4_W4A16=1 python -m sglang.launch_server \
  --model-path Qwen/Qwen3-30B-A3B-Instruct-2507-FP8 \
  --quantization nvfp4_online --dtype bfloat16 \
  --moe-runner-backend flashinfer_megamoe \
  --moe-a2a-backend flashinfer_megamoe \
  --tp-size 4 --dp-size 4 --ep-size 4 --enable-dp-attention

Rebase onto current main (2026-09-28)

Rebased to 6dad1524901dd070d8720f88b58bccfc5f9579f7 at baseline
05627ffec57fe5fef1a4a1619115ef0df6f0aa15; head is recorded in the SiTU section.
Preserve upstream's explicit
MoE backend selection and separate target/draft DP metadata, while retaining
MegaMoE's synchronized counts for capacity-profile selection. Preserve direct
runner rank fields and the physical-GPU cache suffix. The GPU validation
below pins the rebased head and source-built FlashInfer before SiTU wiring;
historical serving results retain their original tested revisions.

SiTU wiring validation (2026-09-29)

Head 561ad447c74bb757a40677ee9ce038f9ca429d2c: macOS/Python 3.12.12,
CPU Torch 2.14.0. Adapter checks cover both precisions, activation constants,
scale storage and W4A16 reload. Serving imports and kernel calls are stubbed.
Eight real FlashInfer config constructions at 90a709ca842996e3431ece0346d97b73693eff94
also passed; activation changes produce distinct tuning-cache keys. Scoped
pre-commit/diff checks passed. No full-model E2E was added or run; prior GPU
results below do not validate K3/SiTU.

From the SGLang checkout:

python - <<'PYTEST'
import pathlib, sys, types, pytest
p = types.ModuleType("sglang")
p.__path__ = [str(pathlib.Path("python/sglang").resolve())]
sys.modules["sglang"] = p
raise SystemExit(pytest.main([
    "test/registered/unit/layers/moe/test_flashinfer_megamoe.py", "-q"]))
PYTEST

26 passed, 1 warning in 0.71s

The existing warning is pytest's unknown asyncio_mode option.

GPU validation before SiTU wiring

• Sources: SGLang 9d38e0530a1e35d1756a7fabf044bc39b77209b8 and FlashInfer
  294fe64b5ce9b7902f842f0dd7f31c0285823d3e (source package reports 0.7.0).
  Both checkouts were clean; file hashes matched. FlashInfer includes merged Correct grafana heatmap. sgl-project/sglang#5019.
• Environment: C2, eight NVIDIA B300 GPUs allocated; unit tests expose GPU 0.
  Image lmsysorg/sglang:nightly-dev-cu13-20260929-79cafec0, driver 590.48.01,
  Python 3.12.3, Torch 2.13.0+cu130/CUDA 13.0, Triton 3.7.1, NVSHMEM 3.4.5.
  Image digest: sha256:0f075735a6bf913a7cc1cd7ece4f526ae8af70fc5b88fdc02d9efe92a61cf3ae.
  Image Torch, CUDA, Triton and
  NVSHMEM packages were preserved. Removed the three image FlashInfer packages,
  installed CuTe DSL and CUDA-13 libraries 4.7.1 plus nccl-extensions==0.1.0,
  and built FlashInfer editable from the pinned source.
• Unit result: the new dummy-scale test fails at the FC1 unit-scale assertion
  with only weight_utils.py restored from 05627ffec57fe5fef1a4a1619115ef0df6f0aa15.
  Restoring the fix gives 56 passed, 29 subtests passed, 15 warnings, with no
  failures, errors or skips. The warning classes are an unknown pytest
  asyncio_mode option and PyTorch script_method deprecation. These registered
  unit suites mock kernel calls; actual native/model results follow.

Reproduce source installation and registered unit suites

Inside the image above, use independent checkouts at the exact source revisions:

export WORKDIR="$PWD"
git clone https://github.com/zianglih/sglang.git "$WORKDIR/sglang"
git -C "$WORKDIR/sglang" checkout --detach 9d38e0530a1e35d1756a7fabf044bc39b77209b8
git clone https://github.com/flashinfer-ai/flashinfer.git "$WORKDIR/flashinfer"
git -C "$WORKDIR/flashinfer" checkout --detach 294fe64b5ce9b7902f842f0dd7f31c0285823d3e
git -C "$WORKDIR/flashinfer" submodule update --init --recursive
python -m pip uninstall -y flashinfer-python flashinfer-cubin flashinfer-jit-cache
python -m pip install --no-deps nvidia-cutlass-dsl==4.7.1 \
  nvidia-cutlass-dsl-libs-cu13==4.7.1 nccl-extensions==0.1.0
(cd "$WORKDIR/flashinfer" && \
  FLASHINFER_BUILD_NO_PIP=1 BUILD_NVEP=0 BUILD_NIXL_EP=0 BUILD_NCCL_EP=0 \
  python -m pip install --no-deps --no-build-isolation -v -e .)
cd "$WORKDIR/sglang"
export PYTHONPATH="$PWD/python"
export PYTHONDONTWRITEBYTECODE=1 PYTHONUNBUFFERED=1
export OMP_NUM_THREADS=1 CUDA_VISIBLE_DEVICES=0
python -B -m pytest -q -ra \
  test/registered/unit/layers/moe/test_flashinfer_megamoe.py \
  test/registered/unit/model_executor/runner/test_flashinfer_autotune_sync.py \
  test/registered/unit/managers/scheduler_components/test_dp_attn.py \
  test/registered/unit/server_args/test_server_args.py::TestFlashinferMegaMoeConfig \
  test/registered/unit/server_args/test_server_args.py::TestCutedslMoeMaxNumTokens \
  test/registered/unit/model_loader/test_weight_utils.py \
  --junitxml=focused-suites.xml

For the expected failing control, keep the final test file and substitute only
the old initializer, then restore the exact fixed file before continuing:

(
  weight_utils=python/sglang/srt/model_loader/weight_utils.py
  cp "$weight_utils" "$WORKDIR/weight_utils.fixed.py"
  trap 'cp "$WORKDIR/weight_utils.fixed.py" "$weight_utils"' EXIT
  git show 05627ffec57fe5fef1a4a1619115ef0df6f0aa15:"$weight_utils" > "$weight_utils"
  python -B -m pytest -q -ra \
    test/registered/unit/model_loader/test_weight_utils.py::TestInitializeDummyWeights \
    --junitxml=dummy-before-fix.xml
)

Unit output: failing assertion/result footer and complete fixed-suite output

Expected failing control, exit 1. Repeated traceback output and the warning summary
(identical to the fixed run) are omitted; exact assertion and result follow:

  File "/hai-workspace/sglang-megamoe-rebase-20260928/sglang/test/registered/unit/model_loader/test_weight_utils.py", line 218, in test_nvfp4_dummy_scales_are_neutral
    self.assertTrue(torch.equal(model.w13_weight_scale_2, torch.ones(4, 2)))
  File "/usr/lib/python3.12/unittest/case.py", line 727, in assertTrue
    raise self.failureException(msg)
AssertionError: False is not true

=========================== short test summary info ============================
FAILED test/registered/unit/model_loader/test_weight_utils.py::TestInitializeDummyWeights::test_nvfp4_dummy_scales_are_neutral
1 failed, 15 warnings in 31.25s

Fixed suites, exit 0:

........................................................                                                            [100%]
=============================== warnings summary ===============================
../../../opt/sglang/lib/python3.12/site-packages/_pytest/config/__init__.py:1464
  /opt/sglang/lib/python3.12/site-packages/_pytest/config/__init__.py:1464: PytestConfigWarning: Unknown config option: asyncio_mode
  
    self._warn_or_fail_if_strict(f"Unknown config option: {key}\n")

../../../opt/sglang/lib/python3.12/site-packages/torch/jit/_script.py:365: 14 warnings
  /opt/sglang/lib/python3.12/site-packages/torch/jit/_script.py:365: DeprecationWarning: `torch.jit.script_method` is deprecated. Please switch to `torch.compile` or `torch.export`.
    warnings.warn(

-- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html
56 passed, 15 warnings, 29 subtests passed in 142.37s (0:02:22)

• Native EP4: both launches exited 0; all eight rank reports passed. Real
  W4A16 kernels, native tuning, graphs and workspace reuse ran with synthetic
  scheduler metadata. Control outputs came from independent fixed-default native
  sessions, not a PyTorch oracle. Cold/warm AUTO calls were 5/0 and graph replays
  were 12 per rank. No timing or performance comparison is claimed. An initial
  hostname-rendezvous attempt stopped before workers; successful runs used 127.0.0.1.

Native EP4 commands and all rank results

Harness setup/smoke_autotune_gpu.py SHA256:
88338c2e2ad1d077e7d51a24f2e8a6df2ed0bdfa9501435c38d763bf4087aadc.

cd "$WORKDIR"
export PYTHONPATH="$WORKDIR/sglang/python:$WORKDIR/flashinfer"
export OMP_NUM_THREADS=1 CUDA_VISIBLE_DEVICES=0,1,2,3
export FLASHINFER_CUDA_ARCH_LIST=10.3a FLASHINFER_NVCC_THREADS=4 MAX_JOBS=4
export FLASHINFER_WORKSPACE_BASE="$WORKDIR/cache/flashinfer"
# This synthetic harness bypasses ServerArgs; set the production defaults explicitly.
export NVSHMEM_REMOTE_TRANSPORT=none NVSHMEM_IB_ENABLE_IBGDA=0
export NVSHMEM_DISABLE_LOCAL_ONLY_PROXY=1
for reuse in 0 1; do
  timeout --kill-after=30s 30m torchrun --nnodes=1 --nproc-per-node=4 \
    --master-addr=127.0.0.1 --master-port=29546 \
    "$WORKDIR/setup/smoke_autotune_gpu.py" \
    --cache "$WORKDIR/results/native-autotune-cache" \
    --reuse-cache "$reuse" --decode-num-tokens 16
done

Complete JSON fields shared by every rank report:

{"passed":true,"capacities":[1,2,4,8,16],"bootstrap_capacity":4,"decode_num_tokens":16,"prefill_num_tokens":8,"startup_contexts":2,"checked_counts":[[0,1,1,2],[0,1,2,7],[8,8,8,8],[0,4,9,16]],"graph_replays":12,"bit_exact":true,"pool_refcounts":[2,1,0]}

Remaining fields, one row per original report:

reuse_cache	rank	auto_calls
false	0	5
false	1	5
false	2	5
false	3	5
true	0	0
true	1	0
true	2	0
true	3	0

Archive SHA256: a8df993571728eba3178b21e82a4c933c93600f7087e6d467af41b03e8f994a3.

Actual model checks, using the source/image pins above and TP=DP=EP=4:
online GSM8K passed 190/200 (0.950) with BF16 KV; serialized NVFP4 reload
also passed with FP8 E4M3 KV. Both use W4A16 with BF16 expert activations/EP transport.
Online dense/excluded paths retain their FP8 source semantics.
Online: 200 examples, 128 threads, completion API, 512-token cap,
score > 0.90, original 4over6 recipe. All four schedulers received the
NVSHMEM defaults from an unset parent environment.
These are correctness runs; concurrent jobs occupied disjoint GPUs, so the emitted
latency/throughput are not controlled performance results.

Raw online metrics and result footer:

Total latency: 10.778 s
Score: 0.950
Output throughput: 2664.291 token/s
metrics={'score:std': np.float64(0.21794494717703372), 'score': np.float64(0.95), 'latency': 10.77847851999104, 'output_throughput': 2664.29069248856}
Ran 1 test in 1122.469s
OK
REGISTERED_E2E_RESULT={"case": "online-gsm8k", "passed": true, "tests_run": 1, "failures": 0, "errors": 0, "skipped": 0}

Reload passed both flush_cache=True/False updates and memory release/resume;
text/token IDs remained identical and logprobs satisfied the original 1e-4 tolerance.

Ran 1 test in 853.822s

OK
REGISTERED_E2E_RESULT={"case": "serialized-reload", "passed": true, "tests_run": 1, "failures": 0, "errors": 0, "skipped": 0}

After source installation, run the original classes; only pinned snapshot
paths and URL/ports change:

cd "$WORKDIR/sglang"
unset NVSHMEM_REMOTE_TRANSPORT NVSHMEM_IB_ENABLE_IBGDA NVSHMEM_DISABLE_LOCAL_ONLY_PROXY
cat > "$WORKDIR/run_model_test.py" <<'PY'
import importlib.util, sys, unittest
from huggingface_hub import snapshot_download
cases = {
    "online": ("backends/test_flashinfer_nvfp4_online_moe_backend.py",
               "TestFlashinferMegaMoeBackendNvFp4OnlineW4A16",
               "Qwen/Qwen3-30B-A3B-Instruct-2507-FP8",
               "5a5a776300a41aaa681dd7ff0106608ef2bc90db", 31000),
    "reload": ("rl/test_update_weights_from_disk_blackwell.py",
               "TestServerUpdateWeightsFromDiskNVFP4W4A16MegaMoE",
               "nvidia/Qwen3-30B-A3B-NVFP4",
               "2538ded2a4edb247b4d2b4a8ba24e44bd4c017c3", 31010),
}
file, name, model, revision, port = cases[sys.argv[1]]
spec = importlib.util.spec_from_file_location("registered", "test/registered/" + file)
module = importlib.util.module_from_spec(spec)
spec.loader.exec_module(module)
case = getattr(module, name)
assert case.model == model
case.model = snapshot_download(model, revision=revision)
module.DEFAULT_URL_FOR_TEST = case.base_url = f"http://127.0.0.1:{port}"
result = unittest.TextTestRunner(verbosity=2).run(
    unittest.defaultTestLoader.loadTestsFromTestCase(case))
raise SystemExit(not result.wasSuccessful())
PY
CUDA_VISIBLE_DEVICES=4,5,6,7 python "$WORKDIR/run_model_test.py" online
CUDA_VISIBLE_DEVICES=4,5,6,7 python "$WORKDIR/run_model_test.py" reload

Dummy loader: W4A16 and W4A4 both passed (exit 0) on GPUs 0–3 with two randomly
initialized layers, one request per DP rank, eight output tokens/request and
finite output logprobs. Both native kernels tuned and captured decode graphs on
all four ranks; every scheduler received the production NVSHMEM defaults.
This checks startup/generation, not accuracy or performance.

{"w4a16":"1","passed":true,"requests":4,"completion_tokens":32,"all_output_logprobs_finite":true,"layers":2}
{"w4a16":"0","passed":true,"requests":4,"completion_tokens":32,"all_output_logprobs_finite":true,"layers":2}

Dummy launch function (1: W4A16; 0: W4A4):

unset NVSHMEM_REMOTE_TRANSPORT NVSHMEM_IB_ENABLE_IBGDA NVSHMEM_DISABLE_LOCAL_ONLY_PROXY
export CUDA_VISIBLE_DEVICES=0,1,2,3
export SGLANG_FLASHINFER_MEGAMOE_IN_KERNEL_FC2_REDUCE=0
export SGLANG_FLASHINFER_MEGAMOE_COMBINE_DTYPE=bf16
export SGLANG_FLASHINFER_NVFP4_PER_TOKEN_ACTIVATION=0
launch_dummy() {
  SGLANG_FLASHINFER_CUTEDSL_NVFP4_W4A16="$1" python -m sglang.launch_server \
  --model-path nvidia/Qwen3-30B-A3B-NVFP4 \
  --revision 2538ded2a4edb247b4d2b4a8ba24e44bd4c017c3 \
  --load-format dummy --json-model-override-args '{"num_hidden_layers":2}' \
  --dtype bfloat16 --tp-size 4 --dp-size 4 --ep-size 4 --enable-dp-attention \
  --fp4-gemm-backend flashinfer_cutedsl --moe-runner-backend flashinfer_megamoe \
  --moe-a2a-backend flashinfer_megamoe --context-length 2048 \
  --chunked-prefill-size 256 --max-running-requests 16 \
  --cuda-graph-max-bs-decode 8 --cuda-graph-backend-prefill disabled \
  --mem-fraction-static 0.4 --host 127.0.0.1 --port "$2" --device cuda
}
launch_dummy 1 31020

After stopping that server, run launch_dummy 0 31030 for W4A4.

No new serving-throughput comparison is reported. All historical benchmark pins
and raw results below remain unchanged.

Concurrent-engine autotune cache isolation

• Fix: 028ae083d2 gives target MegaMoE each physical GPU its own rank-cache filename. Colocated engines on disjoint GPU allocations can tune concurrently without overwriting one another's namespaced records; logical-device remapping preserves the physical identity, and restarts on the same GPUs reuse their files. This follows SGLang's existing device-scoped P2P cache pattern through current_platform.get_device_uuid.
• Previous upstream integration: merge commit 08084aad0e merges main 35eb7cf8d6 without rewriting PR history. It honors #40343 by using the runner's direct rank fields and preserves #39678's target/draft cache load/save lifecycle. Main still uses rank-only filenames, so the separate physical-GPU suffix remains necessary for concurrent engines.
• Overhead and scope: startup path construction only; no new locks, collectives, polling or forward-path work. Cache-disabled runs add a UUID to avoid same-second output collisions. Other backends and draft filenames are unchanged. Old rank-only MegaMoE files remain untouched and the new layout starts cold. Engines sharing physical GPUs must still use separate SGLANG_CACHE_DIR values.
• Validation: macOS arm64, Python 3.12.12, CPU Torch 2.14.0. At merge head 08084aad0e, the actual production module passed the existing autotune CPU checks plus four local isolation checks, including four overlapping process writers and a fresh-process warm restart. Serving/GPU imports were stubbed; native Gloo/filesystem/process operations were real. Existing adapter checks and scoped pre-commit passed. No new test file or documentation change is included; the small same-second regression extends the existing autotune suite. Concurrent GPU serving engines remain untested; the latest native checks cover cold/warm reuse and two contexts per process. Historical GPU results below retain their tested revisions.

Autotune + local isolation: Ran 16 tests in 3.079s; OK (skipped=1, CUDA-only)
Adapter: 20 passed, 1 warning in 1.16s
Scoped pre-commit and git diff --check: passed

Single-node NVSHMEM startup workaround

• Change: supply the following defaults before NVSHMEM initialization for all supported single-node FlashInfer MegaMoE quantization formats, with either no speculative decoding or draft A2A none. This extends the earlier W4A16-only workaround; existing explicit environment values are preserved with setdefault.

  NVSHMEM_REMOTE_TRANSPORT=none
  NVSHMEM_IB_ENABLE_IBGDA=0
  NVSHMEM_DISABLE_LOCAL_ONLY_PROXY=1

• Reason: the source comment links the upstream initialization fix NVIDIA/nvshmem#110. NVSHMEM 3.4.5 allocates its proxy global-exit request/code with cudaMallocHost without initializing the integers. Our four-rank trace observed request 2 before publication and the proxy taking the global-exit path; matched source builds reproduced the startup failure without initialization and completed real GLM-5.2 C8 startup with the two initializations. The supported P2P-only configuration above instead skips proxy creation before those allocations, retaining the installed NVSHMEM wheel. Proxy selection, allocation and early return.

• Limits: all EP GPUs must support direct peer memory access. Device-side NVSHMEM global exit and proxy-based timeout polling are disabled; remote IB/UCX transport is unavailable. This does not change NCCL or UID bootstrap networking. Multi-node jobs and drafts using other A2A backends receive no automatic defaults. Explicit overrides may re-enable the affected proxy path. NVSHMEM option semantics.

• Validation: actual NVFP4 W4A16 and W4A4 dummy-model startup, native tuning, CUDA graphs and generation passed with these defaults; registered online W4A16 GSM8K also passed. No dedicated NVSHMEM behavior tests were added. Multi-node operation, draft A2A and other quantization formats were not tested; historical benchmark numbers retain their original revisions.

Accuracy Tests

GPU-tested source: 6d8a58f177a1488e10e8567a3fa4b4fccd0a26e2; FlashInfer source:
ad0a5e5e78e57070ec7c582efe733cb55cd8839f (FlashInfer sgl-project#5019). Kernel code is unchanged.

The earlier GPU-tested revision was rebased onto upstream
2fa6b94e3440d28dab7460246f0b34426a4e4b80. Preserve upstream's manual generic
cache load/save and model-owned prefill tuning; the Mega namespace publishes
before the shared cache is saved. FlashInfer is unchanged; model_runner.py
and its unit test have no PR diff. Whole-PR pre-commit and GitHub lint passed.
GitHub GPU CI is blocked by the PR's draft status; GPU results below are from C2.

• Environment: C2 hu-pdx-54, B300 SM103, four GPUs per run (TP4/DP4/EP4).
  Image lmsysorg/sglang:nightly-dev-cu13-20260918-20518d85, digest
  sha256:d46a59f4b98658f728a1e006c003ad5ee0628e999fd8b2bef71ac1bb61b814da.
  Driver 590.48.01; Python 3.12.3; image Torch 2.13.0+cu130/CUDA13.0 retained;
  CuTe DSL 4.7.1/compiler CUDA 13.3; NVSHMEM 3.4.5/Python 0.3.1;
  apache-tvm-ffi 0.1.11.
  Removed incompatible image FlashInfer cubin/JIT-cache packages; imported
  the pinned FlashInfer 0.7.0 source and this SGLang checkout.
  Model tests ran concurrently on disjoint four-GPU groups; their timings are
  accuracy-test output, not controlled performance measurements.
• Contract check: SGLANG_FLASHINFER_AUTOTUNE_EXTEND=0 and
  SGLANG_FLASHINFER_MEGAMOE_IN_KERNEL_FC2_REDUCE=0 throughout. Model startup
  prepared the 4096-token profile before capture; no generic extra EXTEND pass.
  Online-model capacities are powers of two through 2048, then 3839 and 4096;
  serialized-model capacities are powers of two through 4096. Separate cold/warm
  native runs exercise actual tuning and cache reuse; all native candidates remain enabled.

Check	Raw result
Adapter / runner / DP / server args	20 / 12 / 7 / 278 passed, plus 96 server-args subtests; zero failures/skips
Native EP4, decode 2 / prefill 8 / bootstrap 4	All four ranks passed; cold/warm AUTO calls 3/0; nine graph replays/rank per run; zero bit mismatches
Native EP4, decode 16 / prefill 8 / bootstrap 4	All four ranks passed; cold/warm AUTO calls 5/0; twelve graph replays/rank per run; zero bit mismatches
Online W4A16 Qwen3 GSM8K, 200 examples, 128 threads, 512 output-token cap	192/200 = 0.960; score > 0.90 unchanged; one test passed
Serialized NVFP4 Qwen3 reload, both cache-flush modes	One registered test passed, both flush-cache variants; text/token IDs unchanged and existing logprob tolerance retained

Native integration uses two layers with the same geometry in separate startup
contexts. It checks empty/uneven ranks, bit-exact outputs against independent
fixed-default sessions, graph replay with runtime scale updates, and pooled
reference counts 2→1→0. Target adaptive/eager sizing has a unit regression:
graph batch 128/eager batch 256 at maximum width 6 requires 1536 tuned rows.
The Qwen model tests are non-MTP; no new full-model MTP performance claim is made.

Reproduce the registered tests

Use the pinned image and independent SGLang/FlashInfer checkouts at the commits
above. Preserve image Torch; install FlashInfer from source without its optional
C++ EP backends, and select the checkouts through PYTHONPATH:

export WORKDIR="$PWD"
python -m pip uninstall -y flashinfer-cubin flashinfer-jit-cache
python -m pip install --no-deps nvidia-cutlass-dsl==4.7.1 \
  nvidia-cutlass-dsl-libs-base==4.7.1 nvidia-cutlass-dsl-libs-core==4.7.1 \
  nvidia-cutlass-dsl-libs-cu13==4.7.1
(cd flashinfer && git submodule update --init --depth=1 \
  3rdparty/cutlass 3rdparty/cccl 3rdparty/spdlog && \
  BUILD_NVEP=0 BUILD_NIXL_EP=0 BUILD_NCCL_EP=0 \
  python -m pip install --no-deps --no-build-isolation -e .)
export PYTHONPATH="$WORKDIR/sglang/python:$WORKDIR/flashinfer"
export FLASHINFER_CUDA_ARCH_LIST=10.3a
export CUTE_DSL_CACHE_DIR="$WORKDIR/cute-cache"
export FLASHINFER_WORKSPACE_BASE="$WORKDIR"
export SGLANG_CACHE_DIR="$WORKDIR/sglang-cache"
export HF_HOME="$WORKDIR/hf-home"
export CUDA_VISIBLE_DEVICES=0,1,2,3 OMP_NUM_THREADS=1 MAX_JOBS=4
export SGLANG_FLASHINFER_AUTOTUNE_EXTEND=0
export SGLANG_FLASHINFER_AUTOTUNE_CACHE=1
export SGLANG_FLASHINFER_MEGAMOE_IN_KERNEL_FC2_REDUCE=0
cd "$WORKDIR/sglang"
python -m pytest test/registered/unit/layers/moe/test_flashinfer_megamoe.py -v
python test/registered/unit/model_executor/runner/test_flashinfer_autotune_sync.py -v
python test/registered/unit/managers/scheduler_components/test_dp_attn.py -v
python test/registered/unit/model_executor/test_model_runner_decode_rows.py -v
python -m pytest test/registered/unit/server_args/test_server_args.py -v

The model wrapper only selects the registered class, an unused localhost port
and the exact cached snapshot below; launch flags, workloads and assertions stay
in the registered test. Run sequentially on the same four visible GPUs:

import importlib.util
import socket
import unittest
from huggingface_hub import snapshot_download

cases = [
    ("test/registered/backends/test_flashinfer_nvfp4_online_moe_backend.py",
     "TestFlashinferMegaMoeBackendNvFp4OnlineW4A16",
     "Qwen/Qwen3-30B-A3B-Instruct-2507-FP8",
     "5a5a776300a41aaa681dd7ff0106608ef2bc90db"),
    ("test/registered/rl/test_update_weights_from_disk_blackwell.py",
     "TestServerUpdateWeightsFromDiskNVFP4W4A16MegaMoE",
     "nvidia/Qwen3-30B-A3B-NVFP4",
     "2538ded2a4edb247b4d2b4a8ba24e44bd4c017c3"),
]
for path, name, model, revision in cases:
    spec = importlib.util.spec_from_file_location("registered_validation", path)
    module = importlib.util.module_from_spec(spec)
    spec.loader.exec_module(module)
    with socket.socket() as sock:
        sock.bind(("127.0.0.1", 0))
        port = sock.getsockname()[1]
    module.DEFAULT_URL_FOR_TEST = f"http://127.0.0.1:{port}"
    case = getattr(module, name)
    case.model = snapshot_download(model, revision=revision)
    case.base_url = module.DEFAULT_URL_FOR_TEST
    result = unittest.TextTestRunner(verbosity=2).run(
        unittest.defaultTestLoader.loadTestsFromTestCase(case)
    )
    assert result.testsRun > 0 and result.wasSuccessful() and not result.skipped

Complete final test summaries and native rank results

unit-adapter:
======================== 20 passed, 1 warning in 10.28s ========================
unit-runner:
Ran 12 tests in 46.539s
OK
unit-dp:
Ran 7 tests in 0.007s
OK
unit-server-args:
============ 278 passed, 16 warnings, 96 subtests passed in 19.74s =============
GSM8K:
Total latency: 9.138 s
Score: 0.960
Output throughput: 3139.069 token/s
Ran 1 test in 1090.604s
OK
REGISTERED_E2E_RESULT={"case": "online-gsm8k", "passed": true, "tests_run": 1, "failures": 0, "errors": 0, "skipped": 0}
serialized reload:
Ran 1 test in 1083.968s
OK
REGISTERED_E2E_RESULT={"case": "serialized-reload", "passed": true, "tests_run": 1, "failures": 0, "errors": 0, "skipped": 0}

Native emitted fields (all 16 reports): bootstrap capacity4, prefill8 and
two startup contexts for every row. Decode2 uses token-count vectors
[0,1,1,2], [0,1,2,7], [8,8,8,8]; decode16 also checks [0,4,9,16].

Run	Rank	AUTO calls	Reuse cache	Capacities	Graph replays	Passed	Bit exact	Pool refs
smoke-decode2-cold	0	3	false	[1, 2, 8]	9	true	true	[2, 1, 0]
smoke-decode2-cold	1	3	false	[1, 2, 8]	9	true	true	[2, 1, 0]
smoke-decode2-cold	2	3	false	[1, 2, 8]	9	true	true	[2, 1, 0]
smoke-decode2-cold	3	3	false	[1, 2, 8]	9	true	true	[2, 1, 0]
smoke-decode2-warm	0	0	true	[1, 2, 8]	9	true	true	[2, 1, 0]
smoke-decode2-warm	1	0	true	[1, 2, 8]	9	true	true	[2, 1, 0]
smoke-decode2-warm	2	0	true	[1, 2, 8]	9	true	true	[2, 1, 0]
smoke-decode2-warm	3	0	true	[1, 2, 8]	9	true	true	[2, 1, 0]
smoke-decode16-cold	0	5	false	[1, 2, 4, 8, 16]	12	true	true	[2, 1, 0]
smoke-decode16-cold	1	5	false	[1, 2, 4, 8, 16]	12	true	true	[2, 1, 0]
smoke-decode16-cold	2	5	false	[1, 2, 4, 8, 16]	12	true	true	[2, 1, 0]
smoke-decode16-cold	3	5	false	[1, 2, 4, 8, 16]	12	true	true	[2, 1, 0]
smoke-decode16-warm	0	0	true	[1, 2, 4, 8, 16]	12	true	true	[2, 1, 0]
smoke-decode16-warm	1	0	true	[1, 2, 4, 8, 16]	12	true	true	[2, 1, 0]
smoke-decode16-warm	2	0	true	[1, 2, 4, 8, 16]	12	true	true	[2, 1, 0]
smoke-decode16-warm	3	0	true	[1, 2, 4, 8, 16]	12	true	true	[2, 1, 0]

Native/unit archive SHA256: 6ff9acf41a44e6b48ef3a15bf118cc7187cb3a346bc31e6c808cde78eb215bc7.
GSM8K archive: b527f4abe8f98504d67fd4f151fb5b748cd93f4616dd1fb6ac09a06211791630.
Reload archive: 314622e034c8109d222196dfeee51fcb704f01cb2699d1e52963060784fb57b0.
All source hashes were unchanged and all owned processes exited.

Scope: no sanitizer or controlled before/after performance rerun in this update.
The older performance measurements below predate native MegaMoE startup tuning.

Historical end-to-end performance, complete raw metrics and reproduction commands

Speed Tests and Profiling

Historical timings: these measurements predate the startup MegaMoE
autotuning added here. They retain their original source pins and external FC2
reduction. They do not establish the current head's end-to-end speedup.

No clear end-to-end speedup over split CuTe DSL with FlashInfer A2A in these runs. Mega averaged 28.095 s / 9634.06 output tok/s, versus 27.985 s / 9671.92 tok/s for that split path (0.996×). Against split CuTe DSL with standard communication (none), Mega showed a modest 1.043× speedup (29.316 s / 9232.68 tok/s for split).

• Environment: c2 hu-pdx-83, same 4× B300 SXM6 AC GPUs (0–3, SM103), sequential servers. Image lmsysorg/sglang:nightly-dev-cu13-20260913-fa663e72, digest sha256:91da70aa3c078b09af1e7501d7fee0827ff2683d0b4707e3c9f7663e2f824f5a. Driver 590.48.01; Python 3.12.3; PyTorch 2.13.0+cu130 / CUDA 13.0; Triton 3.7.1; Transformers 5.12.1; sglang-kernel 0.4.6.post1; CuTe DSL/libraries 4.7.1 (native compiler CUDA 13.3); FlashInfer Python/cubin/JIT-cache 0.6.18 / 0.6.18 / 0.6.18+cu130. Image PyTorch/native SGLang extension retained.
• Pins: effective SGLang source ac2bfe25067a27e59f42a6fbce5145fcfa585455; FlashInfer 083a9fd39e7ddda43a8b912ef50f7a13390287f9; Qwen model/revision shown below. All routed experts use online NVFP4/BF16; dense layers retain FP8. No test-only quantization overrides or expert exclusions; the existing online loader applies its standard weight-quantization recipe.
• Shared settings: TP=DP=EP=4, DP attention, no DP LM head; attention trtllm_mha; dense FP8 backend auto; KV auto → BF16; memory fraction 0.87; chunked prefill 4096/rank; max_prefill_tokens=16384; full decode graphs through batch 512. Prefill graphs explicitly disabled in all arms to match Mega's supported configuration.
• Transport/finalize: split FlashInfer uses BF16 payload all-gather/reduce-scatter for prefill and MNNVL A2A for decode (capacity 4096/rank). Split none retains EP4, gathering before routing and reduce-scattering afterward. Split retains default fused finalize (True); Mega uses external top-k reduction. Reduction grouping differs; matched precision does not imply bitwise parity.
• Workload: 256 random requests, seed 42; default range ratio 0.0 independently samples input/output lengths from 1–8192 / 1–2048, using repeated/truncated ShareGPT text (anon8231489123/ShareGPT_Vicuna_unfiltered, snapshot 192ab2185289094fc556ec8ce5ce1e8e587154ca). These are upper bounds, not fixed lengths. Defaults: unlimited arrival/concurrency, streaming, temperature 0, ignore EOS. 1,089,460 benchmark-reported input / 270,665 output tokens per measured run. Input totals are sampled lengths; decoded text can retokenize differently at the server. Observed request lengths are 35–8155 input and 2–2043 output tokens.
• Method: fresh server and full 256-request warmup before each measured run; each invocation also has its default short warmup. --flush-cache clears warmup prefixes before timing; within-run prefix caching remains enabled. Two samples each, order Mega → split FlashInfer → split standard → split standard → split FlashInfer → Mega (ABCCBA).
• Metrics/scope: speedup = mean split duration / mean Mega duration; pooled output throughput = total output / total duration. Client end-to-end timing in this shared prefill-graph-disabled configuration does not measure each backend's best tuned configuration, isolated MoE/decode speed, or moving FlashInfer head. Two repeats do not establish statistical equivalence or accuracy parity.

Reproduction commands

Use Bash in the pinned image with sources installed/checkpoint cached. Run setup in both shells; choose one server below, then stop it fully before the next sample.

export BENCH_ROOT=/hai-workspace/sglang-flashinfer-w4a16-megamoe
cd "$BENCH_ROOT/sglang"
for benchmark_env_name in ${!SGLANG_@} ${!FLASHINFER_@}; do
  unset "$benchmark_env_name"
done
export PYTHONPATH="$BENCH_ROOT/sglang/python:$BENCH_ROOT/flashinfer"
export HF_HOME="$BENCH_ROOT/hf-home"
export CUDA_VISIBLE_DEVICES=0,1,2,3 OMP_NUM_THREADS=4 MAX_JOBS=4 PYTHONUNBUFFERED=1
export SGLANG_FLASHINFER_CUTEDSL_NVFP4_W4A16=1
export SGLANG_FLASHINFER_NVFP4_PER_TOKEN_ACTIVATION=0 SGLANG_FP4_IGNORED_LAYERS=
MODEL=Qwen/Qwen3-30B-A3B-Instruct-2507-FP8
REVISION=5a5a776300a41aaa681dd7ff0106608ef2bc90db
SERVER_ARGS=(--model-path "$MODEL" --revision "$REVISION"
  --quantization nvfp4_online --dtype bfloat16
  --tp-size 4 --dp-size 4 --ep-size 4 --enable-dp-attention
  --cuda-graph-backend-prefill disabled --host 127.0.0.1 --port 31000)

Mega server:

python -m sglang.launch_server "${SERVER_ARGS[@]}" \
  --moe-runner-backend flashinfer_megamoe --moe-a2a-backend flashinfer_megamoe

Split FlashInfer server:

SGLANG_FLASHINFER_NUM_MAX_DISPATCH_TOKENS_PER_RANK=4096 \
python -m sglang.launch_server "${SERVER_ARGS[@]}" \
  --moe-runner-backend flashinfer_cutedsl --moe-a2a-backend flashinfer \
  --flashinfer-a2a-dispatch-type bf16

Split standard server:

python -m sglang.launch_server "${SERVER_ARGS[@]}" \
  --moe-runner-backend flashinfer_cutedsl --moe-a2a-backend none

After server readiness, set RUN_TAG for this sample and run both client commands:

RUN_TAG=mega-1
RUN_DIR="$BENCH_ROOT/bench/results/$RUN_TAG"
mkdir -p "$BENCH_ROOT/bench/results"
mkdir "$RUN_DIR"
TOKENIZER="$HF_HOME/hub/models--Qwen--Qwen3-30B-A3B-Instruct-2507-FP8/snapshots/$REVISION"
BENCH_ARGS=(--backend sglang --host 127.0.0.1 --port 31000
  --model "$MODEL" --tokenizer "$TOKENIZER" --dataset-name random
  --num-prompts 256 --random-input-len 8192 --random-output-len 2048
  --flush-cache --output-details)
python -m sglang.bench_serving "${BENCH_ARGS[@]}" \
  --output-file "$RUN_DIR/warmup-full.jsonl" > "$RUN_DIR/warmup-full.log" 2>&1
python -m sglang.bench_serving "${BENCH_ARGS[@]}" \
  --output-file "$RUN_DIR/measured.jsonl" > "$RUN_DIR/measured.log" 2>&1
curl --fail --silent http://127.0.0.1:31000/server_info > "$RUN_DIR/server-info.json"

All six measured samples completed 256/256 requests with zero errors and identical per-request input/output lengths. Warmup results are excluded below.

Sample, in execution order	Duration (s)	Output tok/s	Mean TTFT (ms)	Mean TPOT (ms)	Mean ITL (ms)	Mean E2E (ms)
mega-1	28.074	9641.26	4939.48	64.21	14.77	20538.87
split-flashinfer-1	27.853	9717.63	4906.11	70.51	14.97	20716.96
split-none-1	29.077	9308.56	5113.08	67.91	14.98	20937.54
split-none-2	29.555	9158.02	5213.15	69.43	15.40	21483.48
split-flashinfer-2	28.116	9626.64	5154.29	68.03	15.05	21050.65
mega-2	28.116	9626.88	4645.33	67.32	14.83	20314.47

Implementation	Mean duration (s)	Duration min–max (s), n=2	Pooled output tok/s	Mega speedup vs this implementation
Mega	28.095	28.074–28.116	9634.06	1.000×
Split FlashInfer	27.985	27.853–28.116	9671.92	0.996×
Split standard	29.316	29.077–29.555	9232.68	1.043×

Mega's pooled output throughput was 0.39% lower than split FlashInfer and 4.35% higher than split standard. Mega's 28.074–28.116 s duration range overlaps split FlashInfer's 27.853–28.116 s range; the two implementations performed similarly in this workload. The mean TTFT/ITL values were 4792.40/14.80 ms for Mega, 5030.20/15.01 ms for split FlashInfer, and 5163.11/15.19 ms for split standard. TTFT includes queueing and prefill; these client metrics do not isolate kernel speed.

Complete raw Qwen metrics (all six summaries transposed)

Raw metric	mega-1	split-flashinfer-1	split-none-1	split-none-2	split-flashinfer-2	mega-2
Backend:	sglang	sglang	sglang	sglang	sglang	sglang
Traffic request rate:	inf	inf	inf	inf	inf	inf
Max request concurrency:	not set	not set	not set	not set	not set	not set
Successful requests:	256	256	256	256	256	256
Benchmark duration (s):	28.07	27.85	29.08	29.55	28.12	28.12
Total input tokens:	1089460	1089460	1089460	1089460	1089460	1089460
Total input text tokens:	1089460	1089460	1089460	1089460	1089460	1089460
Total generated tokens:	270665	270665	270665	270665	270665	270665
Total generated tokens (retokenized):	270650	270652	270641	270641	270653	270650
Request throughput (req/s):	9.12	9.19	8.80	8.66	9.11	9.11
Input token throughput (tok/s):	38807.27	39114.66	37468.10	36862.16	38748.40	38749.37
Output token throughput (tok/s):	9641.26	9717.63	9308.56	9158.02	9626.64	9626.88
Peak output token throughput (tok/s):	20324.00	23047.00	21891.00	19908.00	21478.00	21801.00
Peak concurrent requests:	256	256	256	256	256	256
Total token throughput (tok/s):	48448.53	48832.29	46776.66	46020.18	48375.04	48376.25
Concurrency:	187.29	190.41	184.34	186.09	191.67	184.97
Mean E2E Latency (ms):	20538.87	20716.96	20937.54	21483.48	21050.65	20314.47
Median E2E Latency (ms):	22163.70	22495.33	22392.37	23068.75	22670.92	21867.31
P90 E2E Latency (ms):	27190.86	26830.78	27933.74	28387.56	27692.97	26938.45
P95 E2E Latency (ms):	27503.32	27319.41	28512.82	28974.86	27922.83	27545.22
P99 E2E Latency (ms):	27848.81	27620.13	28851.15	29330.38	28007.86	27900.61
Mean TTFT (ms):	4939.48	4906.11	5113.08	5213.15	5154.29	4645.33
Median TTFT (ms):	4442.61	4640.23	4568.56	4668.29	4643.73	4389.12
P90 TTFT (ms):	7939.44	8418.78	8392.02	8560.18	8418.45	7975.87
P95 TTFT (ms):	8478.84	8905.14	8890.07	9056.53	8906.33	8439.65
P99 TTFT (ms):	9022.50	9527.00	9498.26	9666.67	9516.29	9016.20
Mean TPOT (ms):	64.21	70.51	67.91	69.43	68.03	67.32
Median TPOT (ms):	15.11	15.23	15.11	15.65	15.34	14.97
P90 TPOT (ms):	30.82	34.35	33.01	34.00	33.26	31.83
P95 TPOT (ms):	49.21	51.13	51.29	52.47	51.70	48.97
P99 TPOT (ms):	344.27	366.99	368.86	374.38	367.03	346.92
Mean ITL (ms):	14.77	14.97	14.98	15.40	15.05	14.83
Median ITL (ms):	11.14	10.30	9.92	11.88	9.46	9.64
P90 ITL (ms):	25.30	25.77	27.87	23.97	28.22	27.36
P95 ITL (ms):	29.75	31.97	32.47	26.89	32.59	32.85
P99 ITL (ms):	36.27	87.80	38.03	33.54	38.28	37.92
Max ITL (ms):	6423.87	8255.06	7038.53	7198.03	7017.00	7719.65

GLM-5.2 NVFP4 W4A16 with bundled BF16 MTP

One retained Mega sample versus the requested image-version FlashInfer Split control, on the same 4×B300: MegaMoE delivered 1,887.80 output tok/s versus 1,356.73 for Split CuTe DSL / standard communication (none): +39.14% throughput / -28.13% duration, or 1.3914× Split/Mega duration. This is a mixed-FlashInfer comparison, not an isolated kernel speedup. The new release-FI Split is primary regardless of direction; all earlier Split samples remain retained, without pooling or fastest-sample selection.

The Split regression remains: release Split is only +0.3419% versus the preceding ad0 Split (1,352.107685 tok/s), -0.7669% versus the earlier ad0 Split (1,367.215440), and -19.1324% versus the old PR's two-Split pool (1,677.718482). Restoring the image's three FI package versions did not recover that historical throughput or remove the mixed MTP acceptance pattern below; this does not establish the cause or a numerical correctness defect.

• Source/image: shared SGLang 6d8a58f177a1488e10e8567a3fa4b4fccd0a26e2. Mega uses FI #5019 ad0a5e5e78e57070ec7c582efe733cb55cd8839f, version 0.7.0; Split imports flashinfer-python==0.6.18, flashinfer-cubin==0.6.18, flashinfer-jit-cache==0.6.18+cu130 from an isolated overlay, with no ad0 source override. Image lmsysorg/sglang:nightly-dev-cu13-20260918-20518d85, amd64 sha256:b518f4f8cd15664cf0f733e9bf4fd2105c9c994882a369b247364394db99f596; host hu-pdx-146, four of eight B300 GPUs used.
• Shared runtime/model: Python 3.12.3, Torch 2.13.0+cu130 / CUDA 13.0, all five CuTe packages 4.7.1, driver 590.48.01; Torch NCCL 2.29.7 versus installed NCCL package 2.30.7. GLM-5.2 target routed experts use NVFP4 weights/BF16 activations; bundled MTP routed weights are BF16. All 282 checkpoint headers checked; config SHA256 e1d5db5fbbb5563af0120260b72f845885a8567fabd0c2a8069852d72d12c334, index ab782ea1432ef96fa4f8931b5732f835e7a74f08e499800ae99c77aa00a13f89; no full-weight payload rehash.
• Shared configuration: TP=DP=EP=4, DP attention and DP/FP32 LM head, memory fraction 0.85, server max-running-requests 64 (16/DP), chunked prefill 4096 / max prefill 16384. FlashMLA sparse DSA prefill/KV decode; target FlashInfer top-k (FLASHINFER_TOPK_ALGO=default), draft sgl-kernel; FP8 E4M3 KV / page 64, target/draft scales 1.0; prefill graphs off, decode graphs on, stream interval 32. EAGLE steps/top-k/draft-tokens = 3/1/4, explicitly unquantized bundled draft with Triton/none. Split uses nonfused finalization, Mega external FC2 reduction. No FlashInfer-A2A Split arm.
• Equal work: 256 full-warmup requests then 256 measured requests per arm, plus one short warmup per invocation: 512 measured + 512 full warmup + 4 short warmups in the reported pair. Random input/output caps 8192/2048, range ratio 0, seed 42, unlimited client rate/concurrency; ShareGPT SHA256 35f0e213ce091ed9b9af2a1f0755e9d39f9ccec34ab281cd4ca60d70f6479ba4. All saved full-phase ordered input/output length arrays match across the retained Mega and three Split samples: input 35–8155, output 2–2043; each measured sample has 1,089,460 input / 270,665 output tokens and 256 successes, zero errors.

One measured sample per arm	Duration (s)	Output tok/s	Output tok/s/GPU	Mean TTFT (ms)	Mean TPOT (ms)	Mean E2E (ms)	Client MTP accept length*
Retained MegaMoE / MegaMoE — FI ad0	143.375716	1887.802257	471.950564	54288.112853	35.156505	85404.133266	3.924492
Primary Split CuTe DSL / none — FI 0.6.18	199.498026	1356.730217	339.182554	66708.272221	39.261592	103397.862711	2.234555

*The client scalar is internal_states[0] in its own /server_info response, whose states arrive in response order. It is not a fixed DP rank or a global weighted mean. Separately saved after-measurement response-order values are Mega [3.924492, 3.931288, 3.935730, 3.991198] and release Split [2.234555, 3.787245, 2.314098, 2.302491]. These are four response values, not rank labels; the three-low/one-high Split pattern persists. Semantic output accuracy is not established by this throughput run.

• Historical Mega comparison: retained Mega 1,887.802257 versus the old PR's two-Mega pooled 1,726.157854 tok/s (+9.3644%); duration 143.375716 s versus old mean 156.801998 s (−8.5626%). Mean TTFT/TPOT/E2E changed −8.2205%/−8.9086%/−8.3397%. Ordered lengths and common settings match, but the old runs used SGLang b594bc02c03011288dee9c1ed99f256ea3e4f763, FI 083a9fd39e7ddda43a8b912ef50f7a13390287f9, image nightly-dev-cu13-20260913-fa663e72 and another physical node/cache history. One new versus two old samples does not isolate autotuning or establish significance; the larger refreshed pair gap accompanies both the faster Mega and slower Split samples.
• Startup/validation: both arms captured batches [1,2,3,4,5,6,7,8,10,12,14,16], target/draft-decode/draft-extend widths 4/1/4. Mega recorded capacities [1,2,4,8,16,32,64,4096]. Release Split started with an empty run cache and saved nine W4A16 profiles [1,2,4,8,16,32,64,128,256] in each of two four-rank namespaces; draft-namespace copies do not prove draft W4A16 execution. The three FI providers changed through pip --target; the base package freeze remained unchanged. All 537 transferred evidence files (including prior runs and preparation), exits, import origins, settings and ordered lengths were rehashed/checked. Earlier orchestration failures and successful ad0 Split samples are retained. No dedicated large-prefill tuning or best-tuned-backend claim.
• Metric limits: startup/tuning/full warmup are excluded from measured duration; queueing and the final acceptance fetch are included. Saved arrays independently reproduce token throughput, TTFT and ITL; final per-request latency arrays are absent, so TPOT/E2E remain client-reported. TTFT starts at the first nonempty text chunk and ITL excludes that chunk, which may contain multiple tokens. Raw SSE/semantic correctness and sustained under-load power/clock traces were not validated. Different FI providers, startup/cache histories and one sample per arm prevent causal or significance claims.

Full raw arrays, logs, source/bootstrap/preflight/import receipts, manifests and prior-run lineage are retained locally and in persistent campaign storage; no public raw-payload link is claimed.

Complete raw measured client summaries

Retained MegaMoE / MegaMoE — FI ad0

============ Serving Benchmark Result ============
Backend:                                 sglang    
Traffic request rate:                    inf       
Max request concurrency:                 not set   
Successful requests:                     256       
Benchmark duration (s):                  143.38    
Total input tokens:                      1089460   
Total input text tokens:                 1089460   
Total generated tokens:                  270665    
Total generated tokens (retokenized):    270644    
Request throughput (req/s):              1.79      
Input token throughput (tok/s):          7598.64   
Output token throughput (tok/s):         1887.80   
Peak output token throughput (tok/s):    3217.00   
Peak concurrent requests:                256       
Total token throughput (tok/s):          9486.44   
Concurrency:                             152.49    
Accept length:                           3.92      
----------------End-to-End Latency----------------
Mean E2E Latency (ms):                   85404.13  
Median E2E Latency (ms):                 89598.05  
P90 E2E Latency (ms):                    134170.21 
P95 E2E Latency (ms):                    136958.95 
P99 E2E Latency (ms):                    141396.36 
---------------Time to First Token----------------
Mean TTFT (ms):                          54288.11  
Median TTFT (ms):                        51814.53  
P90 TTFT (ms):                           112180.58 
P95 TTFT (ms):                           119130.08 
P99 TTFT (ms):                           127158.73 
-----Time per Output Token (excl. 1st token)------
Mean TPOT (ms):                          35.16     
Median TPOT (ms):                        30.77     
P90 TPOT (ms):                           39.73     
P95 TPOT (ms):                           41.98     
P99 TPOT (ms):                           57.07     
---------------Inter-Token Latency----------------
Mean ITL (ms):                           29.51     
Median ITL (ms):                         30.44     
P90 ITL (ms):                            39.95     
P95 ITL (ms):                            47.13     
P99 ITL (ms):                            65.08     
Max ITL (ms):                            967.81    
==================================================

Primary Split CuTe DSL / none — FI 0.6.18

============ Serving Benchmark Result ============
Backend:                                 sglang    
Traffic request rate:                    inf       
Max request concurrency:                 not set   
Successful requests:                     256       
Benchmark duration (s):                  199.50    
Total input tokens:                      1089460   
Total input text tokens:                 1089460   
Total generated tokens:                  270665    
Total generated tokens (retokenized):    235783    
Request throughput (req/s):              1.28      
Input token throughput (tok/s):          5461.01   
Output token throughput (tok/s):         1356.73   
Peak output token throughput (tok/s):    2564.00   
Peak concurrent requests:                256       
Total token throughput (tok/s):          6817.74   
Concurrency:                             132.68    
Accept length:                           2.23      
----------------End-to-End Latency----------------
Mean E2E Latency (ms):                   103397.86 
Median E2E Latency (ms):                 103567.33 
P90 E2E Latency (ms):                    168352.49 
P95 E2E Latency (ms):                    173919.29 
P99 E2E Latency (ms):                    185158.92 
---------------Time to First Token----------------
Mean TTFT (ms):                          66708.27  
Median TTFT (ms):                        62399.30  
P90 TTFT (ms):                           141357.39 
P95 TTFT (ms):                           152763.53 
P99 TTFT (ms):                           163506.77 
-----Time per Output Token (excl. 1st token)------
Mean TPOT (ms):                          39.26     
Median TPOT (ms):                        26.63     
P90 TPOT (ms):                           78.32     
P95 TPOT (ms):                           85.55     
P99 TPOT (ms):                           102.95    
---------------Inter-Token Latency----------------
Mean ITL (ms):                           34.73     
Median ITL (ms):                         25.07     
P90 ITL (ms):                            71.92     
P95 ITL (ms):                            97.32     
P99 ITL (ms):                            145.84    
Max ITL (ms):                            1091.08   
==================================================

Reproduction: pinned setup, provider selection and identical workload

Inside the image above, retain its Torch installation. Use the existing checkpoint and ShareGPT file whose hashes are given above. The base setup used the following package versions and independently pinned source trees (the actual install used predownloaded wheels with --no-index):

export ROOT=/data/home/ziangli/sglang-glm52-bundled-mtp-refresh-20260921
export MODEL=/data/models/ziangli_glm52/GLM-5.2-NVFP4-W4A16
git clone https://github.com/sgl-project/sglang.git "$ROOT/sglang"
git -C "$ROOT/sglang" checkout --detach 6d8a58f177a1488e10e8567a3fa4b4fccd0a26e2
git clone https://github.com/flashinfer-ai/flashinfer.git "$ROOT/flashinfer"
git -C "$ROOT/flashinfer" fetch origin ad0a5e5e78e57070ec7c582efe733cb55cd8839f
git -C "$ROOT/flashinfer" checkout --detach ad0a5e5e78e57070ec7c582efe733cb55cd8839f
git -C "$ROOT/flashinfer" submodule update --init --recursive --depth=1 3rdparty/cutlass 3rdparty/cccl 3rdparty/spdlog
/opt/sglang/bin/python3 -I -m pip --isolated uninstall -y flashinfer-python flashinfer-cubin flashinfer-jit-cache
/opt/sglang/bin/python3 -I -m pip --isolated install --no-deps \
  nvidia-cutlass-dsl==4.7.1 nvidia-cutlass-dsl-libs-base==4.7.1 \
  nvidia-cutlass-dsl-libs-core==4.7.1 nvidia-cutlass-dsl-libs-cu12==4.7.1 \
  nvidia-cutlass-dsl-libs-cu13==4.7.1 nccl-extensions==0.1.0 \
  cupti-python==13.2.0 nvidia-cuda-cupti==13.2.86
/opt/sglang/bin/python3 -I -m pip --isolated install --no-deps --no-build-isolation --no-index -e "$ROOT/flashinfer"

For the release Split control only, install the image's three FI versions into an independent overlay, leaving the base freeze unchanged. Verify the exact wheel hashes before installation:

export OVERLAY="$ROOT/controls/stock-flashinfer-r5/overlay"
export WHEELS="$ROOT/controls/stock-flashinfer-r5/wheels"
mkdir -p "$WHEELS"
curl -fL "https://files.pythonhosted.org/packages/31/7d/3dae45fc293cd5f16cefe4dcf59683d964ca8faf2303ed7fff54f4bca490/flashinfer_python-0.6.18-py3-none-any.whl" -o "$WHEELS/flashinfer_python-0.6.18-py3-none-any.whl"
curl -fL "https://github.com/flashinfer-ai/flashinfer/releases/download/v0.6.18/flashinfer_cubin-0.6.18-py3-none-any.whl" -o "$WHEELS/flashinfer_cubin-0.6.18-py3-none-any.whl"
curl -fL "https://github.com/flashinfer-ai/flashinfer/releases/download/v0.6.18/flashinfer_jit_cache-0.6.18+cu130-cp39-abi3-manylinux_2_28_x86_64.whl" -o "$WHEELS/flashinfer_jit_cache-0.6.18+cu130-cp39-abi3-manylinux_2_28_x86_64.whl"
(cd "$WHEELS" && sha256sum -c - <<'SHA256'
d5d26edb48f8def0bf28b2c94415aef7f0754e2cd9e3cb3d6e7081262d068aad  flashinfer_python-0.6.18-py3-none-any.whl
2dd65c0fcfc6bc44c67f148530de5372979c2e3d260e47935730f94156d4d873  flashinfer_cubin-0.6.18-py3-none-any.whl
428a47a554ade93c30a818e142b781df58582bf056bab94611fb4c906cc366bf  flashinfer_jit_cache-0.6.18+cu130-cp39-abi3-manylinux_2_28_x86_64.whl
SHA256
)
/opt/sglang/bin/python3 -I -m pip --isolated install --no-deps --no-index --no-cache-dir --no-compile --target "$OVERLAY" \
  "$WHEELS/flashinfer_python-0.6.18-py3-none-any.whl" \
  "$WHEELS/flashinfer_cubin-0.6.18-py3-none-any.whl" \
  "$WHEELS/flashinfer_jit_cache-0.6.18+cu130-cp39-abi3-manylinux_2_28_x86_64.whl"

Select one arm at a time; use a separate cache directory for each. The reported release control used a new empty cache, without copied ad0 tactics. The original Mega sample is retained, not rerun. Shared exports and exact server/client arguments follow.

export ROOT=/data/home/ziangli/sglang-glm52-bundled-mtp-refresh-20260921
export MODEL=/data/models/ziangli_glm52/GLM-5.2-NVFP4-W4A16
ARM=release-split # or mega
if [ "$ARM" = mega ]; then
  export RUN_ID=glm52-bundled-mtp-20260921-r1
  export PYTHONPATH="$ROOT/sglang/python:$ROOT/flashinfer"
  RUNNER=flashinfer_megamoe A2A=flashinfer_megamoe
else
  export RUN_ID=glm52-bundled-mtp-20260921-r5-stock-fi
  export PYTHONPATH="$ROOT/sglang/python:$ROOT/controls/stock-flashinfer-r5/overlay"
  export PYTHONDONTWRITEBYTECODE=1
  unset PYTHONHOME
  RUNNER=flashinfer_cutedsl A2A=none
fi
export CACHE="$ROOT/runs/$RUN_ID/compiled-cache"
export HF_HOME="$ROOT/hf-home"
export CUDA_VISIBLE_DEVICES="0,1,2,3"
export OMP_NUM_THREADS="4"
export MAX_JOBS="4"
export PYTHONUNBUFFERED="1"
export SGLANG_CACHE_DIR="$CACHE/sglang"
export FLASHINFER_WORKSPACE_BASE="$CACHE/sglang"
export TRITON_CACHE_DIR="$CACHE/triton"
export TORCHINDUCTOR_CACHE_DIR="$CACHE/inductor"
export TILELANG_CACHE_DIR="$CACHE/tilelang"
export CUDA_CACHE_PATH="$CACHE/cuda"
export CUTE_DSL_CACHE_DIR="$CACHE/cute"
export INDEXER_ROPE_NEOX_STYLE="0"
export SGLANG_DSA_FUSE_TOPK="1"
export SGLANG_DSA_PREFILL_DENSE_ATTN_KV_LEN_THRESHOLD="0"
export SGLANG_DSA_TOPK_FLASHINFER_TIE_BREAK="large"
export NVSHMEM_DISABLE_NCCL="1"
export FLASHINFER_NVFP4_4OVER6="1"
export FLASHINFER_NVFP4_4OVER6_E4M3_USE_256="0"
export FLASHINFER_NVFP4_4OVER6_ERR_MODE="MSE"
export FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH="1"
export TRTLLM_DISABLE_FP4_QUANT_FAST_MATH="1"
export FLASHINFER_DISABLE_FP4_QUANT_FAST_MATH="1"
export SGLANG_FLASHINFER_CUTEDSL_NVFP4_W4A16="1"
export SGLANG_FLASHINFER_NVFP4_PER_TOKEN_ACTIVATION="0"
export SGLANG_MOE_NVFP4_DISPATCH="0"
export SGLANG_FLASHINFER_MOE_FUSED_FINALIZE="0"
export FLASHINFER_TOPK_ALGO="default"
export SGLANG_FLASHINFER_MEGAMOE_IN_KERNEL_FC2_REDUCE="0"

/opt/sglang/bin/python3 -m sglang.launch_server \
  --model-path "$MODEL" --tp-size 4 \
  --dp-size 4 --ep-size 4 \
  --enable-dp-attention --enable-dp-lm-head \
  --enable-fp32-lm-head --mem-fraction-static 0.85 \
  --max-running-requests 64 --attention-backend dsa \
  --dsa-topk-backend flashinfer --dsa-decode-backend flashmla_kv \
  --dsa-prefill-backend flashmla_sparse --kv-cache-dtype fp8_e4m3 \
  --page-size 64 --disable-shared-experts-fusion \
  --moe-dense-tp-size 1 --cuda-graph-backend-prefill disabled \
  --reasoning-parser glm45 --speculative-algorithm EAGLE \
  --speculative-num-steps 3 --speculative-eagle-topk 1 \
  --speculative-num-draft-tokens 4 --speculative-moe-runner-backend triton \
  --speculative-moe-a2a-backend none --speculative-draft-model-quantization unquant \
  --enable-metrics --stream-interval 32 \
  --moe-runner-backend "$RUNNER" --moe-a2a-backend "$A2A" \
  --host 127.0.0.1 --port 31000

These exact arguments resolve chunked prefill4096/max prefill16384 in the pinned source. For each healthy server, run the following once with PHASE=warmup-full, then once with PHASE=measured; the client’s extra short warmup remains enabled. Do not pool full warmup with measured results.

PHASE=measured # warmup-full for the first invocation
/opt/sglang/bin/python3 -m sglang.bench_serving --backend sglang --host 127.0.0.1 --port 31000 \
  --model "$MODEL" --tokenizer "$MODEL" \
  --dataset-name random --dataset-path "$ROOT/datasets/ShareGPT_V3_unfiltered_cleaned_split.json" \
  --num-prompts 256 --random-input-len 8192 --random-output-len 2048 \
  --random-range-ratio 0 --seed 42 --request-rate inf \
  --warmup-requests 1 --flush-cache --output-details --output-file "$PHASE.jsonl"

Metric code: final acceptance fetch and benchmark duration; response arrival-order collection. This reported pair is one retained ad0 Mega and one image-version FI Split sample. The two earlier ad0 Split samples remain separate diagnostic evidence; this is not ABBA or a fastest-result pool.

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.

Review and Merge Process

1. Ping Merge Oncalls to start the process. See the PR Merge Process.
2. Get approvals from CODEOWNERS and other reviewers.
3. Trigger CI tests with comments or contact authorized users to do so.
   • Common commands include /tag-and-rerun-ci, /tag-run-ci-label, /rerun-failed-ci
4. After green CI and required approvals, ask Merge Oncalls or people with Write permission to merge the PR.

---

CI States

Latest PR Test (Base): ❌ Run #36545802564
Latest PR Test (Extra): ❌ Run #36545801990
Latest PR Test (AMD ROCm 10): ❌ Run #36545802424