[vorapolsiloai]

Motivation

ROCm (with NPU and XPU) keeps one ViT call per request in chunked prefill; the cross-request batch was excluded on ROCm in #32415 after a CI regression. With --mm-enable-dp-encoder, run_dp_sharded_mrope_vision_model splits the images of one ViT call across the attention-TP ranks. Called once per request, a single-image request lands on rank 0 while the other ranks wait, so the DP encoder never splits any work. Batching across requests is what gives it work to distribute.

Modifications

• hip_dp_encoder_batches_vit() (multimodal/mm_utils.py): true on ROCm when get_mm().mm_enable_dp_encoder and get_parallel().attn_tp_size > 1 (the width the DP encoder shards over).
• mm_schedule._batch_vit_across_requests(): CUDA batches as before; NPU/XPU keep the per-request path; ROCm batches only when the helper above is true. No new environment variable.
• Qwen3-VL (and its subclasses, e.g. Qwen3.5 / Qwen4-Exp) gets a precompile_kernels_after_loading hook that, under the same condition, runs the vision tower once on dummy inputs on every rank: grids (1,32,32) and (1,48,64), plus one call with _VECTORIZED_VL_POS_EMBED_MIN_IMAGES images so the vectorized pos-embed path loads its torch.unique kernels. Outputs are discarded. Without it, under the batched path a non-zero rank ran its first ViT call (including the rotary torch.compile) inside the first measured image burst: rank 1 compiled for ~18 s and TTFT p99 rose 128%. The startup warmup request has one image, so it only warms rank 0.
• Qwen4ExpForConditionalGeneration.precompile_kernels_after_loading now calls super() first.

Why TP1 and non-DP-encoder ROCm stay per-request

At attention-TP 1 there is nothing to distribute (the server already logs that the DP encoder "has no data-parallel work to distribute" at TP=1), and in our measurements batching at TP1 together with the ViT CUDA graph regressed (-1.8% / -2.6%). Without the DP encoder the behaviour that motivated #32415 still applies, so those paths are unchanged.

Accuracy

Outputs are not bit-exact against the per-request path, because GEMM tiling depends on M: embedding cosine 0.99996 mean, minimum token cosine 0.989. MMMU was checked on our internal stack; MMMU on this head is still to be measured.

Benchmark

Qwen3.8-27B, same-node A/B on MI350X screens (not MI355X), TP2 with --mm-enable-dp-encoder:

Variant	TP2 image throughput	TTFT
batched ViT, no warm-up	+8.6%	-8.3%
batched ViT + per-rank warm-up (confirmation run)	+11.7%

MI355X numbers and MMMU to be measured on this head.

Relation to #44727

