feat(pytorch): add glm5.3-flash support

[qescccczmr]

Summary

Add GLM-5.3 Flash text/image/video support to the PyTorch engine, including MTP and combined data/expert parallelism. Extend existing LMDeploy components for model loading, hybrid cache layout, multimodal preprocessing, and CUDA Graph execution.

• Reuse DeepSeek MLA/MoE loaders, standard TP projections, FA3, HcPrePost, sparse-index Top-K, and blocked-FP8 MoE. GLM sparse attention uses the common FlashMLA backend.
• Keep KDA as a thin backend adapter around public FLA operations and the shared recurrent/causal-convolution kernels. Preserve channelwise gates and per-token speculative state checkpoints.
• Reuse common vision LayerNorm/RoPE operations with FP32 computation, and preserve clamped SwiGLU and routed-expert scaling.
• Reuse the MTP predictor, rejection sampler, cache planner, and target-to-draft multimodal embedding handoff.

Optimizations

• KPool: fuse multi-token verification and persistent-ring checkpoints; skip compression for unclosed pools; bound prefill score workspace. Broadcast selected groups before local expansion, share decode metadata across target layers, and fuse draft raw-cache gather/write. Preserve padded score strides in Top-K instead of copying the full score capacity.
• Sparse MLA: reuse FlashMLA with the compact 512-value NoPE cache; fuse index remap/padding and reuse the shared Hadamard query rotation.
• KDA: consume strided convolution/recurrent inputs and improve small-batch scheduling. Merge f_a/g_a and batch f_b/g_b, reducing projection launches 6→4 while retaining checkpoint/TP loading and LoRA fallback.
• mHC: fuse pre-reduction/RMSNorm and post/FP32-input preparation, retaining the original BF16 rounding. Compile large-prefill input conversion while preserving FP32 GEMM and RMS statistics.
• DeepEP decode: fuse clamped activation and FP8 quantization for valid expert rows only, preserving the BF16 rounding before quantization and FP32 routed scaling.
• MTP/communication: reuse key/score projections; preserve FP32 MoE/TP reductions. NCCL is the default; optional FlashInfer all-reduce supports FP32 and retains fallback for unsupported sizes.

Measured optimization gains

H200, TP4/EP1/DP1, BS1, MTP5, BF16 KV, output 5 tokens, unless specified. Serving timing excludes active profiling. Speedup = before/after; reduction = (before−after)/before. Each stage has its own control; gains are not additive or multiplicative.

End-to-end staged measurements

Change / workload	Before	After	Speedup	Latency reduction
KPool verification fusion, 2K total request, 10 samples	315.26 ms	254.91 ms	1.24×	19.15%
KPool verification fusion, 2K TPOT	17.60 ms	9.09 ms	1.94×	48.34%
Index/mHC/KDA/shared-communicator stage, 2K total, 30 samples	245.21 ms	234.28 ms	1.05×	4.46%
Query-rotation/post-expand prefill stage, 8K TTFT, 60 samples per variant	477.13 ms	441.15 ms	1.08×	7.54%
New: sparse MLA remap + padding, 8K TTFT, 30 paired samples	440.30 ms	438.90 ms	1.0032×	0.32%
New: sparse MLA remap + padding, 8K total request, 30 paired samples	470.82 ms	469.91 ms	1.0019×	0.19%
New (2026-09-29): mHC input conversion, 8K TTFT, 30 paired samples	446.23 ms	433.35 ms	1.030×	2.89%
New (2026-09-29): mHC input conversion, 8K total request, 30 paired samples	475.91 ms	462.68 ms	1.029×	2.78%
KDA batched gates + shared KPool metadata + mHC FP32 preparation, 2K TPOT, 30 pairs	8.353 ms	7.539 ms	1.108×	9.75%
KDA batched gates + shared KPool metadata + mHC FP32 preparation, 2K total request, 30 pairs	264.205 ms	261.389 ms	1.011×	1.07%
KDA batched gates + shared KPool metadata + mHC FP32 preparation, 8K TPOT, 30 pairs	7.292 ms	6.857 ms	1.063×	5.96%
KDA batched gates + shared KPool metadata + mHC FP32 preparation, 8K total request, 30 pairs	461.067 ms	459.403 ms	1.004×	0.36%
DP8/EP8: masked expert activation/quantization + strided KPool scores, K=1, mean TPOT	18.784 ms	16.512 ms	1.138×	12.10%
DP8/EP8: masked expert activation/quantization + strided KPool scores, K=3, mean TPOT	16.045 ms	12.942 ms	1.240×	19.34%
DP8/EP8: masked expert activation/quantization + strided KPool scores, K=5, mean TPOT	14.479 ms	11.832 ms	1.224×	18.28%
Optional prefill CUDA graphs, DP8/EP8 K=1, native mean TPOT	16.327 ms	15.610 ms	1.046×	4.39%
Optional prefill CUDA graphs, DP8/EP8 K=3, native mean TPOT	12.931 ms	12.149 ms	1.064×	6.05%
Optional prefill CUDA graphs, DP8/EP8 K=5, native mean TPOT	11.684 ms	10.939 ms	1.068×	6.38%

