Summary
Several in-tree Triton kernels declare runtime-varying, batch-derived values as tl.constexpr. Constexpr values are baked into the Triton JIT cache key, so each new value forces a full kernel recompile (hundreds of ms) at request time. This is the same bug class we recently root-caused in OpenAI's triton_kernels (triton-lang/triton#10872, fixes triton-lang/triton#10875/#10884), where it produced recurring ~1s TTFT stalls in gpt-oss serving — there we measured warm-phase p99 TTFT dropping 682ms → 316ms (= p50) from demoting two such params.
The instances below were found by a systematic audit of vLLM main (0762f2afeb74) and verified by tracing kernel param → launch site → producing expression. They have not been individually benchmarked; the mechanism and cost are established in the triton_kernels case.
Highest impact: MAX_MM_RANGES — unified attention recompiles per distinct per-batch image count
- Kernel param:
vllm/v1/attention/ops/triton_unified_attention.py:222 (consumed as an unrolled loop bound in triton_attention_helpers.py)
- Launch:
triton_unified_attention.py:1133 with max_mm_ranges = mm_prefix_range.shape[1]
- Producer:
vllm/v1/attention/backends/utils.py:69 — max_ranges = max(len(r) for r in range_lists), i.e. the per-batch max number of multimodal prefix ranges (≈ images), rebuilt every scheduler step
- Effect: on
is_mm_prefix_lm models (Gemma3-style) under TRITON_ATTN (default on ROCm), a request with a previously-unseen image count joining the batch recompiles the entire unified-attention kernel — one of the largest kernels in tree — mid-serving, per-value, unbounded. Same pattern in int4_per_token_head.py:360.
- Suggested fix: runtime loop bound (masked, non-unrolled) or pad to a small fixed set (e.g. pow2-bucket the range count).
Other confirmed instances (same pattern, smaller kernels)
prefill_tokens_with_context — vllm/v1/attention/ops/triton_merge_attn_states.py:71, launch :39 defaults to the per-batch prefill token count → effectively recompiles per prefill batch on the Triton path (ROCm/XPU/CUDA-fp8). The CUDA C++ twin correctly takes this as a runtime arg — the Triton port just needs the same treatment._pack_seq_kernel N/_unpack_seq_triton_kernel B — vllm/v1/attention/ops/common.py:265,391 — per-value (token/req counts) and unused in the kernel bodies; pure cache-key pollution on the DSA indexer non-uniform decode path. Lmax (batch max decode len) also keys per-value.NUM_KV_SPLITS (TritonMLA decode, triton_decode_attention.py:95,308,587) is already deliberately pow2-bucketed (triton_mla.py:41) — noting it here because the bucket count still ratchets ~10× across context growth on the decode hot path; clamping to a fixed small set (as turboquant does) would eliminate the residual recompiles.- Bounded-but-churning:
_count_expert_num_tokens BLOCK_SIZE = next_power_of_2(min(numel,1024)) (fused_moe/utils.py:49, trivial fix: constant 1024); LoRA-shrink SPLIT_K/BLOCK_K flip at batch<128 (lora/ops/triton_ops/utils.py:221); _fill_logprob_token_ids_kernel NUM_TOPK = per-batch max requested logprobs (v1/worker/gpu/sample/logprob.py:196).
Happy to split these into separate issues/PRs if maintainers prefer — the MAX_MM_RANGES and merge_attn_states ones look like straightforward demotions.
🤖 Found via systematic audit with Claude Code; pattern established and measured in triton-lang/triton#10872.
Summary
Several in-tree Triton kernels declare runtime-varying, batch-derived values as
tl.constexpr. Constexpr values are baked into the Triton JIT cache key, so each new value forces a full kernel recompile (hundreds of ms) at request time. This is the same bug class we recently root-caused in OpenAI'striton_kernels(triton-lang/triton#10872, fixes triton-lang/triton#10875/#10884), where it produced recurring ~1s TTFT stalls in gpt-oss serving — there we measured warm-phase p99 TTFT dropping 682ms → 316ms (= p50) from demoting two such params.The instances below were found by a systematic audit of vLLM main (
0762f2afeb74) and verified by tracing kernel param → launch site → producing expression. They have not been individually benchmarked; the mechanism and cost are established in the triton_kernels case.Highest impact:
MAX_MM_RANGES— unified attention recompiles per distinct per-batch image countvllm/v1/attention/ops/triton_unified_attention.py:222(consumed as an unrolled loop bound intriton_attention_helpers.py)triton_unified_attention.py:1133withmax_mm_ranges = mm_prefix_range.shape[1]vllm/v1/attention/backends/utils.py:69—max_ranges = max(len(r) for r in range_lists), i.e. the per-batch max number of multimodal prefix ranges (≈ images), rebuilt every scheduler stepis_mm_prefix_lmmodels (Gemma3-style) under TRITON_ATTN (default on ROCm), a request with a previously-unseen image count joining the batch recompiles the entire unified-attention kernel — one of the largest kernels in tree — mid-serving, per-value, unbounded. Same pattern inint4_per_token_head.py:360.Other confirmed instances (same pattern, smaller kernels)
prefill_tokens_with_context—vllm/v1/attention/ops/triton_merge_attn_states.py:71, launch:39defaults to the per-batch prefill token count → effectively recompiles per prefill batch on the Triton path (ROCm/XPU/CUDA-fp8). The CUDA C++ twin correctly takes this as a runtime arg — the Triton port just needs the same treatment._pack_seq_kernelN/_unpack_seq_triton_kernelB—vllm/v1/attention/ops/common.py:265,391— per-value (token/req counts) and unused in the kernel bodies; pure cache-key pollution on the DSA indexer non-uniform decode path.Lmax(batch max decode len) also keys per-value.NUM_KV_SPLITS(TritonMLA decode,triton_decode_attention.py:95,308,587) is already deliberately pow2-bucketed (triton_mla.py:41) — noting it here because the bucket count still ratchets ~10× across context growth on the decode hot path; clamping to a fixed small set (asturboquantdoes) would eliminate the residual recompiles._count_expert_num_tokensBLOCK_SIZE = next_power_of_2(min(numel,1024))(fused_moe/utils.py:49, trivial fix: constant 1024); LoRA-shrinkSPLIT_K/BLOCK_Kflip atbatch<128(lora/ops/triton_ops/utils.py:221);_fill_logprob_token_ids_kernelNUM_TOPK= per-batch max requested logprobs (v1/worker/gpu/sample/logprob.py:196).Happy to split these into separate issues/PRs if maintainers prefer — the
MAX_MM_RANGESandmerge_attn_statesones look like straightforward demotions.🤖 Found via systematic audit with Claude Code; pattern established and measured in triton-lang/triton#10872.