[Bugfix] Avoid shape-specialized Qwen3-VL pos-embed JIT

[lesj0610]

Purpose

_bilinear_pos_embed_kernel could still JIT during a real Qwen3-VL image request when the image grid differed from the shape used during profiling. The kernel does not annotate H, W, h_scale, or w_scale as tl.constexpr, but Triton still specializes scalar arguments unless they are explicitly excluded.

After review, this PR only marks the integer grid dimensions (H and W) as do_not_specialize. The floating-point scale arguments remain specialized because they participate directly in interpolation arithmetic and may affect generated-code performance.

This PR intentionally does not add a synthetic warmup. There is no model-specific dummy invocation to migrate to the shared warmup contract; normal multimodal profiling and kernel warmup routing remain unchanged.

Changes

• Set do_not_specialize=["H", "W"] on the fused bilinear position-embedding kernel.
• Preserve default specialization for h_scale and w_scale.
• Add unit coverage for the specialization contract and keep the existing output-correctness matrix.
• Merge the latest upstream main.

Test Plan

CUDA_VISIBLE_DEVICES=1 .venv/bin/python -m pytest \
  tests/kernels/core/test_vit_bilinear_pos_embed.py -q

pre-commit run --files \
  vllm/model_executor/models/qwen3_vl.py \
  tests/kernels/core/test_vit_bilinear_pos_embed.py

The kernel was also benchmarked on GPU 1 (RTX 3090) with 100 warmup iterations and 300 measured iterations for the baseline, the updated H/W-only policy, and the previous four-argument policy.

Test Result

• Qwen3-VL bilinear position-embedding tests: 19 passed.
• All applicable pre-commit hooks passed.
• Measured mean kernel latency in milliseconds:

dtype	grid	specialize all	do_not_specialize H/W	do_not_specialize all four
fp32	16x16	0.011936	0.011786	0.011661
fp32	48x48	0.032924	0.032880	0.032772
fp32	60x80	0.042403	0.042379	0.042174
bf16	16x16	0.007822	0.007659	0.007843
bf16	48x48	0.015873	0.016040	0.015872
bf16	60x80	0.022730	0.022691	0.022497

Against the fully specialized baseline, the H/W-only mean was lower in five cases and 1.1% higher for bf16 48x48. This matrix did not identify a consistent regression; the result is limited to the tested RTX 3090 shapes and dtypes.

AI assistance: Codex and Claude Fable 5 were used during implementation and PR preparation.

[LopezCastroRoberto]

Hey @lesj0610 Have you checked that using do_not_specialize does not affect performance? If those variables are part of the compile-key of _bilinear_pos_embed_kernel that specialization might be necessary. I would start by checking if (H, W, h_scale, w_scale) are defined as tl.constexpr in _bilinear_pos_embed_kernel.

I mention this because we are doing an effort to achieve zero JIT compilation during runtime, and would be great if you can migrate this warmup to the shared warmup contract we introduced in #47456

More info in: #49349, with some reference examples already merged. Thanks!

[claude]

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.