• Historical rows retain their original controls and rounded values. Query-rotation/post-expand did not improve 2K TTFT (218.23→220.52 ms). Sparse MLA's 8K TTFT gain has a negative paired 95% interval; its 0.19% total-latency change is not statistically established.
• mHC input conversion (19307597): 3 warmup +30 measured pairs, 8K TTFT/total improve in 30/30 pairs; paired 95% intervals are [−14.90,−11.38]/[−14.36,−11.65] ms. No established 2K or TPOT gain.
• Combined decode optimization: direct A/B of shared KPool metadata, batched gates and mHC FP32 preparation against 932285b0 plus the already validated f_a/g_a merge and MTP projection reuse. Three warmup +30 measured pairs, alternating order, separate captured graphs. TPOT paired 95% intervals are [−0.822,−0.246] ms (2K) and [−0.527,−0.373] ms (8K); both total-latency intervals also exclude zero. TTFT intervals cross zero.

The earlier TP stack at ec2ec448 matches the preceding LMDeploy path on generated token IDs and 96 full-vocabulary logits tensor pairs (24,161,280 elements; max error 0). TP/LoRA loading, ragged cache/rejection, changed-input graph replay, and existing HC tests pass. This is implementation parity, not a new full task-accuracy run.

Cumulative Q/K-length metadata fusion passes exact full-model logits (102 tensors / 25,090,560 values) and generated-token checks, but paired TPOT 95% intervals cross zero: [−0.116,+0.103] ms / [−0.130,+0.037] ms. It is excluded from the final stack; no new production files or source changes are retained.

DP/EP staged rows use input 2048, output 1024, global concurrency 32, 8 H200s, TP1/DP8/EP8, BF16 KV and FP32 recurrent state: ec2ec448 → ed1eea00, 32 warmup + 256 measured requests per K and revision. These are sequential fresh-service measurements, not paired confidence intervals. The dominant gain is masked activation/quantization; the additional KPool-copy removal changed mean TPOT by only 0.42% / 1.63% / 0.34% for K=1/3/5, so its small serving effect is not statistically established. K=5 throughput fell 0.93% in that incremental run.

Optional prefill CUDA graphs: 209a7552 → b2665817, same 64 warmup + 256 measured requests per K. These three rows are fresh-service comparisons. A separate controlled K=5 experiment (six alternating pairs, 64 requests/arm/pair; 384/arm) measured 9.495 → 9.028 ms mean TPOT; mean paired reduction 4.92%, paired bootstrap 95% interval [3.29%, 6.56%], all six pairs improved. Burst throughput was variable (pooled −0.35%); a repeatable throughput gain is not established by that paired experiment. The paired experiment uses the same 2048-row path; the final revision additionally keeps partial buckets eager to preserve BF16 projection rounding. Paired burst metrics and the independent final serving table are different experiments.

Operator-level gains

Isolated GPU measurements; these are not whole-model serving speedups. Historical and newer controls are separate experiments.

Optimization / workload	Before	After	Speedup	Latency reduction
KPool sequence update/compression/cache write, B1/S6 (805 → 4 kernels)	1130.10 µs	8.38 µs	134.86×	99.26%
KDA adapter, B1/S6	37.86 µs	14.66 µs	2.58×	61.28%
Query rotation, 8K rows	2055.00 µs	45.00 µs	45.67×	97.81%
HC post-expand, 8K rows	323.00 µs	179.00 µs	1.80×	44.58%
New: prefill group broadcast + local metadata/expansion, TP4/4096 rows	0.3953 ms	0.3428 ms	1.15×	13.29%
New: prefill group broadcast + local metadata/expansion, TP4/8192 rows	0.5865 ms	0.4414 ms	1.33×	24.75%
New: valid-aware KPool update/compress/cache write, B1/S1 (no closed pool)	4.088 µs	3.363 µs	1.22×	17.74%
New: valid-aware KPool update/compress/cache write, B1/S6 (no gain)	5.761 µs	5.819 µs	0.99×	-1.00%
New: valid-aware KPool update/compress/cache write, B64/S6	6.835 µs	6.603 µs	1.04×	3.39%
New: sparse MLA index remap + padding, 8K/BS1, 7 → 1 kernels	119.66 µs	49.82 µs	2.40×	58.36%
New: sparse MLA index remap + padding, 8K/ragged BS2, 7 → 5 kernels	120.07 µs	66.79 µs	1.80×	44.37%
2026-09-29: KDA f_b/g_b batched projection, TP4/B1/S6	5.443 µs	3.151 µs	1.73×	42.10%
2026-09-29: mHC post + FP32 preparation + existing GEMM, 6 rows	16.444 µs	13.342 µs	1.23×	18.86%
Cumulative Q/K-length metadata fusion (not retained), BS1/int64	5.823 µs	1.012 µs	5.75×	82.62%
Native-head sparse MLA without KV splitting (slower; not retained), H16/NoPE512, 6 rows	59.147 µs	93.551 µs	0.63×	−58.17%
Native-head sparse MLA with KV splitting + FP32 merge (numerical checks failed; not retained), H16/NoPE512, 6 rows	59.130 µs	12.090 µs	4.89×	79.55%
DeepEP valid-row clamped activation + FP8 quantization, DP8 K=1 shape	84.087 µs	3.950 µs	21.29×	95.30%
DeepEP valid-row clamped activation + FP8 quantization, DP8 K=3 shape	167.120 µs	5.864 µs	28.50×	96.49%
DeepEP valid-row clamped activation + FP8 quantization, DP8 K=5 shape	254.469 µs	9.691 µs	26.26×	96.19%
KPool score copy + deterministic Top-K → strided Top-K, DP8 K=1 shape	42.309 µs	19.032 µs	2.22×	55.02%
KPool score copy + deterministic Top-K → strided Top-K, DP8 K=3 shape	63.375 µs	19.277 µs	3.29×	69.58%
KPool score copy + deterministic Top-K → strided Top-K, DP8 K=5 shape	84.082 µs	19.299 µs	4.36×	77.05%
Extend mHC post-expand + rounded FP32 preparation, 24 rows, including following FP32 projection	17.908 µs	14.864 µs	1.205×	17.00%
Extend mHC post-expand + rounded FP32 preparation, 32 rows, including following FP32 projection	19.017 µs	15.757 µs	1.207×	17.14%

