Skip to content

[JIT][ROCm] Fully warmup GLM-5.3-Flash - #60919

Merged
dllehr-amd merged 22 commits into
vllm-project:mainfrom
simondanielsson:feature/warmup-glm5next
Nov 4, 2026
Merged

dllehr-amd merged 22 commits into
vllm-project:mainfrom
simondanielsson:feature/warmup-glm5next

Conversation

@simondanielsson

@simondanielsson simondanielsson commented Oct 9, 2026 •

Copy link
Copy Markdown
Contributor

Overview

It used to JIT several triton kernels during runtime. This PR

  1. registers those kernels for warmup.
  2. fixes a bug where even after registering for warmup, triton re-compiled during runtime on ROCm

Explanation for the triton fix::

  • Triton on AMD specializes kernels on buffer size being below or above 2GiB, i.e int32 addressable.
  • When <2GiB, this allows the compiler to lower to buffer load/store intrinsics which require 32-bit offsets
  • The warmup code previously assumed all buffers passed during warmup were <2GiB.
  • In kpool_tail_seed_dispatch we pass the kv cache which can be large. During warmup it specialized on a <2GiB KV cache, whereas at runtime it gets passed a larger KV cache and re-compiles without buffer loads. See https://github.com/ROCm/triton/blob/0ec280cf80dd91e9a86887981a670f2d4541a32b/third_party/amd/backend/compiler.py#L76-L110 for details
  • This PR removes that assumption and instead passes the real buffer size

Claims

  • No runtime compilation on GLM-5.3-Flash, i.e. doesn't crash with --jit-monitor-mode error

Validation

Expect lots of runtime compilation on main, none on this branch

VLLM_ROCM_USE_AITER=1 \
VLLM_ROCM_USE_AITER_TRITON_SPARSE_MLA=1 \
VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS=1 \
VLLM_USE_BREAKABLE_CUDAGRAPH=1 \
VLLM_USE_RUST_FRONTEND=1 \
vllm serve zai-org/GLM-5.3-Flash \
    --tensor-parallel-size 8 \
    --max-num-seqs 512 \
    --gpu-memory-utilization 0.8 \
    --attention-backend ROCM_AITER_MLA_SPARSE \
    --max-num-batched-tokens 16384 \
    --jit-monitor-mode warn
    --jit-monitor-verbose
curl http://localhost:8000/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{
    "model": "zai-org/GLM-5.3-Flash",
    "messages": [
      {"role": "system", "content": "You are a helpful assistant."},
      {"role": "user", "content": "Tell me a short joke. Think for 50 tokens before answeing"}
    ],
    "temperature": 0.7,
    "max_tokens": 100
  }'

Main: 53 compile keys and runtime compilation

JIT kernel warmup (53 compile keys): 100%|██████████| 9/9 [00:42<00:00,  4.75s/kernel, _DecoratedTritonJitKernel (8 keys)]

...