Independent of the ViT CUDA graph PR (#44727). When both are enabled, each cross-request batch composition is a distinct graph layout there, bounded by its graph cache cap of 32.

Tests

• test/registered/unit/multimodal/test_hip_dp_encoder_vit_batching.py (CPU): the gate across HIP/NPU/XPU/CUDA, DP encoder on/off, attention-TP 1/2; two single-image requests encode in one call only under the gate; the post-load hook warms the tower once with the three dummy calls only under the gate, and is a no-op for language-model-only checkpoints.
• test/registered/chunked_prefill/test_mm_chunked_embedding_unit.py: switched from patching the removed mm_schedule._is_hip to override_platform(is_hip=...).

---

CI States

Latest PR Test (Base): ❌ Run #40494072702
Latest PR Test (Extra): ❌ Run #40494072546
Latest PR Test (AMD ROCm 10): ❌ Run #40494072515

[eppaneamd]

Measurements of this PR on GLM-5.3-Flash (Glm5NextForConditionalGeneration), TP 2, --mm-enable-dp-encoder, on 2x MI355X (gfx950).

Summary

• Once the vision tower is warm, this PR is faster: in the second image burst, TTFT mean −12.8 %, P99 −14.0 %, output throughput +9.3 %.
• On a fresh server, the first image burst is slower: TTFT is 0.4 to 1.3 s higher at every decile. With this PR, ranks other than 0 run the vision tower for the first time inside that burst.
• A GLM vision warm-up (branch below) moves the vision compiles before the server reports ready, for 10 s more launch time. The first burst's TTFT is then 0.9 to 1.4 s lower at every decile and within 0.12 s of the second burst.
• With the warm-up, no compile was logged after ready for image sizes from 112x112 to 2464x2520.

Setup

• Image lmsysorg/sglang-rocm:v0.5.22-rocm724-mi35x-20261024; SGLang imported through PYTHONPATH from the python/ tree of this PR (006d29d3a4, since rebased to 96d4c41944 with the same change) or of its parent 9b46f27090 (base).
• amd/GLM-5.3-Flash-Quark-MXFP4 at revision b5688f25491202978c19c4d036eef579f61bbe07.
• A fresh container per build, HIP_VISIBLE_DEVICES=0,1. Each server is first warmed with text-only requests, so only the vision tower is cold, then gets two image bursts of 64 requests at concurrency 64 (512x512, 1024 input / 128 output tokens, seed 42). Every request completed.
• One launch per build. This PR's first burst differs by about 10 % between the two runs below (mean 3606.3 vs 3286.0 ms), which bounds how far single-launch differences can be read.

Server command (the same for every build)

python3 -m sglang.launch_server \
  --model-path amd/GLM-5.3-Flash-Quark-MXFP4 \
  --revision b5688f25491202978c19c4d036eef579f61bbe07 \
  --served-model-name amd/GLM-5.3-Flash-Quark-MXFP4 \
  --quantization quark --tp-size 2 \
  --attention-backend dsa --dsa-prefill-backend tilelang --dsa-decode-backend tilelang \
  --linear-attn-backend triton --kv-cache-dtype bfloat16 \
  --moe-runner-backend aiter --disable-shared-experts-fusion \
  --disable-radix-cache --max-running-requests 64 \
  --cuda-graph-backend-decode full --cuda-graph-max-bs-decode 64 \
  --mem-fraction-static 0.92 \
  --model-loader-extra-config '{"enable_multithread_load":true,"num_threads":8}' \
  --trust-remote-code --tool-call-parser glm47 --reasoning-parser glm45 \
  --enable-multimodal --limit-mm-data-per-request '{"image":1}' \
  --mm-process-config '{"image":{"min_image_tokens":16,"max_image_tokens":8000}}' \
  --mm-preprocess-cache-size-mb 0 --mm-enable-dp-encoder --mm-attention-backend aiter_attn \
  --context-length 8192 --host 127.0.0.1 --port 30000

This PR vs base

Build	Burst	Mean TTFT (ms)	P99 TTFT (ms)	Output tok/s
base	first	2718.5	3798.0	1184.6
base	second	2318.9	3347.0	1271.4
this PR	first	3606.3	4539.7	1006.3
this PR	second	2023.2	2877.9	1389.1

In this PR's first burst, rank 1's only logged compiles are its first apply_rotary_pos_emb_native and swiglu_clamped compiles (TORCH_LOGS=recompiles, this PR's build only). That fits the TTFT shift; the attribution is not established.

With a GLM vision warm-up

eppaneamd/sglang glm-vision-warm-up (c99d0a7cf8, on top of this PR; cherry-picks cleanly onto 96d4c41944) adds a GLM-5-Next precompile_kernels_after_loading that runs the vision tower on two dummy grids when hip_dp_encoder_batches_vit() holds. It reuses Qwen3-VL's warm-up loop, moved into mm_utils.warm_up_vision_tower. The timings below are from an earlier version with a third, two-image grid, which logged no compile and is dropped in the commit. TORCH_LOGS=recompiles on both builds:

Build	Burst	Mean TTFT (ms)	P99 TTFT (ms)	Output tok/s
this PR	first	3286.0	4199.6	1109.6
this PR	second	2391.6	3247.0	1295.9
this PR + warm-up	first	2011.2	2932.3	1370.7
this PR + warm-up	second	1870.8	2761.6	1416.8

With the warm-up, all vision compiles end before the server reports ready (launch to health 241 s, against 231 s without; 5 s poll). Without it, 6 compiles fell in the bursts (2 on rank 0, 4 on rank 1).

Image sizes

On the two-grid commit, one fresh server with the same flags got image bursts through the chat API at 112x112, 512x512, 1024x1024, 1920x1080, mixed sizes and 2464x2520 (up to 9 of the largest at once), from the 16-token floor to near the 8000-token ceiling of --mm-process-config (7920 image tokens; larger limits not tested). No compile or recompile was logged after the server reported ready, and all 93 requests completed.