The valid-aware path still checkpoints every valid token; B1/S6 shows no gain. The new projection/mHC microbenchmarks preserve numerical outputs. Separate full-model Torch Profiler attribution records 34 fewer target-verification kernels for batched gates and 89 fewer for mHC preparation; their timings are not added to serving gains.

The split-K prototype preserves selected KV and FP32 partial reduction, passes operator references/strided-cache/graph tests, but changes unconstrained 2K generation. Fixed-history target/draft logits reach max errors 2.2734 / 1.6932 (2K/8K); target verification alone reaches 1.53125. It is not enabled, and no end-to-end speedup is claimed. This also shows that removing head padding alone is insufficient.

The mHC row-count microbenchmarks preserve the original BF16 rounding before producing the FP32 input for the next HC projection; seven alternating timing pairs per shape, CUDA Graph replay, TF32 disabled. The DeepEP microbenchmarks use 36 local experts, intermediate width 2048, and padded expert rows 512/1024/1536; valid-row counts are fixed and recorded. The KPool microbenchmarks use 8/16/24 rows, 929,984 score columns and 768 valid groups. These isolate wasted work; neither predicts the serving speedup. Removing the score copy does not bound DeepGEMM’s underlying decode-logits allocation.

Capacity improvement: bounded KPool scoring

Workload / measured resource	Before	After	Memory reduction
8K query rows / 128K history, largest FP32 score payload	1 GiB	≤512 MiB	≥50%
8K query rows / 256K history, largest FP32 score payload	2 GiB	≤512 MiB	≥75%

KPool reuses the configurable DSA score budget, flattens KV once, selects deterministic Top-K per query chunk, and reserves score workspace explicitly. The payload limit includes DeepGEMM alignment; it is not a bound on all allocator memory. Full-model 128K/256K requests completed with identical selected indices. This is a capacity fix, not a latency claim.

FlashKDA and the four-way KDA input merge remain excluded after numerical failures. One-warp recurrence and explicit prefill-copy removal showed no benefit; native-head FA3 lacks a validated LMDeploy backend. FP32 reduction semantics are unchanged.

LMDeploy vs vLLM

Tensor parallel: TP4 / EP1 / DP1

Retained TP reference (2026-09-30): LMDeploy ec2ec448; vLLM 606d124b. This table predates the additional DP/EP optimizations below; it is not a fresh TP benchmark of the new head. Cumulative Q/K-length metadata fusion and native-head sparse MLA remain excluded for unconfirmed serving gains and failed numerical checks, respectively. Both services were restarted and run sequentially on the same four H200 GPUs (4–7), checkpoint and prompt manifest, using the unmodified vLLM streaming HTTP client. vLLM follows the official Hopper recipe.

TP4/EP1/DP1, BS1, MTP5 on both engines, BF16 KV, FP32 recurrent-state buffers, 5 output tokens, max batch 16, prefill chunk 2048, session length 270336, prefix cache off, greedy/seed 42. 3 warmups + 30 measured requests per input/engine; independent medians. Profiling was inactive during these measurements.