(RustFrontend pid=3291896) INFO 10-20 16:14:59 [server/src/lib.rs:212] starting vLLM server model="zai-org/GLM-5.3-Flash"
(RustFrontend pid=3291896) INFO 10-20 16:14:59 [server/src/lib.rs:324] OpenAI server is ready to accept requests bind_address="0.0.0.0:8000" scheme="http" model="zai-org/GLM-5.3-Flash"
(EngineCore pid=3291730) INFO 10-20 16:14:59 [kernel.py:476] Final IR op priority after setting platform defaults: IrOpPriorityConfig(rms_norm=['aiter', 'vllm_c', 'native'], fused_add_rms_norm=['aiter', 'vllm_c', 'native'], gelu_and_mul_sparse=['native'])
(EngineCore pid=3291730) INFO 10-20 16:14:59 [compilation.py:317] Enabled custom fusions: norm_quant, act_quant, allreduce_rms, mla_dual_rms_norm, qk_norm_rope_kvcache
(MainProcess pid=3289259) INFO 10-20 16:14:59 [utils.py:551] Waiting for API servers to complete ...
(RustFrontend pid=3291896) INFO 10-20 16:23:19 [runtime.rs:48] capping request runtime worker threads, set VLLM_RS_REQUEST_WORKER_THREADS to override available_parallelism=384 capped_worker_threads=32
(RustFrontend pid=3291896) INFO 10-20 16:23:19 [default/mod.rs:244] chat_completions{request_id=chatcmpl-f0713618}: using reasoning parser parser_name="glm47"
(Worker_TP7 pid=3295066) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: BuildPrefillChunkMetadataKernel.kernel (constexprs={BLOCK_SIZE=1024, COMPRESS_RATIO=4, DCP_INTERLEAVE=1, DCP_WORLD=1}; signature={BLOCK_SIZE='constexpr', COMPRESS_RATIO='constexpr', DCP_INTERLEAVE='constexpr', DCP_RANK='i32', DCP_WORLD='constexpr', cu_compressed_seq_len_ke_ptr='*i32', cu_compressed_seq_len_ks_ptr='*i32', cu_compressed_seq_lens_ptr='*i32', query_slice_start='i32', query_slice_stop='i32', query_start_loc_ptr='*i32', row_start_cu_compressed_seq_lens_ptr='*i32', token_to_seq_ptr='*i32', uncompressed_seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=7, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.v1.attention.backends.mla.indexer.BuildPrefillChunkMetadataKernel.kernel", "signature": {"query_start...}; key="[('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('i...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP6 pid=3294843) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: BuildPrefillChunkMetadataKernel.kernel (constexprs={BLOCK_SIZE=1024, COMPRESS_RATIO=4, DCP_INTERLEAVE=1, DCP_WORLD=1}; signature={BLOCK_SIZE='constexpr', COMPRESS_RATIO='constexpr', DCP_INTERLEAVE='constexpr', DCP_RANK='i32', DCP_WORLD='constexpr', cu_compressed_seq_len_ke_ptr='*i32', cu_compressed_seq_len_ks_ptr='*i32', cu_compressed_seq_lens_ptr='*i32', query_slice_start='i32', query_slice_stop='i32', query_start_loc_ptr='*i32', row_start_cu_compressed_seq_lens_ptr='*i32', token_to_seq_ptr='*i32', uncompressed_seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=6, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.v1.attention.backends.mla.indexer.BuildPrefillChunkMetadataKernel.kernel", "signature": {"query_start...}; key="[('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('i...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP5 pid=3294656) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: BuildPrefillChunkMetadataKernel.kernel (constexprs={BLOCK_SIZE=1024, COMPRESS_RATIO=4, DCP_INTERLEAVE=1, DCP_WORLD=1}; signature={BLOCK_SIZE='constexpr', COMPRESS_RATIO='constexpr', DCP_INTERLEAVE='constexpr', DCP_RANK='i32', DCP_WORLD='constexpr', cu_compressed_seq_len_ke_ptr='*i32', cu_compressed_seq_len_ks_ptr='*i32', cu_compressed_seq_lens_ptr='*i32', query_slice_start='i32', query_slice_stop='i32', query_start_loc_ptr='*i32', row_start_cu_compressed_seq_lens_ptr='*i32', token_to_seq_ptr='*i32', uncompressed_seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=5, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.v1.attention.backends.mla.indexer.BuildPrefillChunkMetadataKernel.kernel", "signature": {"query_start...}; key="[('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('i...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP4 pid=3294470) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: BuildPrefillChunkMetadataKernel.kernel (constexprs={BLOCK_SIZE=1024, COMPRESS_RATIO=4, DCP_INTERLEAVE=1, DCP_WORLD=1}; signature={BLOCK_SIZE='constexpr', COMPRESS_RATIO='constexpr', DCP_INTERLEAVE='constexpr', DCP_RANK='i32', DCP_WORLD='constexpr', cu_compressed_seq_len_ke_ptr='*i32', cu_compressed_seq_len_ks_ptr='*i32', cu_compressed_seq_lens_ptr='*i32', query_slice_start='i32', query_slice_stop='i32', query_start_loc_ptr='*i32', row_start_cu_compressed_seq_lens_ptr='*i32', token_to_seq_ptr='*i32', uncompressed_seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=4, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.v1.attention.backends.mla.indexer.BuildPrefillChunkMetadataKernel.kernel", "signature": {"query_start...}; key="[('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('i...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP3 pid=3294281) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: BuildPrefillChunkMetadataKernel.kernel (constexprs={BLOCK_SIZE=1024, COMPRESS_RATIO=4, DCP_INTERLEAVE=1, DCP_WORLD=1}; signature={BLOCK_SIZE='constexpr', COMPRESS_RATIO='constexpr', DCP_INTERLEAVE='constexpr', DCP_RANK='i32', DCP_WORLD='constexpr', cu_compressed_seq_len_ke_ptr='*i32', cu_compressed_seq_len_ks_ptr='*i32', cu_compressed_seq_lens_ptr='*i32', query_slice_start='i32', query_slice_stop='i32', query_start_loc_ptr='*i32', row_start_cu_compressed_seq_lens_ptr='*i32', token_to_seq_ptr='*i32', uncompressed_seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=3, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.v1.attention.backends.mla.indexer.BuildPrefillChunkMetadataKernel.kernel", "signature": {"query_start...}; key="[('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('i...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP2 pid=3292443) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: BuildPrefillChunkMetadataKernel.kernel (constexprs={BLOCK_SIZE=1024, COMPRESS_RATIO=4, DCP_INTERLEAVE=1, DCP_WORLD=1}; signature={BLOCK_SIZE='constexpr', COMPRESS_RATIO='constexpr', DCP_INTERLEAVE='constexpr', DCP_RANK='i32', DCP_WORLD='constexpr', cu_compressed_seq_len_ke_ptr='*i32', cu_compressed_seq_len_ks_ptr='*i32', cu_compressed_seq_lens_ptr='*i32', query_slice_start='i32', query_slice_stop='i32', query_start_loc_ptr='*i32', row_start_cu_compressed_seq_lens_ptr='*i32', token_to_seq_ptr='*i32', uncompressed_seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=2, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.v1.attention.backends.mla.indexer.BuildPrefillChunkMetadataKernel.kernel", "signature": {"query_start...}; key="[('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('i...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP1 pid=3292257) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: BuildPrefillChunkMetadataKernel.kernel (constexprs={BLOCK_SIZE=1024, COMPRESS_RATIO=4, DCP_INTERLEAVE=1, DCP_WORLD=1}; signature={BLOCK_SIZE='constexpr', COMPRESS_RATIO='constexpr', DCP_INTERLEAVE='constexpr', DCP_RANK='i32', DCP_WORLD='constexpr', cu_compressed_seq_len_ke_ptr='*i32', cu_compressed_seq_len_ks_ptr='*i32', cu_compressed_seq_lens_ptr='*i32', query_slice_start='i32', query_slice_stop='i32', query_start_loc_ptr='*i32', row_start_cu_compressed_seq_lens_ptr='*i32', token_to_seq_ptr='*i32', uncompressed_seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=1, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.v1.attention.backends.mla.indexer.BuildPrefillChunkMetadataKernel.kernel", "signature": {"query_start...}; key="[('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('i...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP0 pid=3292060) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: BuildPrefillChunkMetadataKernel.kernel (constexprs={BLOCK_SIZE=1024, COMPRESS_RATIO=4, DCP_INTERLEAVE=1, DCP_WORLD=1}; signature={BLOCK_SIZE='constexpr', COMPRESS_RATIO='constexpr', DCP_INTERLEAVE='constexpr', DCP_RANK='i32', DCP_WORLD='constexpr', cu_compressed_seq_len_ke_ptr='*i32', cu_compressed_seq_len_ks_ptr='*i32', cu_compressed_seq_lens_ptr='*i32', query_slice_start='i32', query_slice_stop='i32', query_start_loc_ptr='*i32', row_start_cu_compressed_seq_lens_ptr='*i32', token_to_seq_ptr='*i32', uncompressed_seq_lens_ptr='*i32'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=0, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.v1.attention.backends.mla.indexer.BuildPrefillChunkMetadataKernel.kernel", "signature": {"query_start...}; key="[('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('*i32', 'DS'), ('i...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP6 pid=3294843) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: _kpool_tail_seed_kernel (constexprs={BLOCK_D=128, HEAD_DIM=128, KPOOL=4, KPOOL_HEAD=512, RING=4, TAIL_BLOCK_ELEMS=3734016}; signature={BLOCK_D='constexpr', HEAD_DIM='constexpr', KPOOL='constexpr', KPOOL_HEAD='constexpr', RING='constexpr', TAIL_BLOCK_ELEMS='constexpr', key_ptr='*bf16', n_tokens='i32', score_ptr='*bf16', tail_ptr='*bf16', tslot_ptr='*i64'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=6, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.models.glm5next.amd.ops.kpool_compress._kpool_tail_seed_kernel", "signature": {"key_ptr": "*bf16", "s...}; key="[('*bf16', 'DS'), ('*bf16', 'DS'), ('*i64', 'DS'), ('*bf16', 'D'), ('i32', ''), ('constexpr', 3734016), ('constexpr'...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP5 pid=3294656) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: _kpool_tail_seed_kernel (constexprs={BLOCK_D=128, HEAD_DIM=128, KPOOL=4, KPOOL_HEAD=512, RING=4, TAIL_BLOCK_ELEMS=3734016}; signature={BLOCK_D='constexpr', HEAD_DIM='constexpr', KPOOL='constexpr', KPOOL_HEAD='constexpr', RING='constexpr', TAIL_BLOCK_ELEMS='constexpr', key_ptr='*bf16', n_tokens='i32', score_ptr='*bf16', tail_ptr='*bf16', tslot_ptr='*i64'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=5, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.models.glm5next.amd.ops.kpool_compress._kpool_tail_seed_kernel", "signature": {"key_ptr": "*bf16", "s...}; key="[('*bf16', 'DS'), ('*bf16', 'DS'), ('*i64', 'DS'), ('*bf16', 'D'), ('i32', ''), ('constexpr', 3734016), ('constexpr'...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP7 pid=3295066) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: _kpool_tail_seed_kernel (constexprs={BLOCK_D=128, HEAD_DIM=128, KPOOL=4, KPOOL_HEAD=512, RING=4, TAIL_BLOCK_ELEMS=3734016}; signature={BLOCK_D='constexpr', HEAD_DIM='constexpr', KPOOL='constexpr', KPOOL_HEAD='constexpr', RING='constexpr', TAIL_BLOCK_ELEMS='constexpr', key_ptr='*bf16', n_tokens='i32', score_ptr='*bf16', tail_ptr='*bf16', tslot_ptr='*i64'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=7, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.models.glm5next.amd.ops.kpool_compress._kpool_tail_seed_kernel", "signature": {"key_ptr": "*bf16", "s...}; key="[('*bf16', 'DS'), ('*bf16', 'DS'), ('*i64', 'DS'), ('*bf16', 'D'), ('i32', ''), ('constexpr', 3734016), ('constexpr'...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP4 pid=3294470) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: _kpool_tail_seed_kernel (constexprs={BLOCK_D=128, HEAD_DIM=128, KPOOL=4, KPOOL_HEAD=512, RING=4, TAIL_BLOCK_ELEMS=3734016}; signature={BLOCK_D='constexpr', HEAD_DIM='constexpr', KPOOL='constexpr', KPOOL_HEAD='constexpr', RING='constexpr', TAIL_BLOCK_ELEMS='constexpr', key_ptr='*bf16', n_tokens='i32', score_ptr='*bf16', tail_ptr='*bf16', tslot_ptr='*i64'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=4, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.models.glm5next.amd.ops.kpool_compress._kpool_tail_seed_kernel", "signature": {"key_ptr": "*bf16", "s...}; key="[('*bf16', 'DS'), ('*bf16', 'DS'), ('*i64', 'DS'), ('*bf16', 'D'), ('i32', ''), ('constexpr', 3734016), ('constexpr'...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP3 pid=3294281) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: _kpool_tail_seed_kernel (constexprs={BLOCK_D=128, HEAD_DIM=128, KPOOL=4, KPOOL_HEAD=512, RING=4, TAIL_BLOCK_ELEMS=3734016}; signature={BLOCK_D='constexpr', HEAD_DIM='constexpr', KPOOL='constexpr', KPOOL_HEAD='constexpr', RING='constexpr', TAIL_BLOCK_ELEMS='constexpr', key_ptr='*bf16', n_tokens='i32', score_ptr='*bf16', tail_ptr='*bf16', tslot_ptr='*i64'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=3, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.models.glm5next.amd.ops.kpool_compress._kpool_tail_seed_kernel", "signature": {"key_ptr": "*bf16", "s...}; key="[('*bf16', 'DS'), ('*bf16', 'DS'), ('*i64', 'DS'), ('*bf16', 'D'), ('i32', ''), ('constexpr', 3734016), ('constexpr'...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP2 pid=3292443) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: _kpool_tail_seed_kernel (constexprs={BLOCK_D=128, HEAD_DIM=128, KPOOL=4, KPOOL_HEAD=512, RING=4, TAIL_BLOCK_ELEMS=3734016}; signature={BLOCK_D='constexpr', HEAD_DIM='constexpr', KPOOL='constexpr', KPOOL_HEAD='constexpr', RING='constexpr', TAIL_BLOCK_ELEMS='constexpr', key_ptr='*bf16', n_tokens='i32', score_ptr='*bf16', tail_ptr='*bf16', tslot_ptr='*i64'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=2, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.models.glm5next.amd.ops.kpool_compress._kpool_tail_seed_kernel", "signature": {"key_ptr": "*bf16", "s...}; key="[('*bf16', 'DS'), ('*bf16', 'DS'), ('*i64', 'DS'), ('*bf16', 'D'), ('i32', ''), ('constexpr', 3734016), ('constexpr'...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP1 pid=3292257) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: _kpool_tail_seed_kernel (constexprs={BLOCK_D=128, HEAD_DIM=128, KPOOL=4, KPOOL_HEAD=512, RING=4, TAIL_BLOCK_ELEMS=3734016}; signature={BLOCK_D='constexpr', HEAD_DIM='constexpr', KPOOL='constexpr', KPOOL_HEAD='constexpr', RING='constexpr', TAIL_BLOCK_ELEMS='constexpr', key_ptr='*bf16', n_tokens='i32', score_ptr='*bf16', tail_ptr='*bf16', tslot_ptr='*i64'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=1, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.models.glm5next.amd.ops.kpool_compress._kpool_tail_seed_kernel", "signature": {"key_ptr": "*bf16", "s...}; key="[('*bf16', 'DS'), ('*bf16', 'DS'), ('*i64', 'DS'), ('*bf16', 'D'), ('i32', ''), ('constexpr', 3734016), ('constexpr'...). This causes a latency spike; consider extending warmup to cover this shape/config.
(Worker_TP0 pid=3292060) WARNING 10-20 16:23:19 [jit_monitor.py:137] Triton kernel JIT compilation during inference: _kpool_tail_seed_kernel (constexprs={BLOCK_D=128, HEAD_DIM=128, KPOOL=4, KPOOL_HEAD=512, RING=4, TAIL_BLOCK_ELEMS=3734016}; signature={BLOCK_D='constexpr', HEAD_DIM='constexpr', KPOOL='constexpr', KPOOL_HEAD='constexpr', RING='constexpr', TAIL_BLOCK_ELEMS='constexpr', key_ptr='*bf16', n_tokens='i32', score_ptr='*bf16', tail_ptr='*bf16', tslot_ptr='*i64'}; extra_compile_info={configs=[{(0,): [['tt.divisibility', 16], ['tt.pointer_range', 32]], (1,): [['tt.divisibility', 16], ['tt.pointer_range', 32]..., device=0, enable_fp_fusion=True, extern_libs=(('ocml', '/usr/local/lib/python3.12/dist-packages/triton/backends/amd/lib/ocml.bc'), ('ockl', '/usr/local/lib/python..., is_warmup=False, launch_cooperative_grid=False, num_ctas=1, num_stages=2, num_warps=4, specialization_data='{"name": "vllm.models.glm5next.amd.ops.kpool_compress._kpool_tail_seed_kernel", "signature": {"key_ptr": "*bf16", "s...}; key="[('*bf16', 'DS'), ('*bf16', 'DS'), ('*i64', 'DS'), ('*bf16', 'D'), ('i32', ''), ('constexpr', 3734016), ('constexpr'...). This causes a latency spike; consider extending warmup to cover this shape/config.
(RustFrontend pid=3291896) INFO 10-20 16:23:29 [log_stats.rs:273] Avg prompt tput: 3.4 toks/s, Avg generation tput: 10.0 toks/s, Reqs Running: 0, Waiting: 0, GPU KV cache used: 0.0%, Prefix cache hit rate: 0.0%

Here we can see the issue: at runtime one of the parameters are specialized as ('*bf16', 'D') (i.e., not DS, so without an S suffix) which means it's >2GiB. Here's the code that does this.

This branch: 167 compile keys and no runtime jit

JIT kernel warmup (167 compile keys):  73%|███████▎  | 11/15 [00:27<00:06,  1.69s/kernel, BuildPrefillChunkMetadataKernel (72 keys)]

...

(RustFrontend pid=3211463) INFO 10-20 16:00:14 [server/src/lib.rs:212] starting vLLM server model="zai-org/GLM-5.3-Flash"
(RustFrontend pid=3211463) INFO 10-20 16:00:14 [server/src/lib.rs:324] OpenAI server is ready to accept requests bind_address="0.0.0.0:8000" scheme="http" model="zai-org/GLM-5.3-Flash"
(EngineCore pid=3211294) INFO 10-20 16:00:14 [kernel.py:476] Final IR op priority after setting platform defaults: IrOpPriorityConfig(rms_norm=['aiter', 'vllm_c', 'native'], fused_add_rms_norm=['aiter', 'vllm_c', 'native'], gelu_and_mul_sparse=['native'])
(EngineCore pid=3211294) INFO 10-20 16:00:14 [compilation.py:317] Enabled custom fusions: norm_quant, act_quant, allreduce_rms, mla_dual_rms_norm, qk_norm_rope_kvcache
(RustFrontend pid=3211463) INFO 10-20 16:07:16 [runtime.rs:48] capping request runtime worker threads, set VLLM_RS_REQUEST_WORKER_THREADS to override available_parallelism=384 capped_worker_threads=32
(RustFrontend pid=3211463) INFO 10-20 16:07:16 [default/mod.rs:244] chat_completions{request_id=chatcmpl-14baf43f}: using reasoning parser parser_name="glm47"
(RustFrontend pid=3211463) INFO 10-20 16:07:24 [log_stats.rs:273] Avg prompt tput: 3.4 toks/s, Avg generation tput: 10.0 toks/s, Reqs Running: 0, Waiting: 0, GPU KV cache used: 0.0%, Prefix cache hit rate: 0.0%

This branch with --jit-monitor-mode error: no crash and ok output

{"id":"chatcmpl-a5f240f6","object":"chat.completion","created":1792521728,"model":"zai-org/GLM-5.3-Flash","choices":[{"index":0,"message":{"role":"assistant","content":null,"reasoning":"The user wants a short joke and asked me to think for 50 tokens before answering. Let me think of a good short joke.\n\nSome options:\n1. Why don't scientists trust atoms? Because they make up everything!\n2. What do you call a fake noodle? An impasta!\n3. Why did the scarecrow win an award? He was outstanding in his field!\n\nLet me pick one that's classic and reliably funny. The atoms one is a solid choice - it's short"},"logprobs":null,"finish_reason":"length","stop_reason":null,"token_ids":null}],"usage":{"prompt_tokens":34,"total_tokens":134,"completion_tokens":100,"prompt_tokens_details":null,"completion_tokens_details":{"reasoning_tokens":100}},"metrics":null,"system_fingerprint":null,"prompt_logprobs":null,"prompt_token_ids":null,"kv_transfer_params":null,"ec_transfer_params":null}

Details


Pull Request Checklist
  • I used vLLM's /pr-checklist skill. (Mandatory for agents, optional for humans).

  • AI assistance was used during the creation of this PR.

  • Design Fit: Minimizes impact on core components, reuses existing functionality, and justifies added complexity.

  • Testing and Validation: Validates the change and ensures any added tests are meaningful and reliable, with CI coverage or documented CI resource constraints and validation performed outside CI.

  • Code Quality and Style: Keeps code and comments clear and concise, and updates relevant documentation and examples.

  • Pull Request Contents: Includes a brief summary and relevant links, supports claims with evidence, explains root causes and implementation trade-offs, and follows the contributing guide.

Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
simondanielsson and others added 7 commits October 20, 2026 09:49
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
@simondanielsson simondanielsson changed the title [JIT][ROCm] Warmup GLM-5.3-Flash [JIT][ROCm] Fully warmup GLM-5.3-Flash Oct 20, 2026
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>

@claude claude Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Claude Code Review

This pull request is from a fork — automated review is disabled. A repository maintainer can comment @claude review to run a one-time review.

@GodRishUniverse

GodRishUniverse commented Oct 22, 2026 •

Copy link
Copy Markdown
Contributor

@simondanielsson #60084 is also open and is approved. Please check that and collaborate. Thanks!

@simondanielsson

simondanielsson commented Oct 22, 2026 •

Copy link
Copy Markdown
Contributor Author

@simondanielsson #60084 is also open and is approved. Please check that and collaborate. Thanks!

@GodRishUniverse Thanks! IIUC that fix seems orthogonal to the changes here but indeed it's on the same topic

This PR doesn't try to convince triton that the k and v is indeed <2GiB. Instead it adds the option to specialize the kernel on >2GiB buffers during warmup instead of assuming it's always 32-bit addressable

@GodRishUniverse

Copy link
Copy Markdown
Contributor

@simondanielsson #60084 is also open and is approved. Please check that and collaborate. Thanks!

@GodRishUniverse Thanks! IIUC that fix seems orthogonal to the changes here but indeed it's on the same topic

Cool. Thanks for the quick response!

kv_cache = k_cache.kv_cache
page_size = kv_cache.shape[1]
buf_numel_per_page = kv_cache.stride(0)
buf_ptr_range_bytes = triton_ptr_range(kv_cache)

@simondanielsson simondanielsson Oct 23, 2026 •

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Note: This is where the specialization otherwise broke. If not specifying that the buffer it would warmup assuming <2GiB, and at runtime re-compile for >2GiB

@mergify mergify Bot removed the needs-rebase label Oct 23, 2026
@okorzh-amd

Copy link
Copy Markdown
Contributor

LGTM. Added comment on the test.

Comment thread vllm/models/glm5next/common/kda.py Resolved
if backend_name != "DEEPSEEK_V32_INDEXER":
if backend_name not in (
"DEEPSEEK_V32_INDEXER",
"ROCM_AITER_MLA_SPARSE",

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you double check skipping this warmup is okay/needed?

@simondanielsson simondanielsson Oct 28, 2026 •

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Skipping is OK but not really related, so I removed it here.

Details if curious:

That specific warmup is only needed for models using Deepseek SWA (dsv4 and dsv41), so I think the actual proper fix is

if backend_name in (ROCM_FLASHMLA_SPARSE_DSV4, ROCM_FLASHMLA_SPARSE_DSV4)

So I'd vouch for removing my change and instead defer this change to another PR as it'd require more validation.

simondanielsson and others added 3 commits October 28, 2026 11:46
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
@simondanielsson simondanielsson added the ready ONLY add when PR is ready to merge/full CI is needed label Oct 29, 2026
@simondanielsson

Copy link
Copy Markdown
Contributor Author

/ci run

@github-actions

Copy link
Copy Markdown

✅ Triggered Buildkite CI #98361 for commit f69558138601.

…next

Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
@simondanielsson

Copy link
Copy Markdown
Contributor Author

/ci run

@github-actions

Copy link
Copy Markdown

✅ Triggered Buildkite CI #98429 for commit 00c873fa4766.

@eppaneamd

Copy link
Copy Markdown

Optional, not blocking: the kpool kernels can warm with 2 compiles instead of 8

This PR warms 8 variants of the two kpool kernels: _kpool_softmax_rotate_write_cache_kernel for loc_ptr aligned or not, and _kpool_tail_seed_kernel for tslot_ptr aligned or not × n_tokens = 1, a multiple of 16, or other. On gfx950 they compile to the same code, except that the n_tokens == 1 variant compares the mask against a constant instead of the argument (same instruction count). loc_ptr and tslot_ptr are only read by scalar loads, and n_tokens only feeds the i + KPOOL < n_tokens mask.

Proposed fix: opt the two kernels out of these specializations:

@triton.jit(do_not_specialize_on_alignment=["loc_ptr"])
def _kpool_softmax_rotate_write_cache_kernel(...)

@triton.jit(do_not_specialize=["n_tokens"], do_not_specialize_on_alignment=["tslot_ptr"])
def _kpool_tail_seed_kernel(...)

Measured through this PR's warm-up at its head on MI355X: 2 + 6 compiles become 1 + 1, and the warm-up of the two kernels takes 538 ms instead of 769–775 ms per process. In both cases, 24 prefill-write batches after the warm-up compile nothing and write identical cache bytes. With JIT warm-up off (--kernel_config.enable_jit_warmup=False, or --enforce-eager unless fault tolerance is enabled), serving compiles at most 2 variants instead of up to 8.

Setup
  • MI355X (gfx950), vllm/vllm-openai-rocm:nightly-rocm72-3966924d56 (Triton 3.8.0), this PR's vllm tree with and without the decorators on PYTHONPATH. get_warmup_keys and warmup of kpool_compress_dispatch and kpool_tail_seed_dispatch, then the AMD indexer's prefill writes for 0–3 decode and 1–1024 prefill tokens at GLM-5.3-Flash index shapes. Two runs of each in the order PR, decorators, decorators, PR, with an empty Triton cache per run.
  • ISA: gfx950 code of the decorated kernels compared with all 8 original variants, at the 4 warps they launch with.
  • Scope: the AMD kernels only; the NVIDIA twins in nvidia/ops/kpool_compress.py are not checked.

simondanielsson and others added 3 commits October 30, 2026 14:02
Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
…next

Signed-off-by: simondanielsson <simon.danielsson99@hotmail.com>
@simondanielsson

Copy link
Copy Markdown
Contributor Author

/ci run

@github-actions

Copy link
Copy Markdown

✅ Triggered Buildkite CI #98806 for commit 3beb49357534.

@dllehr-amd dllehr-amd left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good to me. Thanks @simondanielsson!

@dllehr-amd
dllehr-amd merged commit beaa726 into vllm-project:main Nov 4, 2026
253 of 254 checks passed
@github-project-automation github-project-automation Bot moved this from Todo to Done in AMD Nov 4, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

glm ready ONLY add when PR is ready to merge/full CI is needed rocm Related to AMD ROCm warmup Related to the warmup system

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

6 participants