Input	Metric	LMDeploy	vLLM
2,048	TTFT	234.60 ms	138.45 ms
2,048	TPOT	7.28 ms	7.24 ms
2,048	Total request	264.15 ms	167.49 ms
8,192	TTFT	892.82 ms	557.40 ms
8,192	TPOT	7.08 ms	6.68 ms
8,192	Total request	921.83 ms	582.79 ms

This replaces the previous chunk=8192 table; it is a configuration rerun, not a new optimization gain. The 8K input now requires four prefill chunks in both traces. Historical staged A/B and operator gains above keep their original controls. Neither input establishes a TPOT win over vLLM in this rerun.

TPOT = (last SSE choices response − first SSE choices response)/(5−1). Each median is computed independently. The client includes LMDeploy’s final empty-text finish response; in the diagnostic captures this adds 1.03–1.22 ms to the response window (0.26–0.30 ms to TPOT). MTP can return multiple tokens per SSE event. These are short-output serving metrics, not single-forward GPU latency, steady-state throughput or token-parity results. The same prompt does not always produce identical text across engines; vLLM’s 8K captures also differ in acceptance/outputs.

LMDeploy uses FP32 NCCL/FLA, with FlashInfer/symmetric-memory all-reduce and FlashKDA off; vLLM retains native backends. Memory budgets retain their native meanings: LMDeploy free-memory fraction 0.8; vLLM total-memory utilization 0.9.

Native API-server launch commands

Use the matching installed environments and the FP8 checkpoint at revision 3f1971b7b5f7a528c9c4ef6212c8785298a8c24a. Run the two services sequentially. These commands start the services; BS1, greedy sampling and max_tokens=5 are client request settings.

LMDeploy CLI

MODEL=/path/to/GLM-5.3-Flash
CUDA_VISIBLE_DEVICES=4,5,6,7 NCCL_NVLS_ENABLE=0 FLA_FLASH_KDA=0 \
LMDEPLOY_ENABLE_FLASHINFER_ALLREDUCE=0 \
LMDEPLOY_ENABLE_SYMM_MEM_ALLREDUCE=0 \
lmdeploy serve api_server "$MODEL" \
  --backend pytorch --tp 4 --ep 1 --dp 1 \
  --dtype bfloat16 --quant-policy 0 \
  --max-batch-size 16 --session-len 270336 \
  --max-prefill-token-num 2048 --cache-max-entry-count 0.8 \
  --distributed-executor-backend ray \
  --speculative-algorithm deepseek_mtp --speculative-num-draft-tokens 5 \
  --model-name glm53 --server-name 127.0.0.1 --server-port 23982 \
  --trust-remote-code --log-level INFO

from lmdeploy.messages import PytorchEngineConfig, SpeculativeConfig
from lmdeploy.serve.openai.api_server import serve

serve(
    model_path="/path/to/GLM-5.3-Flash", model_name="glm53", backend="pytorch",
    server_name="127.0.0.1", server_port=23982, trust_remote_code=True, log_level="INFO",
    backend_config=PytorchEngineConfig(
        dtype="bfloat16", tp=4, ep=1, dp=1, max_batch_size=16,
        session_len=270336, max_prefill_token_num=2048, cache_max_entry_count=0.8,
        enable_prefix_caching=False, distributed_executor_backend="ray",
        mp_engine_backend="ray"),
    speculative_config=SpeculativeConfig(method="deepseek_mtp", num_speculative_tokens=5),
)

BF16 KV follows dtype=bfloat16 with KV quantization disabled; GLM’s recurrent-state allocation is FP32. The local measured environment also used a DeepGEMM masked-GEMM compatibility alias and propagated backend/profile environment variables to Ray workers; ensure a compatible DeepGEMM build and consistent worker environments when reproducing.

vLLM CLI

MODEL=/path/to/GLM-5.3-Flash
CUDA_VISIBLE_DEVICES=4,5,6,7 OMP_NUM_THREADS=1 \
VLLM_CUSTOM_SCOPES_FOR_PROFILING=1 \
vllm serve "$MODEL" \
  --host 127.0.0.1 --port 23983 --served-model-name glm53 \
  --tensor-parallel-size 4 --dtype bfloat16 \
  --kv-cache-dtype bfloat16 --mamba-ssm-cache-dtype float32 \
  --max-model-len 270336 --max-num-seqs 16 --max-num-batched-tokens 2048 \
  --gpu-memory-utilization 0.9 --no-enable-prefix-caching --enable-chunked-prefill \
  --speculative-config '{"method":"mtp","num_speculative_tokens":5}' \
  --seed 42 --tool-call-parser glm47 --reasoning-parser glm47 --enable-auto-tool-choice

vLLM uses its default DP1 and no expert-parallel enable flag (EP1). Leave the profiler inactive for timing. CPU+CUDA captures below were collected separately, after all formal requests.

Torch Profiler / Perfetto: TP decode windows

Separate CPU+CUDA diagnostic captures, rank 0 / trial 0 consistently. Actual Perfetto screenshots below. Client timestamps are overlaid after capture: decode_window Duration / 4 = TPOT for 5 output tokens. Native CPU forward scopes and GPU verification spans are different measurements; profiled values do not replace the unprofiled table.

Engine	Input	TTFT	Client post-first-response window	TPOT (window / 4)	Second target verification GPU span
LMDeploy	2,048	320.019 ms	39.725 ms	9.931 ms	13.345 ms
LMDeploy	8,192	1224.736 ms	40.801 ms	10.200 ms	13.595 ms
vLLM	2,048	186.481 ms	38.247 ms	9.562 ms	8.438 ms
vLLM	8,192	727.148 ms	36.988 ms	9.247 ms	8.448 ms

2,048-token input: LMDeploy and vLLM decode windows

LMDeploy

vLLM

8,192-token input: LMDeploy and vLLM decode windows

LMDeploy

vLLM

Four-rank traces, measured client overlays and source hashes are recorded in the screenshots and SQL provenance.

Data and expert parallel: DP8 / EP8 / TP1

2026-09-30, LMDeploy b2665817; vLLM 606d124b. Eight H200 GPUs, identical checkpoint and fixed ShareGPT-derived input IDs: 2048 input / 1024 output tokens / global concurrency 32. K is the number of MTP draft tokens (1/3/5); verification can contain K+1 tokens. BF16 KV, FP32 recurrent state, max batch 32 per rank, prefill chunk 2048 per rank, session length 8192, prefix caching off, greedy/seed42 and ignore-EOS for a fixed output length.

Fresh native-CLI reruns use 64 warmup + 256/256 successful measured requests, four concurrent client workers per DP rank. The unmodified vLLM HTTP streaming client supplies timing; active profiling is excluded. LMDeploy uses native DeepEP Normal for globally mixed/prefill steps and Low Latency for global decode; vLLM explicitly uses deepep_low_latency. This is an engine comparison, not an isolated communication-backend A/B.

Engine	K	Output tok/s	Mean TTFT	Mean TPOT	MTP acceptance	Mean accepted length, including bonus
LMDeploy	1	1971.26	328.82 ms	15.610 ms	94.32%	1.943
LMDeploy	3	2373.09	327.05 ms	12.149 ms	81.36%	3.441
LMDeploy	5	2370.39	343.64 ms	10.939 ms	68.72%	4.436
vLLM	1	2178.85	374.48 ms	14.104 ms	94.44%	1.944
vLLM	3	2546.96	335.09 ms	11.580 ms	82.18%	3.465
vLLM	5	2722.61	343.95 ms	10.348 ms	69.03%	4.451

Mean TPOT remains 10.68% / 4.91% / 5.71% higher than vLLM for K=1/3/5; parity has not been reached. Acceptance = accepted/proposed draft tokens; mean accepted length = 1 + accepted/draft iterations. TPOT uses the first-to-last SSE response interval divided by 1023; ITL is SSE-event spacing, since MTP may emit several tokens per event. The original screenshot’s prompt manifest was unavailable, so this verifies its stated configuration rather than claiming identical acceptance values.

Latest LMDeploy measurements enable all retained optimizations together, including mHC FP32 preparation up to 32 rows and optional prefill CUDA graphs (--piecewise-cudagraph-max-tokens 2048). LMDeploy was freshly restarted for every K; vLLM retains the preceding native-CLI reference with the same 64-warmup protocol, checkpoint, input manifest and workload. Earlier DP/EP staged rows used 32 warmups. The final LMDeploy P99 TPOT values are 19.690 / 17.249 / 20.627 ms for K=1/3/5.

Native API-server commands. Use each engine’s tested environment and installed source revision. Run engines sequentially; set K=1, 3 or 5. The LMDeploy CLI starts eight rank APIs on ports 24100–24107 automatically; the benchmark sends four streams to each rank. vLLM uses one public API and the X-data-parallel-rank header to preserve the same distribution.

MODEL=/path/to/GLM-5.3-Flash
K=3
export CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7
export OMP_NUM_THREADS=1
export LMDEPLOY_DP_MASTER_ADDR=127.0.0.1
export LMDEPLOY_DP_MASTER_PORT=24400
export FLA_FLASH_KDA=0
export LMDEPLOY_ENABLE_FLASHINFER_ALLREDUCE=0
export LMDEPLOY_ENABLE_SYMM_MEM_ALLREDUCE=0
export NCCL_NVLS_ENABLE=0
lmdeploy serve api_server "$MODEL" \
  --backend pytorch --tp 1 --dp 8 --ep 8 \
  --dtype bfloat16 --quant-policy 0 \
  --max-batch-size 32 --session-len 8192 \
  --max-prefill-token-num 2048 --cache-max-entry-count 0.8 \
  --piecewise-cudagraph-max-tokens 2048 \
  --distributed-executor-backend ray \
  --speculative-algorithm deepseek_mtp --speculative-num-draft-tokens "$K" \
  --model-name glm53 --server-name 0.0.0.0 --server-port 24100 \
  --trust-remote-code --log-level INFO

MODEL=/path/to/GLM-5.3-Flash
K=3
export CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7
export OMP_NUM_THREADS=1
unset FLA_FLASH_KDA NCCL_NVLS_ENABLE
export VLLM_CUSTOM_SCOPES_FOR_PROFILING=1
export NVSHMEM_QP_DEPTH=8192
export NVSHMEM_REMOTE_TRANSPORT=none
vllm serve "$MODEL" \
  --host 0.0.0.0 --port 24200 --served-model-name glm53 \
  --tensor-parallel-size 1 --data-parallel-size 8 --data-parallel-size-local 8 \
  --enable-expert-parallel --all2all-backend deepep_low_latency \
  --data-parallel-rpc-port 24401 \
  --dtype bfloat16 --kv-cache-dtype bfloat16 --mamba-ssm-cache-dtype float32 \
  --max-model-len 8192 --max-num-seqs 32 --max-num-batched-tokens 2048 \
  --gpu-memory-utilization 0.9 --no-enable-prefix-caching --enable-chunked-prefill \
  --seed 42 --speculative-config "{\"method\":\"mtp\",\"num_speculative_tokens\":$K}" \
  --tool-call-parser glm47 --reasoning-parser glm47 --enable-auto-tool-choice

NVSHMEM_REMOTE_TRANSPORT=none is for this single-node NVLink run. vLLM uses a compatible DeepEP V1 build (567632d); LMDeploy uses its existing c997a4e build. Memory fractions have engine-specific meanings (LMDeploy 0.8 of free memory; vLLM 0.9 of total).

Accuracy

MMMU-Pro vision/test

Backend	Final score
LMDeploy, NCCL	1333/1730 (77.0520%)
vLLM	1329/1730 (76.8208%)

LMDeploy is +0.1720 percentage points from the 76.88% reference. Both engines finish with 0 request errors and 0 remaining length stops.

Protocol: all 1,730 vision/test items from dataset revision 563f3e84bb3b90893083a1f039cfa13077f2302b; original image bytes; identical image-first prompts; pinned NeMo prompt and answer-extraction helpers (778f31a). Temperature 1.0, top_p 0.95, top_k disabled, maximum 327,680 new tokens, seed 42, reasoning effort max. TP4/EP1/DP1/MTP5, client concurrency 8, maximum batch 16, session length 344,064, prefill chunk 8,192; LMDeploy uses NCCL. Both backends use FP8 checkpoint 3f1971b7b5f7a528c9c4ef6212c8785298a8c24a on disjoint groups of four H200 GPUs. NCCL_NVLS_ENABLE=0 is an explicit benchmark setting for both.

The 76.88% reference is a BF16 baseline reported in the NVIDIA GLM-5.3 Flash model card. This FP8 evaluation is not an exact reproduction of that baseline or its unpublished harness details.

Final scores include one fresh-service retry of the single LMDeploy length-stopped item, using the same sampling parameters and 327,680-token output cap. All length-selected replacements are retained regardless of correctness; incorrect answers alone do not trigger retries.

GPQA Diamond and SciCode

Evaluated revision: the scores below were measured on frozen LMDeploy commit 38779f12. The PR has since advanced to 18a8c917, including MoE-reduction changes. This table does not include a full-model accuracy rerun of those later commits.

Same GLM-5.3 Flash FP8 checkpoint as above; LMDeploy 38779f12, vLLM 606d124b. Both run TP4/EP1/DP1/MTP5 on separate four-H200 groups. LMDeploy uses NCCL; FlashInfer/symmetric-memory all-reduce is disabled. Temperature 1.0, top_p 0.95, top_k disabled, repetition penalty 1, seed 42, reasoning effort max, one sample per item. Client concurrency 8, maximum batch 16, BF16 KV, FP32 recurrent state, prefix caching off, prefill chunk 8,192, NCCL_NVLS_ENABLE=0.

Benchmark	LMDeploy NCCL	vLLM	LMDeploy − vLLM
GPQA Diamond	184/198 (92.9293%)	182/198 (91.9192%)	+1.0101 pp
SciCode with background: subproblem	151/288 (52.4306%)	143/288 (49.6528%)	+2.7778 pp

Generation starts with a 327,680-token output limit and context 393,216. Length-stopped items are retried once at 655,360 output tokens and context 720,896; SciCode also regenerates their dependent downstream steps. The table includes every selected replacement, regardless of whether it improves the score. Incorrect answers alone do not trigger retries.

SciCode evaluation fix: restore the published Maxwell, Block, and EnlargedBlock class definitions omitted by the pinned helper, then regenerate all 14 dependent subproblems for both engines. The final scores above include these corrected runs. No numerical assertion or tolerance was changed.

GPQA uses all 198 Diamond questions, frozen zero-shot prompts/option permutations validated against the source answers, and the pinned NeMo MCQ extractor. Cross-engine prompt-token count differences: 0 / 198.

SciCode uses all 65 test problems with background, 288 scored subproblems, and the three official prefilled steps excluded from the scored denominator. Main score requires every scored subproblem in that problem to pass. Each next prompt uses that engine’s previous generated code. Dataset revision; NeMo prompt/generation helpers. Original numerical assertions run in an isolated Python 3.11 environment with NumPy 1.26.4, SciPy 1.10.1, SymPy 1.12, h5py 3.11.0; 1,800-second per-step timeout, one BLAS thread, Python/NumPy execution seed 42. Development-reference validation passed 48/50: 70.8 and 78.3 also fail using the original SciCode helpers (large-phase numerical mismatch and wall-clock-dependent trajectory shape respectively). These are development cases; no test case was excluded and no scoring tolerance was relaxed

GSM8K

Backend	Correct / total	Accuracy
LMDeploy, NCCL	1,290 / 1,319	97.8014%
LMDeploy, optional FlashInfer	1,286 / 1,319	97.4981%
vLLM	1,288 / 1,319	97.6497%

TP4/MTP5, client concurrency 32, greedy, maximum 2,048 output tokens, identical prompt token IDs and parser. Errors and length stops remain in the denominator; explicit final answers are scored even at a length stop. Score proximity does not establish token parity or absence of regression.

[yidingcheng0206]

Keep KV-B sharding consistent with the attention TP group
When dp > 1 and attn_tp > 1, dp_disable_tp=True makes kv_b_proj retain all attention heads, while DeepseekV2BMM still shards kc/vc by attn_tp. The inherited process_weights_after_loading() then copies the full KV-B-derived weights directly into these sharded tensors.

[qescccczmr]

Fixed in 3fcabb7. Removed dp_disable_tp=True from GLM's KV-B projection so it uses the existing attention-TP sharding and weight loader, consistent with DeepseekV2BMM KC/VC absorption.

The focused CPU test exercises the real builder/loader/absorption in nine DP/TP/rank combinations and checks the exact local weight slices. This is loader-contract validation, not an EP+DP distributed-service validation; the unrelated EP+DP experiments are not included.

[RunningLeon]

@qescccczmr Hi, is this PR ready to review?

[RunningLeon]

can we use existing _state_select , _state_scatter in here?

lmdeploy/lmdeploy/pytorch/backends/cuda/gated_delta_rule.py

Line 201 in 9adf485

def _state_select(state, state_indices, spec_offsets):

lmdeploy/lmdeploy/pytorch/backends/cuda/gated_delta_rule.py

Line 232 in 9adf485

def _state_scatter(state, state_indices, spec_offsets, src):

[qescccczmr]

Addressed in 3fcabb7. KDA now reuses the existing GDN _state_select / _state_scatter helpers for both AR and MTP. The ordinary AR bank is exposed as a one-slot ring, rather than introducing another state kernel.

The shared helpers mask invalid state/slot IDs: reads return zero and writes are skipped, preventing padded rows from overwriting a live request. Actual CUDA tests cover FP32/BF16, irregular shapes, negative/out-of-range IDs, initialization, dummy/live-row collisions, and replay after updating graph inputs.

[RunningLeon]

mtp may need support as well

[qescccczmr]

Updated in 3fcabb7. KDA restores state using accepted history and saves every verified token's convolution/recurrent checkpoint, including partial rejection, reordering and dummy rows. MTP now also honors index_share_for_mtp_iteration through the existing DSATopKIndicesBuffer; cache updates still run on every draft step. Growing the shared buffer changes the graph key so old allocations are not replayed.

We also found and fixed a sampling-distribution bug: rejection/recovery did not apply AR's top-k/top-p/min-p filters. The fix reuses FusedLogitsProcessor filtering, preserves greedy fast-path behavior, and leaves source logits intact. Focused CUDA tests verify that filtered-out tokens cannot be accepted or recovered.

Mode	LMDeploy	vLLM
Greedy, fixed 256 tokens	62.9341%	62.7595%
Sampling, fixed 256 tokens	54.7305%	53.5599%
Greedy, natural EOS (max 256)	62.4700%	61.5550%
Sampling, natural EOS (max 256)	56.9082%	57.5145%

H200, TP4/EP1, CUDA Graph on, prefix cache off, BF16 KV, BS1, 1024-token prefill chunks. Inputs: 20 GSM8K prompts plus four 8436-token long prompts, each run under all four policies; sampling uses temperature=1, top_p=0.95, seed=42 and disabled top-k. Stop IDs/minimum lengths are matched.

[RunningLeon]

this may lead to poor performance. If fla does not provide this verification kernel, we may need to change the tilelang kernel of gated_delta_rule.

[qescccczmr]

Agreed; addressed in 3fcabb7 without a new kernel file. The existing GDN TileLang recurrent kernel's transposed-state path now supports KDA channelwise decay. Convolution windows are batched through FLA, then all verification tokens use one recurrent-kernel call. The kernel is parallel across batch/head/state tiles; causally dependent timesteps remain sequential inside that call, and each state is checkpointed for partial acceptance. The default GDN scalar-gate path is preserved.

Fresh CUDA regressions cover scalar/channelwise gates, AR versus multi-token state equality, dummy rows, graph replay and the original GDN tests. An earlier module-only CUDA Graph benchmark (3 verification tokens, 8 heads, head dim 128) measured 0.190/0.247/0.349 ms for the serial reference versus 0.057/0.059/0.065 ms for the shared path at BS1/4/16. This is KDA-local evidence, not a whole-model or vLLM performance-parity claim.

[qescccczmr]

@RunningLeon The MTP correctness/reuse changes are ready for another review in 3fcabb7. I am keeping this PR in Draft while whole-model accuracy/performance validation remains incomplete.

This update modifies 11 existing production files, with no new production module/kernel. It addresses the four review comments: attention-TP KV-B sharding, shared state helpers, accepted-history MTP state recovery, and reuse of GDN's TileLang verification kernel. It also fixes rejection sampling's missing AR filters, MTP index sharing and graph-buffer growth, and includes the GLM opt-in FP32 TP / mHC / stable sparse-index changes used for AR/MTP numerical stability. Existing defaults for other model callers are preserved. EP+DP, prefix-cache, FlashInfer communication and compact-MoE scheduling experiments are excluded.

[RunningLeon]

this layout could the be the same as Qwen3.5 and reuse the causal conv kernel when mtp is on.

lmdeploy/lmdeploy/pytorch/configurations/qwen3_5.py

Lines 65 to 66 in af4ad9b

	conv_kernel_size = text_config.linear_conv_kernel_dim + num_spec_tokens
	conv_state_shape = (num_delta_layers, conv_dim, conv_kernel_size)

[qescccczmr]

Updated in 41563c7.

The convolution cache now uses (num_linear_layers, conv_dim, conv_kernel_size + num_spec_tokens), matching Qwen3.5. MTP decode calls the existing causal_conv1d_update with accepted-history cache_seqlens and request state indices. This removes the per-draft convolution-window checkpoints; recurrent/KPool checkpoints remain unchanged. For kernel width 4 and MTP5, convolution-state storage drops from 24 to 9 elements per channel.

The shared kernel now accepts the actual weight/bias dtype, preserving GLM's FP32 convolution weights with BF16 activations. Prefill retains FLA's existing short-chunk handling and maps its chronological window into the compact token ring.

Validation on the isolated patch: 59 CPU checks passed; the CUDA-enabled regression run passed 75 checks, including MTP2/MTP5 partial-acceptance recovery (bit-exact against sequential AR in the tested KDA cases), request reordering/padding, and CUDA Graph replay. The 24 upstream optional Dao-reference cases were skipped because that dependency is unavailable; six additional independent FP16/BF16/FP32 non-circular convolution checks passed. Ruff and git diff --check also passed. These are local operator/contract checks, not a fresh whole-model acceptance or performance benchmark.

[RunningLeon]

[P1] Honor or explicitly reject the new MoE options across backend branches
build_fused_moe() now exposes fp32_acc and output_scale, but only the blocked-FP8 branch forwards them. The BF16 and W8A8 branches silently ignore these arguments, so the same public API has different numerical semantics depending on the selected backend.
This causes a concrete correctness issue for GLM’s BF16 path: the router returns unscaled weights and relies on output_scale=2.5 in the expert reduction. Dropping that argument produces an incorrectly scaled routed-expert contribution.
Please propagate these options through supported implementations, preserving the FP32 weighted reduction → output scaling → output cast order. For implementations that do not support this contract yet, explicitly reject non-default values rather than silently ignoring them. Keep False / 1.0 as defaults to preserve existing behavior, and add regression tests for parameter propagation, numerical correctness, and unsupported-backend rejection.
The scale must apply only to the routed-expert contribution, not to the shared-expert output.

[qescccczmr]

Fixed in 96423321.

• The BF16 fused-MoE path now forwards fp32_acc and output_scale through its build spec and applies them in the CUDA reduction as FP32 weighted reduction → scale → output cast.
• The scale is applied inside the routed-expert reduction, before the caller combines the result with the shared expert output.
• W8A8, static-FP8, W4A16, and DeepEP branches now explicitly reject non-default values instead of silently ignoring them.
• Added propagation, unsupported-backend, and CUDA numerical regression tests in tests/pytorch/nn/test_moe_options.py.

Validation: pytest -q tests/pytorch/nn/test_moe_options.py (4 passed, 1 CUDA test skipped because CUDA is unavailable in this environment); Ruff and compile checks pass.

[RunningLeon]

does glm5.3 need to do this on fp32?

[RunningLeon]

Where do we use this module?

[RunningLeon]

consider dlinfer backend for these two arguments

[RunningLeon]

why add this config?