Skip to content

[ROCm][Kernel] Add HybridW4A16LinearKernel: Triton prefill + HIP skinny decode - #40977

Merged
vllm-bot merged 21 commits into
vllm-project:mainfrom
ROCm:matthias.hybrid-w4a16-kernel
Jul 14, 2026
Merged

vllm-bot merged 21 commits into
vllm-project:mainfrom
ROCm:matthias.hybrid-w4a16-kernel

Conversation

@mgehre-amd

@mgehre-amd mgehre-amd commented Apr 27, 2026

Copy link
Copy Markdown
Contributor

Add a hybrid W4A16 linear kernel for ROCm gfx11/gfx12 that routes between two GEMM implementations based on batch size:

  • M <= 5: HIP wvSplitK_int4_g skinny GEMM (optimized for single-token decode)
  • M > 5: Triton fused dequant GEMM (optimized for prefill/batched inference)

Both paths share a single weight tensor in the AWQ-style packed shuffle layout [N, K//8]. Supports both symmetric (uint4b8, bias=8) and asymmetric (uint4, per-group zero-points) quantization with group sizes 32, 64 or 128.

The Triton kernel includes tuned tile-size heuristics for gfx1151 (Strix Halo, RDNA 3.5) and gfx1201 (RDNA 4) based on real weight shapes from Qwen3-4B and Llama-3.1-8B-AWQ, plus shape-specific overrides.

The HIP skinny decode kernel uses a marlin-style magic-number dequantization trick for bf16 models, avoiding an expensive fp32 round-trip on gfx1151.

Benchmarked on AMD Strix Halo (gfx1151, LPDDR5X-8000 128 GB) with --max-num-seqs 1:

Model: RedHatAI/Qwen3-4B-quantized.w4a16 (compressed-tensors, symmetric)

Config Metric Baseline Kernel Speedup
in=128, out=128 TTFT 82.33 ms 68 ms 1.21x
in=128, out=128 TPOT 54.81 ms 13.69 ms 4.00x
in=1920, out=128 TTFT 1222.24 ms 760 ms 1.61x
in=1920, out=128 TPOT 56.16 ms 15.29 ms 3.67x

Model: Qwen/Qwen2.5-7B-Instruct-AWQ (AWQ, asymmetric)

Config Metric Baseline Kernel Speedup
in=128, out=128 TTFT 118.64 ms 100.09 ms 1.19x
in=128, out=128 TPOT 82.04 ms 21.03 ms 3.90x
in=1920, out=128 TTFT 2004.68 ms 1183.48 ms 1.69x
in=1920, out=128 TPOT 84.69 ms 21.67 ms 3.91x

Model: RedHatAI/gemma-3-4b-it-quantized.w4a16 (compressed-tensors, symmetric)

Config Metric Baseline Kernel Speedup
in=128, out=128 TTFT 85.18 ms 66 ms 1.29x
in=128, out=128 TPOT 53.39 ms 15.59 ms 3.43x
in=1920, out=128 TTFT 1392.00 ms 727 ms 1.91x
in=1920, out=128 TPOT 61.36 ms 24.24 ms 2.53x

Model: RedHatAI/Qwen3-8B-quantized.w4a16 (compressed-tensors, asymmetric)

Config Metric Baseline Kernel Speedup
in=128, out=128 TTFT 150.51 ms 140 ms 1.08x
in=128, out=128 TPOT 116.99 ms 24.55 ms 4.77x
in=1920, out=128 TTFT 3181.00 ms 1529 ms 2.08x
in=1920, out=128 TPOT 119.72 ms 26.12 ms 4.58x

Model: trymirai/SmolLM2-1.7B-Instruct-AWQ (AWQ, asymmetric)

Config Metric Baseline Kernel Speedup
in=128, out=128 TTFT 39.02 ms 46 ms 0.85x
in=128, out=128 TPOT 22.95 ms 7.25 ms 3.17x
in=1920, out=128 TTFT 769.37 ms 413 ms 1.86x
in=1920, out=128 TPOT 26.07 ms 10.53 ms 2.48x

Accuracy (lm_eval gsm8k, 200 samples):
RedHatAI/Qwen3-4B-quantized.w4a16: 0.835 -> 0.855 (no regression)

@mergify mergify Bot added ci/build performance Performance-related issues rocm Related to AMD ROCm labels Apr 27, 2026
@github-project-automation github-project-automation Bot moved this to Todo in AMD Apr 27, 2026

@gemini-code-assist gemini-code-assist Bot 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.

Code Review

This pull request introduces a Hybrid W4A16 kernel for ROCm, optimizing performance by dispatching to a HIP skinny GEMM kernel for decode and a Triton kernel for prefill using a shared ExLlama shuffle weight layout. A critical issue was identified in the ROCm C++ implementation where a hardcoded LDS size macro conflicts with dynamic host-side checks, potentially causing memory corruption on architectures with larger LDS (like MI300X/MI350) when the input size exceeds the static allocation.

Comment thread csrc/rocm/skinny_gemms_int4.cu Outdated

@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.

@mergify

mergify Bot commented May 4, 2026

Copy link
Copy Markdown
Contributor

Hi @mgehre-amd, the pre-commit checks have failed. Please run:

uv pip install pre-commit>=4.5.1
pre-commit install
pre-commit run --all-files

Then, commit the changes and push to your branch.

For future commits, pre-commit will run automatically on changed files before each commit.

Tip

Is mypy failing?
mypy is run differently in CI. If the failure is related to this check, please use the following command to run it locally:
# For mypy (substitute "3.10" with the failing version if needed)
pre-commit run --hook-stage manual mypy-3.10

@mgehre-amd
mgehre-amd force-pushed the matthias.hybrid-w4a16-kernel branch from 20867a2 to 258f0d8 Compare May 4, 2026 16:44
JartX added a commit to JartX/vllm that referenced this pull request May 5, 2026
Mirror Hybrid PR vllm-project#40977's union+pointer-cast data-flow opacity to defeat
clang ROCm 7.2 InstCombine fold (LLVM #76000) that previously dropped
v_dot2_f32_bf16 from our M=1 bf16 path despite multiple inline-asm and
builtin attempts:

  - Activations stored in a fp32-aliased union via uint32 writes
  - Weight magic values (nibble | 0x4300) written into another union
  - Inner dot uses *((bf16x2_t*)(&union.f[b])) pointer-cast to force
    bf16-typed reads, breaking the InstCombine fold chain
  - sum_a computed via a second v_dot2 with bf162(1.0, 1.0) constant,
    avoiding any fp32 widening of activations
  - Bias correction y_b_f * partial + z_b_f * sum_a unchanged

Verified via llvm-objdump --offloading + --disassemble-symbols on the
gemm_q4_kernel_rdna3<__hip_bfloat16, 1> instance: 32 v_dot2_f32_bf16
emits (was 0). VGPR=35.

Bench impact (Qwen3-4B b=1 in128/out128): 847 ms -> 828 ms (-2.2%).
Modest because M=1 decode is memory-bound on gfx1100 — the fp32 FMA
chain we replaced was already fast, v_dot2 saves only a few cycles
per int32 weight word. Closing the rest of the gap to Hybrid (~17%
on small models) requires architectural changes (YTILE-style multi-N
processing, DPP wave reduce instead of atomic CAS).

Signed-off-by: JartX <sagformas@epdcenter.es>
JartX added a commit to JartX/vllm that referenced this pull request May 5, 2026
Two complementary optimizations targeting the 27B b=4 in=2048 production
bench where Hybrid PR vllm-project#40977 was winning by 37%:

1. WMMA v4: 64M × 32N tile (was 64M × 16N in v3).

   Each wave now issues 2 wmmas per K-iter (one for cols 0-15, one for
   cols 16-31, sharing a_frag). With 4 waves × 2 wmmas = 8 wmmas in
   flight per K-iter, the wmma pipeline stays much closer to peak. LDS
   B-tile doubles to [2][16][32] = 2 KB. Dequant doubles too (64 slots
   per K-tile) and is distributed across waves 0,1 (waves 2,3 still
   idle on dequant — full-wave-distribution is a follow-up).

   Microbench (qkv-square K=N=4096) on RX 7900 XTX:
     M=2048: 670 K → 982 K tk/s    (1.47× over v3, 6.8× over v2)
   Larger shapes (gate/up, down, qwen-14B-up) hit ~1.5× over v3 too.

2. v_dot2_f32_bf16 in M_COUNT={2,4,8} bf16 decode path.

   Mirrors the M=1 trick: activations + magic-value weights stored in
   fp32-aliased pack4 union, dot via __builtin_amdgcn_fdot2_f32_bf16
   with bf16x2_t pointer-cast. sum_a[m] computed via second v_dot2
   with bf162(1,1). q_pack built once per col and reused across all
   M_COUNT rows, amortizing dequant cost.

End-to-end (Qwen3 27B-class W4A16 group_size=32, RX 7900 XTX):

  Bench                     v2-base   this      Hybrid   Δ vs Hybrid
  4B b=1 in128/out128       847 ms    817 ms    693 ms   -18 %
  27B b=1 in128/out128      3598 ms   3482 ms   3552 ms  +2 %  ← we win
  27B b=4 in2048/out128     26041 ms  17867 ms  16313 ms -9.5 %

The b=4 bench was the hardest: -31 % from baseline, gap to Hybrid
shrank 37 % → 9.5 %. The remaining 9.5 % is in shapes where N <= 32
(no v4 benefit, falls back to v3) and in decode multi-batch where
v_dot2_f32_bf16's half-rate keeps us at parity rather than ahead.

Dispatch chain (M-major / N-minor):
  M >= 64 && N >= 32 → v4 (4 waves, 64M × 32N, 2 wmma/wave/K-iter)
  M >= 64 && N <  32 → v3 (4 waves, 64M × 16N)
  32 <= M < 64       → v2 (2 waves, 32M × 16N)
  M < 32             → v1 (1 wave, 16M × 16N)

Signed-off-by: JartX <sagformas@epdcenter.es>
JartX added a commit to JartX/vllm that referenced this pull request May 5, 2026
Final WMMA tier for the prefill path: each wave issues 4 wmmas per
K-iter (cols 0-15, 16-31, 32-47, 48-63) sharing a_frag. 4 waves × 4
wmmas = 16 wmmas in flight per K-iter, fully saturating the wmma
pipeline (16-cycle latency × 1 issue/cycle).

LDS B-tile grows to 4 KB; dequant scales to 128 slots/K-tile (64 N ×
2 K-octets) and is distributed across all 4 waves with full lane
utilization (32 lanes × 4 waves = 128 dequants per K-iter, 4×
oversubscribed gives perfect 32-slot-per-wave mapping). Per-thread
VGPR ≈ 170 (a_frag 16 + 4×b_frag 64 + 4×c_acc 64 + locals) — under
the gfx1100 192-VGPR cap.

Microbench (qkv-square K=N=4096) on RX 7900 XTX:
  M=128:   753K → 980K  tk/s   (1.30× over v4)
  M=512:   820K → 1048K tk/s   (1.28×)
  M=1024:  894K → 1131K tk/s   (1.27×)
  M=2048:  982K → 1245K tk/s   (1.27×)

Total speedup vs the v2 baseline at M=2048 across Qwen-class shapes:
  qkv-square  (4096×4096):  144K → 1245K tk/s   (8.6×)
  gate/up     (4096×11008): 144K →  478K tk/s   (3.3×)
  qwen-14B-up (5120×13824): 150K →  296K tk/s   (2.0×)

End-to-end on Qwen3.6 27B-GPTQ-W4A16-G32 (RX 7900 XTX):

  Bench                     v2-base   this      Hybrid   Result
  4B b=1 in128/out128       847 ms    846 ms    693 ms   -22% (decode-bound)
  27B b=1 in128/out128      3598 ms   3439 ms   3552 ms  +3 % WIN
  27B b=4 in2048/out128     26041 ms  14738 ms  16313 ms +10 % WIN

The 27B b=4 in=2048 production bench is the headline result: shrunk
from 37 % behind Hybrid PR vllm-project#40977 to 10 % AHEAD of it (-43 % vs our
own v2 baseline).

Tried distributed-dequant across all 4 waves (16 slots/wave on the
v4 path): regressed 3 % on gate/up M=2048 because dequant is not on
the critical path under the double-buffer; spreading it just adds
LDS bank pressure. Reverted, kept v4's wave-0+1-only dequant.

Final dispatch chain (M-major / N-minor):
  M >= 64 && N >= 64 → v5 (4 waves, 64M × 64N, 4 wmma/wave/K-iter)
  M >= 64 && N >= 32 → v4 (4 waves, 64M × 32N, 2 wmma/wave/K-iter)
  M >= 64 && N <  32 → v3 (4 waves, 64M × 16N)
  32 <= M < 64       → v2 (2 waves, 32M × 16N)
  M < 32             → v1 (1 wave, 16M × 16N)

Signed-off-by: JartX <sagformas@epdcenter.es>
JartX added a commit to JartX/vllm that referenced this pull request May 5, 2026
Re-disasm of the v5 kernel after the fp32-internal dequant landed showed
16 NaN-canonicalisation VALU ops still in the K-loop body, all on the
__float2bfloat16 narrow at the end of the dequant chain. The pattern:

  v_or_b32_e32  v63, 0x400000, v54   // canonical NaN payload
  v_cmp_u_f32_e32 vcc_lo, v40, v40   // unordered compare (is NaN?)
  v_add3_u32 v68, v68, v34, 0x7FFF   // round-half-to-even bias
  v_cndmask_b32_e32                  // pick rounded value or canon NaN
  ds_store_b16_d16_hi                // write upper 16 bits

The NaN check is unreachable in this kernel: the dequant output is a
product of (nibble - zero) ∈ [-15, 15] (a small finite integer) and the
bf16 scale weight (a finite normal number), so the FMA result is always
finite. Add a `f32_to_bf16_no_canon` helper that emits just the rounding
+ truncate (`(fu + 0x7FFF + lsb) >> 16`) and use it for all 8 stores in
dequant_4bit_8_bf16_to_bf16.

Confirmed via re-disassembly: 32 fewer v_cmp_u_f32_e32 / v_cndmask_b32_e32
across the v5 kernel (158 → 132 → 100 cumulative across the two
dequant-path commits today).

Measured (gfx1100, RX 7900 XTX, --tp 1 --max-num-seqs 1 --max-model-len 4096):

  Qwen3.6-27B-GPTQ-W4A16-G32 b=1 in=128 out=128
                          before      after       Δ
    total                 3244 ms     3237 ms    -0.2%
    TTFT                  176.7 ms    173.1 ms   -2.0%
    TPOT                  24.15 ms    24.13 ms   -0.1%

  Qwen3.6-27B-GPTQ-W4A16-G32 b=1 in=1920 out=128
                          before      after       Δ
    total                 5427 ms     5383 ms    -0.8%
    TTFT                  2273 ms     2198 ms    -3.3%
    TPOT                  24.84 ms    25.07 ms   +0.9%   (within bench SD)

  Qwen3-4B-W4A16 b=1 in=1920 out=128
    total                 847.5 ms    844.4 ms   -0.4%

Cumulative TTFT improvement on Qwen3.6-27B in=1920 across today's three
W4A16 perf commits: 2454 ms (baseline) → 2198 ms (-10.4%); gap to
Hybrid PR vllm-project#40977 narrows from +49% to +33%.

Paris/4 generation test passes. No correctness regression observed.

Signed-off-by: JartX <sagformas@epdcenter.es>
JartX added a commit to JartX/vllm that referenced this pull request May 11, 2026
Mirror Hybrid PR vllm-project#40977's union+pointer-cast data-flow opacity to defeat
clang ROCm 7.2 InstCombine fold (LLVM #76000) that previously dropped
v_dot2_f32_bf16 from our M=1 bf16 path despite multiple inline-asm and
builtin attempts:

  - Activations stored in a fp32-aliased union via uint32 writes
  - Weight magic values (nibble | 0x4300) written into another union
  - Inner dot uses *((bf16x2_t*)(&union.f[b])) pointer-cast to force
    bf16-typed reads, breaking the InstCombine fold chain
  - sum_a computed via a second v_dot2 with bf162(1.0, 1.0) constant,
    avoiding any fp32 widening of activations
  - Bias correction y_b_f * partial + z_b_f * sum_a unchanged

Verified via llvm-objdump --offloading + --disassemble-symbols on the
gemm_q4_kernel_rdna3<__hip_bfloat16, 1> instance: 32 v_dot2_f32_bf16
emits (was 0). VGPR=35.

Bench impact (Qwen3-4B b=1 in128/out128): 847 ms -> 828 ms (-2.2%).
Modest because M=1 decode is memory-bound on gfx1100 — the fp32 FMA
chain we replaced was already fast, v_dot2 saves only a few cycles
per int32 weight word. Closing the rest of the gap to Hybrid (~17%
on small models) requires architectural changes (YTILE-style multi-N
processing, DPP wave reduce instead of atomic CAS).

Signed-off-by: JartX <sagformas@epdcenter.es>
JartX added a commit to JartX/vllm that referenced this pull request May 11, 2026
Two complementary optimizations targeting the 27B b=4 in=2048 production
bench where Hybrid PR vllm-project#40977 was winning by 37%:

1. WMMA v4: 64M × 32N tile (was 64M × 16N in v3).

   Each wave now issues 2 wmmas per K-iter (one for cols 0-15, one for
   cols 16-31, sharing a_frag). With 4 waves × 2 wmmas = 8 wmmas in
   flight per K-iter, the wmma pipeline stays much closer to peak. LDS
   B-tile doubles to [2][16][32] = 2 KB. Dequant doubles too (64 slots
   per K-tile) and is distributed across waves 0,1 (waves 2,3 still
   idle on dequant — full-wave-distribution is a follow-up).

   Microbench (qkv-square K=N=4096) on RX 7900 XTX:
     M=2048: 670 K → 982 K tk/s    (1.47× over v3, 6.8× over v2)
   Larger shapes (gate/up, down, qwen-14B-up) hit ~1.5× over v3 too.

2. v_dot2_f32_bf16 in M_COUNT={2,4,8} bf16 decode path.

   Mirrors the M=1 trick: activations + magic-value weights stored in
   fp32-aliased pack4 union, dot via __builtin_amdgcn_fdot2_f32_bf16
   with bf16x2_t pointer-cast. sum_a[m] computed via second v_dot2
   with bf162(1,1). q_pack built once per col and reused across all
   M_COUNT rows, amortizing dequant cost.

End-to-end (Qwen3 27B-class W4A16 group_size=32, RX 7900 XTX):

  Bench                     v2-base   this      Hybrid   Δ vs Hybrid
  4B b=1 in128/out128       847 ms    817 ms    693 ms   -18 %
  27B b=1 in128/out128      3598 ms   3482 ms   3552 ms  +2 %  ← we win
  27B b=4 in2048/out128     26041 ms  17867 ms  16313 ms -9.5 %

The b=4 bench was the hardest: -31 % from baseline, gap to Hybrid
shrank 37 % → 9.5 %. The remaining 9.5 % is in shapes where N <= 32
(no v4 benefit, falls back to v3) and in decode multi-batch where
v_dot2_f32_bf16's half-rate keeps us at parity rather than ahead.

Dispatch chain (M-major / N-minor):
  M >= 64 && N >= 32 → v4 (4 waves, 64M × 32N, 2 wmma/wave/K-iter)
  M >= 64 && N <  32 → v3 (4 waves, 64M × 16N)
  32 <= M < 64       → v2 (2 waves, 32M × 16N)
  M < 32             → v1 (1 wave, 16M × 16N)

Signed-off-by: JartX <sagformas@epdcenter.es>
JartX added a commit to JartX/vllm that referenced this pull request May 11, 2026
Final WMMA tier for the prefill path: each wave issues 4 wmmas per
K-iter (cols 0-15, 16-31, 32-47, 48-63) sharing a_frag. 4 waves × 4
wmmas = 16 wmmas in flight per K-iter, fully saturating the wmma
pipeline (16-cycle latency × 1 issue/cycle).

LDS B-tile grows to 4 KB; dequant scales to 128 slots/K-tile (64 N ×
2 K-octets) and is distributed across all 4 waves with full lane
utilization (32 lanes × 4 waves = 128 dequants per K-iter, 4×
oversubscribed gives perfect 32-slot-per-wave mapping). Per-thread
VGPR ≈ 170 (a_frag 16 + 4×b_frag 64 + 4×c_acc 64 + locals) — under
the gfx1100 192-VGPR cap.

Microbench (qkv-square K=N=4096) on RX 7900 XTX:
  M=128:   753K → 980K  tk/s   (1.30× over v4)
  M=512:   820K → 1048K tk/s   (1.28×)
  M=1024:  894K → 1131K tk/s   (1.27×)
  M=2048:  982K → 1245K tk/s   (1.27×)

Total speedup vs the v2 baseline at M=2048 across Qwen-class shapes:
  qkv-square  (4096×4096):  144K → 1245K tk/s   (8.6×)
  gate/up     (4096×11008): 144K →  478K tk/s   (3.3×)
  qwen-14B-up (5120×13824): 150K →  296K tk/s   (2.0×)

End-to-end on Qwen3.6 27B-GPTQ-W4A16-G32 (RX 7900 XTX):

  Bench                     v2-base   this      Hybrid   Result
  4B b=1 in128/out128       847 ms    846 ms    693 ms   -22% (decode-bound)
  27B b=1 in128/out128      3598 ms   3439 ms   3552 ms  +3 % WIN
  27B b=4 in2048/out128     26041 ms  14738 ms  16313 ms +10 % WIN

The 27B b=4 in=2048 production bench is the headline result: shrunk
from 37 % behind Hybrid PR vllm-project#40977 to 10 % AHEAD of it (-43 % vs our
own v2 baseline).

Tried distributed-dequant across all 4 waves (16 slots/wave on the
v4 path): regressed 3 % on gate/up M=2048 because dequant is not on
the critical path under the double-buffer; spreading it just adds
LDS bank pressure. Reverted, kept v4's wave-0+1-only dequant.

Final dispatch chain (M-major / N-minor):
  M >= 64 && N >= 64 → v5 (4 waves, 64M × 64N, 4 wmma/wave/K-iter)
  M >= 64 && N >= 32 → v4 (4 waves, 64M × 32N, 2 wmma/wave/K-iter)
  M >= 64 && N <  32 → v3 (4 waves, 64M × 16N)
  32 <= M < 64       → v2 (2 waves, 32M × 16N)
  M < 32             → v1 (1 wave, 16M × 16N)

Signed-off-by: JartX <sagformas@epdcenter.es>
JartX added a commit to JartX/vllm that referenced this pull request May 11, 2026
Re-disasm of the v5 kernel after the fp32-internal dequant landed showed
16 NaN-canonicalisation VALU ops still in the K-loop body, all on the
__float2bfloat16 narrow at the end of the dequant chain. The pattern:

  v_or_b32_e32  v63, 0x400000, v54   // canonical NaN payload
  v_cmp_u_f32_e32 vcc_lo, v40, v40   // unordered compare (is NaN?)
  v_add3_u32 v68, v68, v34, 0x7FFF   // round-half-to-even bias
  v_cndmask_b32_e32                  // pick rounded value or canon NaN
  ds_store_b16_d16_hi                // write upper 16 bits

The NaN check is unreachable in this kernel: the dequant output is a
product of (nibble - zero) ∈ [-15, 15] (a small finite integer) and the
bf16 scale weight (a finite normal number), so the FMA result is always
finite. Add a `f32_to_bf16_no_canon` helper that emits just the rounding
+ truncate (`(fu + 0x7FFF + lsb) >> 16`) and use it for all 8 stores in
dequant_4bit_8_bf16_to_bf16.

Confirmed via re-disassembly: 32 fewer v_cmp_u_f32_e32 / v_cndmask_b32_e32
across the v5 kernel (158 → 132 → 100 cumulative across the two
dequant-path commits today).

Measured (gfx1100, RX 7900 XTX, --tp 1 --max-num-seqs 1 --max-model-len 4096):

  Qwen3.6-27B-GPTQ-W4A16-G32 b=1 in=128 out=128
                          before      after       Δ
    total                 3244 ms     3237 ms    -0.2%
    TTFT                  176.7 ms    173.1 ms   -2.0%
    TPOT                  24.15 ms    24.13 ms   -0.1%

  Qwen3.6-27B-GPTQ-W4A16-G32 b=1 in=1920 out=128
                          before      after       Δ
    total                 5427 ms     5383 ms    -0.8%
    TTFT                  2273 ms     2198 ms    -3.3%
    TPOT                  24.84 ms    25.07 ms   +0.9%   (within bench SD)

  Qwen3-4B-W4A16 b=1 in=1920 out=128
    total                 847.5 ms    844.4 ms   -0.4%

Cumulative TTFT improvement on Qwen3.6-27B in=1920 across today's three
W4A16 perf commits: 2454 ms (baseline) → 2198 ms (-10.4%); gap to
Hybrid PR vllm-project#40977 narrows from +49% to +33%.

Paris/4 generation test passes. No correctness regression observed.

Signed-off-by: JartX <sagformas@epdcenter.es>
@BadrBasowid

Copy link
Copy Markdown
Contributor

@mgehre-amd some tests pass on the second try, if only a few tests are failing, tag @tjtanaa so he can help retry those specific tests.

@mgehre-amd

Copy link
Copy Markdown
Contributor Author

@AndreasKaratzas @tjtanaa, only three test failing. Do you want to rerun those or force merge?

@tjtanaa

tjtanaa commented Jul 8, 2026

Copy link
Copy Markdown
Member

@mgehre-amd I have rerun them. On of the failure is critical where CUDA CI didn't even run.

@mgehre-amd

mgehre-amd commented Jul 8, 2026

Copy link
Copy Markdown
Contributor Author

Failed this time:

tests/kernels/mamba/test_cpu_short_conv.py::test_short_conv_forward_native_prefill
FAILED lora/test_qwenvl.py::test_qwen2vl_multiple_lora_types - AssertionError: Generated text A view of the Tokyo doesn't match expected pattern A closeup shot of the Tokyo Skytree with pink flowers in the foreground.
FAILED lora/test_transformers_model.py::test_ilama_lora - RuntimeError: Engine core initialization failed. See root cause above. Failed core proc(s): {}

test_cpu_short_conv is a known problem on main and was fixed by #47848
The lora test failures are known problems on main and there is a PR to address them: #47798
None of the failed tests are due to this PR

plasticchris added a commit to plasticchris/vllm that referenced this pull request Jul 8, 2026
…1100

Native HIP skinny-decode GEMM (wvSplitK_int4_g) + Triton prefill for W4A16.
Supports compressed-tensors asymmetric uint4 (group_size 32/64/128), which the
symmetric-only RDNA3 kernel (PR vllm-project#41394) rejects. Registered ahead of
TritonW4A16LinearKernel in the ROCm mixed-precision priority list.

Adds on_gfx1151() to platforms/rocm.py (required by the hybrid prefill path).

Decode throughput on 2x RX 7900 XTX, Qwen3.6-27B W4A16-G32 asym, TP=2:
  ctx    50:  16.4 -> 28.1 tok/s (1.71x)
  ctx  8000:  14.5 -> 23.5 tok/s (1.62x)
  ctx 24000:  12.2 -> 18.6 tok/s (1.52x)
  ctx 48000:   9.9 -> 14.1 tok/s (1.42x)

Correctness: tests/kernels/quantization/test_rdna_hybrid_w4a16.py 83/83 pass
(forward parity sym+asym, all shapes/group sizes; process_weights repack).
@mgehre-amd

Copy link
Copy Markdown
Contributor Author

Failed on this rerun:

model_executor/model_loader/test_reload.py::test_reload_weights[inference-optimization/DeepSeek-V3-debug-empty-FP8_DYNAMIC-inference-optimization/DeepSeek-V3-debug-multiply-FP8_DYNAMIC-inference-optimization/DeepSeek-V3-debug-add-FP8_DYNAMIC-1]: AcceleratorError: CUDA error: operation not permitted when stream is capturing on MI300X. This PR only triggers for w4a16 models.
docker-build-arm64-image: denied: The repository with name 'vllm-ci-test-repo' in registry with id 'q9t5s3a7' already has the maximum allowed number of images which is '100000'
docker-build-cpu-image: denied: The repository with name 'vllm-ci-test-repo' in registry with id 'q9t5s3a7' already has the maximum allowed number of images which is '100000'
docker-build-image

@mgehre-amd

mgehre-amd commented Jul 10, 2026

Copy link
Copy Markdown
Contributor Author

After rebase, failed unrelated tests:

FAILED v1/kv_connector/extract_hidden_states_integration/test_extraction.py::test_extract_hidden_states_tp2 - Failed: Timeout >60.0s
FAILED model_executor/model_loader/test_reload.py::test_reload_weights[inference-optimization/DeepSeek-V3-debug-empty-FP8_DYNAMIC-inference-optimization/DeepSeek-V3-debug-multiply-FP8_DYNAMIC-inference-optimization/DeepSeek-V3-debug-add-FP8_DYNAMIC-1]

@mgehre-amd

mgehre-amd commented Jul 13, 2026

Copy link
Copy Markdown
Contributor Author

List of tests failed after this rebase:

FAILED v1/cudagraph/test_cudagraph_mode.py::test_cudagraph_compilation_combo[FA2-FULL_DECODE_ONLY-0-True] - ValueError: Memory of devices devices=[0] not free after dur_s=120.01 (threshold='device=0: 0.100')
FAILED renderers/test_hf.py::test_resolve_content_format_hf_defined[Qwen/Qwen3.5-4B-openai] - ValueError: No example model defined for Qwen/Qwen3.5-4B; please update this file.

First one is already known issue on main with a revert tracked in #47232
Second one is also unrelated to this PR.

@vllm-bot
vllm-bot merged commit 382bbd5 into vllm-project:main Jul 14, 2026
232 of 236 checks passed
@github-project-automation github-project-automation Bot moved this from Todo to Done in AMD Jul 14, 2026
@mgehre-amd
mgehre-amd deleted the matthias.hybrid-w4a16-kernel branch July 14, 2026 06:22
NickLucche pushed a commit to NickLucche/vllm that referenced this pull request Jul 15, 2026
…ny decode (vllm-project#40977)

Signed-off-by: Matthias Gehre <matthias.gehre@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
lcskrishna added a commit to lcskrishna/vllm that referenced this pull request Jul 15, 2026
* AGENTS MD: Add suggestion on how to incorporate tests (#47784)

Signed-off-by: Simon Mo <simon.mo@hey.com>
Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: Benjamin Chislett <chislett.ben@gmail.com>

* [BugFix][LoRA] Refresh punica metadata when LoRA slots are reassigned under an unchanged mapping (#47725)

Signed-off-by: AmeenP <ameenp360@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [Bug] Fix tmp directory for `lm_eval` (#47755)

Signed-off-by: yewentao256 <zhyanwentao@126.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [XPU] Fix topk_sigmoid arg mismatch on XPU (#47858)

Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com>

* [Doc] Fix grammatically incorrect error message in gpu_worker and xpu_worker (#36715)

Signed-off-by: Hongbin10 <jdmjdm1998@163.com>

* [Perf][3/N] Expand Triton kernel warmup coverage, Qwen (#47546)

Signed-off-by: LopezCastroRoberto <rocastro@redhat.com>

* Upgrade tpu-inference to v0.24.0 (#47835)

Signed-off-by: StingLin <sting.lin@cienet.com>

* [Bugfix] Exclude kv_cache_memory_bytes from CacheConfig.compute_hash (#47356)

Signed-off-by: Nils Matteson <nils@thaw.sh>
Signed-off-by: Nils Matteson <nilsmatteson@icloud.com>
Co-authored-by: Nils Matteson <nils@thaw.sh>

* fix: include topic frame in KV events replay response (#45177)

Signed-off-by: RishabhSaini <rishabhsaini01@gmail.com>

* [Bugfix] Fix mamba+dflash for MRV2 (#47698)

* [Bugfix] Allow non-contiguous query in FlashInfer FP8 query quantization (#47908)

Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [Doc] Fix manylinux tag in installation guide (#47913)

Signed-off-by: Nick Cao <ncao@redhat.com>

* [CI Bug Fix] Temp fix for v3.2 accuracy (#47902)

* [Bugfix] Pad Mamba page size instead of scaling block_size in unify_kv_cache_spec_page_size (#45207)

Signed-off-by: Sahil170595 <147995121+Sahil170595@users.noreply.github.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>

* [Doc] Surface the --kv-cache-memory suggestion at INFO and document fast-startup knobs (#47374)

Signed-off-by: Nils Matteson <nils@thaw.sh>
Signed-off-by: Nils Matteson <nilsmatteson@icloud.com>
Co-authored-by: Nils Matteson <nils@thaw.sh>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* Add tuned selective_state_update config for AMD Instinct MI355 (#47767)

Signed-off-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Co-authored-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>

* [BugFix] Fix ModelOpt quantization inference for fused siblings (#47445)

Signed-off-by: jasonlizhengjian <jasonlizhengjian@gmail.com>

* fix(security): bound completion prompt list to prevent unbounded engine fan-out (#47845)

Signed-off-by: jperezde <jperezde@redhat.com>

* [Bugfix] Patch Hopper MXFP4 OOB scales reads leading to NaN (#47910)

Signed-off-by: Yifan Zong <yzong@redhat.com>

* [CI/Build] Accept ready-run-all-tests label in pre-commit gate (#47897)

Signed-off-by: AmeenP <ameen@primeintellect.ai>
Co-authored-by: AmeenP <ameen@primeintellect.ai>
Co-authored-by: Claude <noreply@anthropic.com>

* [ROCm][CI] Minimize comment in RocmAttention q_scale check (#47731)

Signed-off-by: stefankoncarevic <stefan.koncarevic@amd.com>

* [ROCm][Bugfix] Key sparse-MLA persistent metadata on per-request context lengths (#47766)

Signed-off-by: Rohan Potdar <rohan.potdar@amd.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [Core] Persist and reuse the memory-profiling result across boots (opt-in) (#47388)

Signed-off-by: Nils Matteson <nils@thaw.sh>
Signed-off-by: Nils Matteson <nilsmatteson@icloud.com>
Co-authored-by: Nils Matteson <nils@thaw.sh>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [CPU][Bugfix] Fix flaky ShortConv prefill test on ARM (uninitialized weights) (#47848)

Signed-off-by: Rahul Vishwakarma <Rahul.Vishwakarma2@ibm.com>

* [Frontend] Add endpoint plugins framework (#47454)

Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com>

* [Bugfix] Avoid blocking model launching when no system ffmpeg available for TorchCodec (#47888)

Signed-off-by: Isotr0py <Isotr0py@outlook.com>

* [Core] Pass request context to CPU offload cache policy touch (#47744)

Signed-off-by: jacklin78911-collab <jacklin78911@gmail.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>

* [Core][Engine] only materialize tokens when thinking budget is in req (#47053)

Signed-off-by: walterbm <walter.beller.morales@gmail.com>

* [Bugfix][Pooling] Align CrossEncoder token type ids after truncation (#47772)

Signed-off-by: Ting Sun <suntcrick@gmail.com>

* [Minimax-M3] Using tok_sparse_select from MSA instead of triton kernels (#47502)

Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [ROCm][Perf] MXFP8 dense-linear + grouped-MoE GEMM optimizations for MiniMax-M3 (#46117)

Signed-off-by: amd-ethany <amd-ethany@users.noreply.github.com>
Co-authored-by: amd-ethany <amd-ethany@users.noreply.github.com>

* [Bugfix] DSV4 TP16 garbage output (#47493)

Signed-off-by: Jeff Ma <jeffjma@umich.edu>
Co-authored-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com>

* [Perf] Minimax M3 - Support cross-layer allreduce-norm fusion (#47631)

Signed-off-by: wzhao18 <wzhao18.sz@gmail.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>

* [Bugfix][Frontend] Fix http_requests_total metric recording some 4xx errors as 5xx (#44303)

Signed-off-by: Zach Zhu <zzqshu@126.com>

* [Bugfix] Register VLLM_BUILD_* and VLLM_IMAGE_TAG provenance env vars (#45313)

Signed-off-by: Nicklas Frahm <nicklas.frahm@gmail.com>

* [Bugfix] Avoid leaking Pydantic repr in tool_choice error message (#47028)

Signed-off-by: muhammadfawaz1 <135441198+muhammadfawaz1@users.noreply.github.com>
Co-authored-by: Mahad Durrani <114791389+mahadrehmann@users.noreply.github.com>

* [ROCm] fixed aiter master flag and expert parallelism compatibility on minimax-m3-mxfp8 (#47158)

Signed-off-by: Hongxia Yang <hongxia.yang@amd.com>

* fix: hash speculative draft model config (#47911)

Signed-off-by: Ace Eldeib <aeldeib@coreweave.com>
Signed-off-by: Ace Eldeib <alexeldeib@gmail.com>

* [Model][HunyuanVL] Use native transformers processor and adapt to transformers 5.13 (#47872)

Co-authored-by: manayang <manayang@tencent.com>

* [Docs] `kv_sharing_fast_prefill` correction (#47044)

Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai>

* Remove unused _get_kv_cache_config_deepseek_v4 alias (#47969)

Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai>

* [Test] Skip DeepEP MoE layer tests without P2P access (#47946)

Signed-off-by: Tyler Michael Smith <tyler@tylermsmith.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>

* [XPU] [Fusion passes] Disable fuse_rope_kvcache_cat_mla & qk_norm_rope_ fusion on XPU (#47962)

Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>

* [Misc][Docs]  Add human-readable integer support for more cli-args (#47608)

Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai>

* [Frontend] Support bad_words in the /v1/completions endpoint (#46793)

Signed-off-by: sungbin1015 <sbin@solbox.com>
Co-authored-by: Claude <noreply@anthropic.com>

* [ROCm] Add tuned selective_state_update float16 config for AMD Instinct MI300X (#47945)

Signed-off-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Co-authored-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>

* [Model] Support MOSS-Transcribe-Diarize (#47729)

Signed-off-by: gcanlin <canlinguosdu@gmail.com>

* [ROCm] Add tuned selective_state_update float32 config for AMD Instinct MI300X (#47947)

Signed-off-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Co-authored-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>

* Remove router weight upcast for DSv2-related models (#47970)

Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>

* [Bugfix] Allocate HY V3 expert_bias in float32 to prevent silent downcasting (#47797)

Signed-off-by: 辰言 <oncwnuIWp30GguOyJ615Fqj8H-yc@git.weixin.qq.com>
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
Co-authored-by: 辰言 <oncwnuIWp30GguOyJ615Fqj8H-yc@git.weixin.qq.com>
Co-authored-by: Jee Jee Li <jeejeelee@inferact.ai>
Co-authored-by: Jee Jee Li <pandaleefree@gmail.com>

* Add tuned selective_state_update float32 config for AMD Instinct MI355 (#47943)

Signed-off-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Co-authored-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>

* [Rust Frontend] Handle `continue_final_message` with renderer sentinel (#47844)

Signed-off-by: Bugen Zhao <i@bugenzhao.com>

* [CI/Build][BugFix][The Rock] Fix get_ssm_device_name to return sanitized, usable filename (#47781)

Signed-off-by: Randall Smith <Randall.Smith@amd.com>

* [Bugfix][V1] Free out-of-window blocks on the processed-token basis under async scheduling (#47728)

Signed-off-by: Saddss <28726669061@qq.com>
Signed-off-by: Nick Hill <nickhill123@gmail.com>
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai>
Co-authored-by: Saddss <28726669061@qq.com>
Co-authored-by: Nick Hill <nickhill123@gmail.com>

* Improvement of Docker image build for IBM Power using prebuilt wheels from IBM published devpi index (#46017)

Signed-off-by: vivek sharma <vivsharm@redhat.com>
Signed-off-by: puneetsharma21 <puneet.sharma21@ibm.com>
Signed-off-by: Puneet Sharma <puneet.sharma21@ibm.com>
Co-authored-by: vivek sharma <vivsharm@redhat.com>
Co-authored-by: Puneet Sharma <puneet.sharma21@ibm.com>
Co-authored-by: depthfirst-app[bot] <184448029+depthfirst-app[bot]@users.noreply.github.com>

* [Bug] Fix Batched DeepGEMM (#47884)

Signed-off-by: Robert Shaw <robertgshaw2-redhat@dgx-b200-02.mgmt.accl-001.lab.rdu2.dc.redhat.com>
Co-authored-by: Robert Shaw <robertgshaw2-redhat@dgx-b200-02.mgmt.accl-001.lab.rdu2.dc.redhat.com>
Co-authored-by: Claude <noreply@anthropic.com>

* Fix NVML capability lookup for visible devices (#47892)

Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>

* [Spec Decode] Support hybrid (SWA + full attention) DFlash drafters (#47914)

Signed-off-by: mgoin <mgoin64@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Fix FlashAttention MLA prefill V unpadding (#42642)

Signed-off-by: Martin Vit <martin@voipmonitor.org>
Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Matthew Bonanni <mbonanni@redhat.com>

* [Bugfix][DCP] Cast LSE to fp32 in a2a combine to fix bf16 bitcast crash (#47801)

Signed-off-by: Shawn Tsai <shawnyht@gmail.com>

* [Bugfix] Re-enable benchmarking of librispeech dataset. (#47033)

Signed-off-by: Anna Mayne <anna.mayne@arm.com>

* [CI] BugFix Eval Small Models Distributed test for DiffusionGemma (#47980)

Signed-off-by: Markov Ilya <markovilya19@gmail.com>
Co-authored-by: Markov Ilya <markovilya19@gmail.com>

* updated flash_attn GIT_TAG to point to torch Stable ABI FA3 commit (#47995)

Signed-off-by: Chris Leonard <chleonar@redhat.com>

* [Core] Move MRV1 `late_interaction_runner.py` out of MRV2 subtree (#48014)

Signed-off-by: Nick Hill <nickhill123@gmail.com>

* [Perf] Integrate TRTLLM BF16 MoE Modular Kernel  (#45182)

Signed-off-by: Kaihang Jiang <kaihangj@login-lyris02.lyris.clusters.nvidia.com>

* [ROCm][Bugfix] Fix empty-tensor .max() crash in AITER FA (#47894)

Signed-off-by: Djordje Ramic <djoramic@amd.com>

* Allow FlashInfer A2A backends for TRTLLM FP8 MoE Modular (#46661)

Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>

* [CI] GSM8K eval integration test for KV offloading (#46893)

Signed-off-by: Tyler Michael Smith <tyler@tylermsmith.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Signed-off-by: Itay Etelis <itay.etelis@ibm.com>
Signed-off-by: Itay Etelis <92247226+Etelis@users.noreply.github.com>
Signed-off-by: Or Ozeri <oro@il.ibm.com>
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-authored-by: Itay Etelis <92247226+Etelis@users.noreply.github.com>
Co-authored-by: Itay Etelis <itay.etelis@ibm.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [Feat] Add runtime monitor for post-warmup TileLang compilation (#46718)

Signed-off-by: LopezCastroRoberto <rocastro@redhat.com>
Signed-off-by: Roberto L. Castro <38211239+LopezCastroRoberto@users.noreply.github.com>

* [docs] Fix the docs build (#48008)

Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com>
Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>

* [ROCm] Fix pooling startup workspace lock (#47912)

Signed-off-by: Andreas Karatzas <akaratza@amd.com>

* Fix embed scaling + CUDA graphs in Transformers modelling backend (#48010)

Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>

* [Bugfix] Use int8 workspace for FlashInfer MLA decode (#48046)

Signed-off-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [KVConnector][NIXL] Support pipeline-parallel prefill in push mode (#45880)

Signed-off-by: zixi-qi <zixi@inferact.ai>
Co-authored-by: Nick Hill <nickhill123@gmail.com>

* [ROCm][CI][MoE] Fix double-transpose of fused w3 expert weights (#47874)

Signed-off-by: stefankoncarevic <stefan.koncarevic@amd.com>
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [Bugfix][ROCm] Change AttentionCGSuppoort in TritonMLA to UNIFORM_SINGLE_TOKEN_DECODE (#47144)

Signed-off-by: Dino Music <Dino.Music@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [XPU][LoRA] Fix torch.compile DEVICE_LOST by avoiding view-mutation in LoRA shrink (#47944)

Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>

* [KV Offloading] Add free block iterator for CPU offload scheduling (#47849)

Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com>

* [CPU] Fix Qwen-Next SSM type for AMX GDN (#48073)

Signed-off-by: jiang1.li <jiang1.li@intel.com>

* [Bugfix] Fix Qwen3-ASR transcription streaming postprocessing (#42478)

Signed-off-by: JooHo Lee <BWAAEEEK@users.noreply.github.com>
Signed-off-by: JooHo Lee <jooho414@gmail.com>
Co-authored-by: JooHo Lee <BWAAEEEK@users.noreply.github.com>

* Remove TeleChatForCausalLM  (#47989)

Signed-off-by: Xianbao QIAN <xianbao.qian@gmail.com>
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>

* Add Intel XPU Docker release pipeline (#47880)

Signed-off-by: wenjun.liu <wenjun.liu@intel.com>
Signed-off-by: jun,du <jun.du@intel.com>
Co-authored-by: jun,du <jun.du@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>

* [Bugfix] Fix race condition in KVBlockZeroer (#48085)

Signed-off-by: Benjamin Chislett <bchislett@nvidia.com>
Signed-off-by: Benjamin Chislett <chislett.ben@gmail.com>
Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>

* Sanitize server file paths from validation error responses (#46415)

Signed-off-by: muhammadfawaz1 <135441198+muhammadfawaz1@users.noreply.github.com>
Co-authored-by: Mahad Durrani <114791389+mahadrehmann@users.noreply.github.com>

* [P/D][Bugfix] Fix PD async KV load lookahead handling for MTP spec decode (#46694)

Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com>
Co-authored-by: Nick Hill <nickhill123@gmail.com>

* Remove PersimmonForCausalLM and FuyuForCausalLM model architectures (#48096)

Signed-off-by: Xianbao QIAN <xianbao.qian@gmail.com>
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Signed-off-by: Cyrus Leung <tlleungac@connect.ust.hk>
Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk>

* Migrate Olmo and Olmo2 to the Transformers modeling backend (#48100)

Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>

* [Bugfix][Model Runner V2] Order uniform decodes first so spec decodes aren't misclassified as prefills (#47381)

Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>
Signed-off-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Nick Hill <nickhill123@gmail.com>

* [CI] Annotate built Docker image tags on the Buildkite build page (#48101)

Signed-off-by: khluu <khluu000@gmail.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

* Pin PyNvVideoCodec to tested 2.0.4 wheel (#48056)

* [Bugfix] Guard CUDA-only rms_norm_per_block_quant in FUSED_OPS for non-CUDA builds (#47296)

Signed-off-by: Tsvika Shapira <tsvika@moonmath.ai>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [ROCm][CI] Set all timeout_in_minutes to 180 (#48146)

Signed-off-by: charlifu <charlifu@amd.com>

* [Bugfix] Preserve tensor causal metadata for grouped attention (#48135)

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Codex <codex@openai.com>

* [KVConnector] MultiConnector: give every sub-connector the request's real blocks in `update_state_after_alloc` (#46865)

Signed-off-by: deng451e <838677410@qq.com>

* [Bugfix][MRV2] Reset num_accepted_tokens on add_request in all modes (#48132)

Signed-off-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [ROCM][DSV32][Perf][MTP] Enable UNIFORM_BATCH CG mode in rocm_aiter_mla_sparse (#45149)

Signed-off-by: Teemu Virolainen <teemu.virolainen@amd.com>
Co-authored-by: TJian <tunjian.tan@embeddedllm.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [KV Connector][Mooncake] Apply SWA lookup mask before hashing/key build (#47317)

Signed-off-by: Zhewen Li <zhewenli@inferact.ai>
Co-authored-by: Zhewen Li <zhewenli@inferact.ai>
Co-authored-by: Codex <codex@openai.com>
Co-authored-by: Claude <noreply@anthropic.com>

* [ROCm] Synchronize sparse MLA metadata before graph replay (#47404)

Signed-off-by: zihaomu <zmu@amd.com>

* [CI] Increase extract hidden states TP2 timeout (#48161)

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>

* [ROCm] Revert Part of `[ROCm] Fix pooling startup workspace lock` #47912 (#48154)

Signed-off-by: Micah Williamson <micah.williamson@amd.com>

* [CI] Fix cargo-deny config flag ordering (#48170)

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>

* Correct model layer aliasing for Bert style models (#43896)

* update marlin M size for EP (#48144)

Signed-off-by: gnovack <novackgm@gmail.com>

* [ROCm][CI] Move remaining engine/samplers AMD steps to mi325_1 (#48169)

Signed-off-by: pei.zhang <pei.zhang@amd.com>
Co-authored-by: Claude <noreply@anthropic.com>

* [bugfix] bge-m3-sparse-plugin mismatch requests (#48112)

Signed-off-by: augusto.yjh <augusto.yjh@antgroup.com>

* [CI/Build][AMD] Fix ROCm OOM in eagle_correctness_heavy by reserving CUDA graph memory (#47366)

Signed-off-by: pei.zhang <pei.zhang@amd.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [kv_offload] Emit tier-owned BlockStored events from FS/OBJ secondary tiers (#47923)

Signed-off-by: Change72 <changg@nvidia.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>

* DCP supports hybrid attention (#40996)

Signed-off-by: YanXu <yancey.yx@alibaba-inc.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
Co-authored-by: Jingyi Yang <girasoleyang@gmail.com>

* [Core][KV events] Report prefix-cache-reused blocks in full report mode (#45261)

Signed-off-by: Lei Gong <gonglei25@huawei.com>
Co-authored-by: Lei Gong <gonglei25@huawei.com>
Co-authored-by: Claude <noreply@anthropic.com>

* [Feature][Parser] Support include_reasoning param for non-Harmony models (#44301)

Signed-off-by: Alberto Perdomo <aperdomo@redhat.com>
Co-authored-by: Chauncey <chaunceyjiang@gmail.com>

* [Perf] fuse more rmsnorm and all-reduce in qwen3.5 (#46998)

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>

* [XPU] Enable v1/sample tests on XPU CI (#44472)

Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>

* [CI] Right-size test-area timeouts from nightly durations (#48186)

Signed-off-by: khluu <khluu000@gmail.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

* Add XPU nightly and release image publishing to DockerHub (#48126)

Signed-off-by: wenjun.liu <wenjun.liu@intel.com>
Signed-off-by: jun,du <jun.du@intel.com>
Co-authored-by: jun,du <jun.du@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>

* [Rust Frontend] Integrate MM video support (#47959)

Signed-off-by: Bugen Zhao <i@bugenzhao.com>

* [Model][CI/Build] Cosmos3: enable registry tests and register Cosmos3-Super (#48211)

Signed-off-by: Mingfei Guo <1800012773@pku.edu.cn>

* [CI] Add TORCH_NIGHTLY=1 build mode (run full suite on torch nightly) (#47180)

Co-authored-by: Andrey Talman <atalman@users.noreply.github.com>
Co-authored-by: Kevin H. Luu <khluu000@gmail.com>

* Deepstream video backend (#42424)

Signed-off-by: Viranjan Pagar <vpagar@nvidia.com>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: Roger Wang <hey@rogerw.io>

* [Rust Frontend] Add roundtrip fixtures for more chat parsers (#47883)

Co-authored-by: Bugen Zhao <i@bugenzhao.com>
Signed-off-by: reidliu41 <reid201711@gmail.com>
Signed-off-by: Bugen Zhao <i@bugenzhao.com>

* [Misc] Remove dead code in ViT functionality test (#48220)

Signed-off-by: Isotr0py <Isotr0py@outlook.com>

* [Bugfix][Spec Decode] Fix DFlash draft/target layer-count mismatch (#48113)

Signed-off-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [Model] Migrate MistralLarge3ForCausalLM to AutoWeightsLoader (#48153)

Signed-off-by: Yuchen Fan <functionhx@gmail.com>

* [Refactor] Remove unused rocm kernel `combine_topk_swa_indices_ragged` (#48158)

Signed-off-by: yewentao256 <zhyanwentao@126.com>

* [Bugfix] Fix turboquant FP8 cast failure for BF16 models on Ampere GPUs (#39988)

Signed-off-by: Xu Zhou <xuzhou9417@163.com>
Co-authored-by: Xu Zhou <xuzhou9417@163.com>
Co-authored-by: Hoseung Kim <ghyutjik123@gmail.com>

* fix: correct load_weights track logic and enable weight integrity for… (#41811)

Signed-off-by: Yipeng Hu <i26268@metax-tech.com>
Signed-off-by: HuYiPeng <144002351+MynameFelix@users.noreply.github.com>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: Yipeng Hu <i26268@metax-tech.com>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>

* [Build/CI] Build arm64 PR and postmerge image builds for Blackwell SM10x and SM110 (#48041)

Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>

* [Model] Add LongCat-Flash-Lite (n-gram embedding) (#47857)

Signed-off-by: mgoin <mgoin64@gmail.com>

* [ROCm] Enable DeepSeek-V4 DSpark speculative decoding on AMD (MI350X / MI355X, gfx950) (#47419)

Signed-off-by: larryli2-amd <larryli2@amd.com>
Signed-off-by: larryli2-amd <Larry.Li@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [Bugfix] Fix FlashMLA dense fp8 metadata crash (num_sm_parts clamp) (#48045)

Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Claude <noreply@anthropic.com>

* handle topk_ids padding in align sum kernel (#47785)

Signed-off-by: gnovack <novackgm@gmail.com>

* [Bugfix][Test] Register Qwen/Qwen3.5-4B example model (#48276)

Signed-off-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [Bugfix] Fix thinking_token_budget not enforced after natural </think> re-entry (#45984)

Signed-off-by: Ashwin Giridharan <girida@amazon.com>

* Add VLLM_FLASHINFER_AUTOTUNE_SKIP_OPS and skip CuTeDSL fp4_gemm autotuning by default (#48268)

Signed-off-by: mgoin <mgoin64@gmail.com>
Signed-off-by: Michael Goin <mgoin64@gmail.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>

* [Quantization] Bound peak memory when repacking FP4 MoE weights for Marlin (#47851)

Signed-off-by: Joe Rowell <joerowell4@gmail.com>
Signed-off-by: mgoin <mgoin64@gmail.com>
Co-authored-by: mgoin <mgoin64@gmail.com>

* [BugFix] weights processing peak memory reduction for nvfp4 MoE layers (#46276)

Signed-off-by: Jimmy Lee <hirejimmylee@gmail.com>

* [BugFix] Fix packed HND KV cache reshape for FlashAttention (#47314)

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [Misc] Use meta tensor for KV cache stride calculation (#47316)

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [Logs] DP Supervisor Log Improvement (#48278)

Signed-off-by: Robert Shaw <robertgshaw2-redhat@ip-172-31-18-125.us-east-2.compute.internal>
Co-authored-by: Robert Shaw <robertgshaw2-redhat@ip-172-31-18-125.us-east-2.compute.internal>

* [Revert] [Build] Update vllm ...builds FA3 with torch stable API (#48269)

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>

* Bump Transformers version to 5.13.0 (#47867)

Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>

* [XPU]remove is_xxx from moe class and bump up kernels (#48079)

Signed-off-by: mayuyuace <qiming1.zhang@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>

* [CI] Point CI at Transformers release rather than release branch (#48328)

Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>

* FP32 router GEMV optimization (#48335)

Signed-off-by: peiyuanz <peiyuanz@inferact.ai>
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
Co-authored-by: peiyuanz <peiyuanz@inferact.ai>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: zhouzhou <zhouzhou@zhouzhoudeMacBook-Pro.local>

* [XPU][UT]Fix InternS1ProForConditionalGeneration AssertionError (#48232)

Signed-off-by: Lai, Yejing <yejing.lai@intel.com>

* [2/N][KV-Cache Layout Refactor] Pack K/V into the content dim across attention backends (#44455)

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Signed-off-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
Co-authored-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Nicolò Lucchesi <nlucches@redhat.com>

* fix(entrypoints): stop resolve_items leaking in-flight media fetch tasks on partial failure (#48333)

Signed-off-by: ErenAta16 <erena6466@gmail.com>

* fix(processor): route MiMo-V2-Omni media fetch through MediaConnector (#43117)

Signed-off-by: Ievgen Bondarenko <ibondarenko@student.sierracollege.edu>
Signed-off-by: Ievgen (Jack) Bondarenko <ibondarenko@student.sierracollege.edu>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>

* [CI] split tests to reduce CI time (#48219)

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>

* [Bugfix][LoRA] Support ark_linear base layer in _get_lora_device (#47690)

Signed-off-by: AlejandroParedesLT <alejandroparedeslatorre@gmail.com>

* [CI][CPU] Add Qwen2-VL multimodal tests for CPU backend and fix incompatibilities (#48072)

Signed-off-by: Zhenhui Zhao <zhenhui.zhao@intel.com>

* [2/N][Core] support partial prefix cache hit for hybrid model (#46384)

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai>

* Runtime Draft Weight Update for Speculative Decoding (#46725)

Signed-off-by: vx120 <893600387@qq.com>
Signed-off-by: vx120 <57470515+vx120@users.noreply.github.com>
Signed-off-by: aoshen02 <aoshen@inferact.ai>
Co-authored-by: crp0128 <191679376@qq.com>
Co-authored-by: aoshen02 <aoshen@inferact.ai>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

* [Perf][Qwen] Replace MOE all-reduce with reduce-scatter (#47006)

Signed-off-by: gcanlin <canlinguosdu@gmail.com>
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Co-authored-by: yewentao256 <zhyanwentao@126.com>

* [Frontend] Add /abort_requests to the RLHF dev API router (#47173)

Signed-off-by: aoshen02 <aoshen@inferact.ai>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [Bugfix] Guard mixed-dtype allreduce RMSNorm quant fusions (#48330)

Signed-off-by: hcenteno <hugo.centeno@estudiantat.upc.edu>

* Make tiering offload region DP-replica aware (#47987)

Signed-off-by: Liran Schour <lirans@il.ibm.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>

* [CI][2/N] reduce CI time (#48394)

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>

* [EC Connector] Add EC Transfer Params (#42433)

Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>

* [CI][Rust Frontend] Pin cargo tool versions (#48222)

Signed-off-by: Bugen Zhao <i@bugenzhao.com>

* [Bugfix][KV Offloading] Fix stale transfer_jobs after reset_cache + harden job completion (#48102)

Signed-off-by: Alex <alex.tech.lab@outlook.com>

* [Kernel] Implement CUDA kernel for ReLUSquaredActivation (relu^2) (#39058)

Signed-off-by: Tanish Malekar <tanishmalekar32@gmail.com>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>

* [ROCm][MiniMax-M3] Add AITER sparse paged attention (#47287)

Signed-off-by: Tan Pin Siang <tanpinsiang@gmail.com>
Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com>
Co-authored-by: vllmellm <vllm.ellm@embeddedllm.com>
Co-authored-by: Hongxia Yang <hongxia.yang@amd.com>
Co-authored-by: Jun Kang Chow <junkangchow@gmail.com>
Co-authored-by: tjtanaa <tunjian.tan@embeddedllm.com>

* [Model] Optimize Qwen3.5 on H20 (#48350)

Signed-off-by: zzt <zengzetang.zzt@antgroup.com>

* [Bugfix][ROCm] Keep TP all_gather on base-class collective (#48446)

Signed-off-by: fai <fangzhouai@gmail.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

* Re-disable CUDA graph memory profiling on ROCm (#48440)

Signed-off-by: Rohan Potdar <rohan.potdar@amd.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* [CPU][Spec Decode] Support DFlash speculative decoding for GDN models on CPU (#46090)

Signed-off-by: guybd <guy.boudoukh@intel.com>
Signed-off-by: Guy Boudoukh <guy.boudoukh@intel.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>

* [Bugfix][UT]Fix EagleMiniCPMForCausalLM meet TypeError (#48452)

Signed-off-by: Lai, Yejing <yejing.lai@intel.com>

* [CPU] Create Proper Numa topology for s390x (#40714)

Signed-off-by: Rehan Khan <Rehan.Khan7@ibm.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>

* [BugFix] Restore full tokens for Qwen MTP When MoE SP (#48429)

Signed-off-by: gcanlin <canlinguosdu@gmail.com>

* [ROCm][CI] Cache Rust builds by source inputs (#46527)

Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>
Co-authored-by: Matthew Wong <Matthew.Wong2@amd.com>

* [Distributed][Perf] Enable FlashInfer MNNVL allreduce RMS quant fusion (#48064)

Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>

* [Attention] Make sliding-window support an explicit backend capability (#48011)

Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [Bugfix][KV Cache] Don't route uniform-page-size MLA+SWA models into DeepseekV4 packing (#48256)

Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai>
Co-authored-by: Claude <noreply@anthropic.com>

* [Core] Support fp32 lm_head for generation models via head_dtype (RFC #48305 §3.6) (#48390)

Signed-off-by: Karthik Kothuri <karthikkothuri2009@gmail.com>
Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: wang.yuqi <yuqi.wang@daocloud.io>

* [CI] Add SPDX license header to Rust/Protobuf sources (#48472)

Signed-off-by: Bugen Zhao <i@bugenzhao.com>

* [Mypy Fix] Split mypy work (#48490)

Signed-off-by: yewentao256 <zhyanwentao@126.com>

* lower memory required for capturing cudagraphs for large cudagraph sizes (#48483)

Signed-off-by: Omer Ullman Argov <118735753+omera-nv@users.noreply.github.com>

* remove force channels_last in Idefics3MultiModalProcessor (#48467)

Signed-off-by: Yan Ma <yan.ma@intel.com>

* [XPU][CI]Adjust timeout_in_minutes in Intel GPU CI (#48418)

Signed-off-by: zengxian <xiangdong.zeng@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>

* [BugFix][ModelRunner V2] Fix stale attn metadata in speculator prefill cudagraph capture (#48261)

Signed-off-by: Nick Hill <nickhill123@gmail.com>

* [Misc]  Improve Matryoshka pooling dimensions validation (#48057)

Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [EC Connector] CPU Offloading EC Connector (#47423)

Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
Signed-off-by: Or Ozeri <oro@il.ibm.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>

* [ROCm][CI] Transformers: pass only one of input_ids/inputs_embeds (#48258)

Signed-off-by: Stefan Koncarevic <stefan.koncarevic@amd.com>
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>

* [Doc] Add DeepseekV32ForCausalLM to supported_models.md (#48293)

Signed-off-by: Gavin Morris <gmorriscs@gmail.com>

* [Bugfix][Frontend] Flush engine reasoning parser at engine-reasoning → tool streaming boundary (#47606)

Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com>
Co-authored-by: Ben Browning <56071+bbrowning@users.noreply.github.com>

* [ROCm][CI] Unblock `AMD: Language Models Test (Extended Pooling)` (#48513)

Signed-off-by: Micah Williamson <micah.williamson@amd.com>

* Add DCP + Eagle support for Tokenspeed MLA backends (#48180)

Signed-off-by: Pavani Majety <pmajety@nvidia.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
Co-authored-by: Jingyi Yang <girasoleyang@gmail.com>

* Added sliding window attention support for qwen-eagle3 architecture (#47568)

Signed-off-by: shanjiaz <zsjwpianpian@gmail.com>

* [Core] Preserve Marconi caching with selective hybrid cache retention (#47782)

Signed-off-by: Nick Hill <nickhill123@gmail.com>

* [BugFix] Initialize model_config for Qwen3-VL MoE (#44863)

Signed-off-by: wenpengw-nv <wenpengw@nvidia.com>
Co-authored-by: Roger Wang <hey@rogerw.io>

* [Feat] Add Support for BertForMaskedLM to vLLM (#48463)

Signed-off-by: atalhens <sneh.lata@nutanix.com>

* [Frontend] Expose logprob_token_ids on Python OpenAI endpoints (#43463)

Signed-off-by: Lang Zhao <lang.zhao@galileo.ai>
Co-authored-by: Claude <noreply@anthropic.com>

* [BugFix] Correct OTEL span start time for Dynamo compilation (#40698)

Signed-off-by: emricksini-h <emrick.birivoutin@hcompany.ai>
Co-authored-by: Simon Mo <simon.mo@hey.com>

* add pad-aware swiglu limit kernel (#48287)

Signed-off-by: gnovack <novackgm@gmail.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [Quantization][INC][ARK] Support INT2 XPU WOQ Linear (#47521)

Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com>
Signed-off-by: Zhenzhong Xu <zhenzhong.xu@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>

* [1/N] Add dense MHA path for sparse MLA short sequences (#47327)

Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

* up FI fp8 moe topk to 32 (#44462)

* [ROCm][MiniMax-M2] Dispatch fused QK-norm + AllReduce via AITER (#44849)

Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com>
Co-authored-by: Pawel Kowalski <pawel.kowalski@amd.com>

* [ROCm][Kernel] Add HybridW4A16LinearKernel: Triton prefill + HIP skinny decode (#40977)

Signed-off-by: Matthias Gehre <matthias.gehre@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [Perf][ROCm] Fix GDN KKT warmup regression on RDNA by avoiding fp32 tl.dot (#45000)

Signed-off-by: Saeid Rostami <srostami@amd.com>
Signed-off-by: nemanjaudovic <nudovic@amd.com>

* [ROCm][MiniMax-M3][Spec Decode] Support speculative decode with AITER sparse PA (#47984)

Signed-off-by: Tan Pin Siang <tanpinsiang@gmail.com>
Co-authored-by: vllmellm <vllm.ellm@embeddedllm.com>
Co-authored-by: Hongxia Yang <hongxia.yang@amd.com>
Co-authored-by: Jun Kang Chow <junkangchow@gmail.com>
Co-authored-by: TJian <tunjian.tan@embeddedllm.com>

* [Bugfix] Skip minimax_m3 tool parser tests when Rust extension is absent (#48523)

Signed-off-by: Matt Woodson <mwoodson@redhat.com>

* [Misc] Clean up "swap_space" (#48549)

Signed-off-by: wangxiyuan <wangxiyuan1007@gmail.com>

* [ROCm] Retune MI355 selective_state_update float16 config on the unified effective_batch grid (#48372)

Signed-off-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Co-authored-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>

* [Perf][Feat] Add generic cuteDSL LL BF16 router (GEMM) (#42562)

Signed-off-by: LopezCastroRoberto <rocastro@redhat.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Signed-off-by: Roberto L. Castro <38211239+LopezCastroRoberto@users.noreply.github.com>
Co-authored-by: Lucas Wilkinson <lwilkins@redhat.com>

* fix(security): guard lm-format-enforcer regex compile with timeout (#47595)

Signed-off-by: jperezde <jperezde@redhat.com>

* Add Cosmos3 Edge Reasoner model (#48291)

Signed-off-by: Bartosz Stefaniak <bstefaniak@nvidia.com>
Co-authored-by: Bartosz Stefaniak <bstefaniak@nvidia.com>

* [Bugfix] Return 400 instead of 500 when multimodal data is sent to a text-only model (#48473)

Signed-off-by: Hoang Nguyen Tien <hoang.nguyentien.2601@gmail.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

* [XPU][CI] Add tests/v1/e2e/general/test_correctness_sliding_window.py in Intel GPU CI (#47231)

Signed-off-by: zengxian <xiangdong.zeng@intel.com>
Signed-off-by: xiangdong <40376367+zxd1997066@users.noreply.github.com>

* [Test] Enable KV cache events for HMA models in CPU offloading test (#47754)

Signed-off-by: Itay Etelis <92247226+Etelis@users.noreply.github.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>

* [Misc] Rename VLLM_TRITON_ATTN_USE_TD to VLLM_TRITON_USE_TD (#45781)

Signed-off-by: Artur Fierka <artur.fierka@intel.com>
Co-authored-by: Claude <noreply@anthropic.com>

* [Bugfix][Security] Fix concurrent sparse invariant race bypassing CVE remediation (#48583)

Signed-off-by: jperezde <jperezde@redhat.com>

* [Model] Enable LoRA support for tower and connector in LlavaNextVideo (#48594)

Signed-off-by: gangula-karthik <gkarthik923@gmail.com>

* [NIXL] Avoid reading expired blocks in bidirectional turn-2 read (#47021)

Signed-off-by: Tomer Gilad <tgilad@nvidia.com>
Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai>
Co-authored-by: NickLucche <nicolo.lucchesi@mistral.ai>

* [Bugfix] Include inline per-token-head scales in offloaded page transfer width (#48411)

Signed-off-by: Itay Etelis <itay.etelis@ibm.com>
Signed-off-by: Itay Etelis <Itay.etelis@gmail.com>
Signed-off-by: Itay Etelis <92247226+Etelis@users.noreply.github.com>
Co-authored-by: Itay Etelis <itay.etelis@ibm.com>
Co-authored-by: Itay Etelis <Itay.etelis@gmail.com>

* [Bugfix] Gemma4 parser: classify channel-less output consistently in streaming and non-streaming (#48262)

Signed-off-by: Adhithya Balakrishnan <adhithya.b2004@gmail.com>
Co-authored-by: Ben Browning <56071+bbrowning@users.noreply.github.com>

* [CI Bug] Fully solve accuracy issue for DSv3.2 + MTP + Sequence Parallel (#48036)

Signed-off-by: yewentao256 <zhyanwentao@126.com>

* [Bugfix] Make MLA+SWA check the layer's backend, not the model config (#48520)

Signed-off-by: mgoin <mgoin64@gmail.com>

* [Doc] Sync four function docstrings with their signatures (#45437)

Signed-off-by: Daoyuan Li <94409450+DaoyuanLi2816@users.noreply.github.com>
Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [KV Offload] Split cpu_cache_usage_perc into write/read usage gauges (#47666)

Signed-off-by: srinivas_oo7 <sklinkedin0120@gmail.com>
Co-authored-by: srinivas_oo7 <sklinkedin0120@gmail.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>

* [ROCm] Retune MI355 selective_state_update float32 config on the unified effective_batch grid (#48373)

Signed-off-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Co-authored-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>

* [Reasoning] Optimize TPOT for thinking budget when used with speculative decoding (#46662)

Signed-off-by: rishitdholakia13 <rishit+github@cohere.com>
Co-authored-by: Cursor <cursoragent@cursor.com>

* [CI] Build macOS arm64 CPU wheel natively on the macmini queue (#48289)

Signed-off-by: mgoin <mgoin64@gmail.com>

* [ROCm][CI] fix flashinfer import check (#48647)

Signed-off-by: Divakar Verma <divakar.verma@amd.com>

* [CI][Bugfix] Fix FlashAttention reported MLA dimension support (#48631)

Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>

* Log fully resolved pooling config at startup (#48030)

Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* [Bugfix] MoRIIO toy P/D proxy: add /health (#45222)

Signed-off-by: Chaemin Lim <chaemin.lim@mangoboost.io>
Signed-off-by: Edwin Lim <edwinlim0919@gmail.com>
Co-authored-by: Edwin Lim <edwin.lim@mangoboost.io>
Co-authored-by: Jaeyoun Kim <jaeyoun.kim@mangoboost.io>
Co-authored-by: Edwin Lim <edwinlim0919@gmail.com>

* [Bugfix][CI] Fix test_head_dtype quant_method test on ROCm (#48654)

Signed-off-by: Micah Williamson <micah.williamson@amd.com>

* fix: size FlashInfer prefill workspace to batch head footprint (#48428)

Signed-off-by: Joe Rowell <joerowell4@gmail.com>

* [Bugfix][R3] Exclude draft routers from expert capture (#48622)

Signed-off-by: aoshen02 <aoshen@inferact.ai>

* [Bugfix] Preserve unloaded non-persistent buffers during layerwise reload (#44371)

Signed-off-by: Joan Velja <joan.velja22@gmail.com>
Co-authored-by: Dakai An <77474977+andakai@users.noreply.github.com>

* [Perf] Remove redundant repeat and copy for dsv4, 1.8% E2E TPOT improvement. (#48137)

Signed-off-by: yewentao256 <zhyanwentao@126.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

* add pad-aware reduce path (#48385)

Signed-off-by: gnovack <novackgm@gmail.com>

* [ROCm][CI] Remove mxfp4 test skips after `amd-quark` 0.12 release (#47330)

Signed-off-by: Micah Williamson <micah.williamson@amd.com>
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com>
Co-authored-by: Andreas Karatzas <Andreas.Karatzas@amd.com>
Co-authored-by: fxmarty-amd <felmarty@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [Core][LoRA] Support fp32 lm_head (head_dtype) on the LoRA path (#48525)

Signed-off-by: Karthik Kothuri <karthikkothuri2009@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>

* [CI][AMD] Configure MI300 tests for native execution without DinD (#48387)

Signed-off-by: Andreas Karatzas <akaratza@amd.com>

* [Quant] Enable humming w[2-7]a[4,8] inference with compressed-tensors (#46390)

Signed-off-by: HDCharles <charlesdavidhernandez@gmail.com>

* [LoRA][1/N] Integrate flashinfer MoE LoRA for BF16 model (#48632)

Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>

* [Security] Replace diskcache to eliminate pickle deserialization (#44549)

Signed-off-by: Russell Bryant <rbryant@redhat.com>
Co-authored-by: Claude <noreply@anthropic.com>

* Build with ABI stable FlashMLA (#48174)

Signed-off-by: Jane Xu <janeyx@meta.com>
Signed-off-by: Shengqi Chen <i@harrychen.xyz>
Co-authored-by: Shengqi Chen <i@harrychen.xyz>

* [Bugfix] Set kv_quant_mode on the generic MLA KV-cache spec (#48379)

Signed-off-by: Mikhail Kostryukov <mike@triptrack.net>
Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Matthew Bonanni <mbonanni@redhat.com>

* [Bugfix][Spec Decode] Support heterogeneous QK fusion geometry (#48671)

Signed-off-by: aoshen02 <aoshen@inferact.ai>
Co-authored-by: OpenAI Codex <codex@openai.com>

* [ROCm] Run init test engine in-process to avoid KV-cache OOM (#48527)

Signed-off-by: Djordje Ramic <djoramic@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [ROCm][CI] fix test_common.py (#48676)

Signed-off-by: charlifu <charlifu@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [Rust Frontend] Integrate MM audio support (#48554)

Co-authored-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>

* [CI/Build] Split release artifact annotations by type (#48600)

Signed-off-by: khluu <khluu000@gmail.com>
Co-authored-by: OpenAI Codex <noreply@openai.com>

* [ROCm] Add tuned selective_state_update config for AMD MI350 (#48159)

Signed-off-by: Giuseppe Grossi <ggrossi@amd.com>

* fix flaky multi example connector consistency (#48206)

Signed-off-by: aarushjain29 <aarushi.jain2@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>

* [ROCm][Bugfix] Enable the fp32 head_dtype torch.mm fast path on ROCm (#48688)

Signed-off-by: Turner Jabbour <doubleujabbour@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>

---------

Signed-off-by: Simon Mo <simon.mo@hey.com>
Signed-off-by: AmeenP <ameenp360@gmail.com>
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com>
Signed-off-by: Hongbin10 <jdmjdm1998@163.com>
Signed-off-by: LopezCastroRoberto <rocastro@redhat.com>
Signed-off-by: StingLin <sting.lin@cienet.com>
Signed-off-by: Nils Matteson <nils@thaw.sh>
Signed-off-by: Nils Matteson <nilsmatteson@icloud.com>
Signed-off-by: RishabhSaini <rishabhsaini01@gmail.com>
Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
Signed-off-by: Nick Cao <ncao@redhat.com>
Signed-off-by: Sahil170595 <147995121+Sahil170595@users.noreply.github.com>
Signed-off-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Signed-off-by: jasonlizhengjian <jasonlizhengjian@gmail.com>
Signed-off-by: jperezde <jperezde@redhat.com>
Signed-off-by: Yifan Zong <yzong@redhat.com>
Signed-off-by: AmeenP <ameen@primeintellect.ai>
Signed-off-by: stefankoncarevic <stefan.koncarevic@amd.com>
Signed-off-by: Rohan Potdar <rohan.potdar@amd.com>
Signed-off-by: Rahul Vishwakarma <Rahul.Vishwakarma2@ibm.com>
Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Signed-off-by: jacklin78911-collab <jacklin78911@gmail.com>
Signed-off-by: walterbm <walter.beller.morales@gmail.com>
Signed-off-by: Ting Sun <suntcrick@gmail.com>
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Signed-off-by: amd-ethany <amd-ethany@users.noreply.github.com>
Signed-off-by: Jeff Ma <jeffjma@umich.edu>
Signed-off-by: wzhao18 <wzhao18.sz@gmail.com>
Signed-off-by: Zach Zhu <zzqshu@126.com>
Signed-off-by: Nicklas Frahm <nicklas.frahm@gmail.com>
Signed-off-by: muhammadfawaz1 <135441198+muhammadfawaz1@users.noreply.github.com>
Signed-off-by: Hongxia Yang <hongxia.yang@amd.com>
Signed-off-by: Ace Eldeib <aeldeib@coreweave.com>
Signed-off-by: Ace Eldeib <alexeldeib@gmail.com>
Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai>
Signed-off-by: Tyler Michael Smith <tyler@tylermsmith.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>
Signed-off-by: sungbin1015 <sbin@solbox.com>
Signed-off-by: gcanlin <canlinguosdu@gmail.com>
Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>
Signed-off-by: 辰言 <oncwnuIWp30GguOyJ615Fqj8H-yc@git.weixin.qq.com>
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
Signed-off-by: Randall Smith <Randall.Smith@amd.com>
Signed-off-by: Saddss <28726669061@qq.com>
Signed-off-by: Nick Hill <nickhill123@gmail.com>
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
Signed-off-by: vivek sharma <vivsharm@redhat.com>
Signed-off-by: puneetsharma21 <puneet.sharma21@ibm.com>
Signed-off-by: Puneet Sharma <puneet.sharma21@ibm.com>
Signed-off-by: Robert Shaw <robertgshaw2-redhat@dgx-b200-02.mgmt.accl-001.lab.rdu2.dc.redhat.com>
Signed-off-by: mgoin <mgoin64@gmail.com>
Signed-off-by: Martin Vit <martin@voipmonitor.org>
Signed-off-by: Shawn Tsai <shawnyht@gmail.com>
Signed-off-by: Anna Mayne <anna.mayne@arm.com>
Signed-off-by: Markov Ilya <markovilya19@gmail.com>
Signed-off-by: Chris Leonard <chleonar@redhat.com>
Signed-off-by: Kaihang Jiang <kaihangj@login-lyris02.lyris.clusters.nvidia.com>
Signed-off-by: Djordje Ramic <djoramic@amd.com>
Signed-off-by: Itay Etelis <itay.etelis@ibm.com>
Signed-off-by: Itay Etelis <92247226+Etelis@users.noreply.github.com>
Signed-off-by: Or Ozeri <oro@il.ibm.com>
Signed-off-by: Roberto L. Castro <38211239+LopezCastroRoberto@users.noreply.github.com>
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Signed-off-by: zixi-qi <zixi@inferact.ai>
Signed-off-by: Dino Music <Dino.Music@amd.com>
Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com>
Signed-off-by: jiang1.li <jiang1.li@intel.com>
Signed-off-by: JooHo Lee <BWAAEEEK@users.noreply.github.com>
Signed-off-by: JooHo Lee <jooho414@gmail.com>
Signed-off-by: Xianbao QIAN <xianbao.qian@gmail.com>
Signed-off-by: wenjun.liu <wenjun.liu@intel.com>
Signed-off-by: jun,du <jun.du@intel.com>
Signed-off-by: Benjamin Chislett <bchislett@nvidia.com>
Signed-off-by: Benjamin Chislett <chislett.ben@gmail.com>
Signed-off-by: Cyrus Leung <tlleungac@connect.ust.hk>
Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>
Signed-off-by: khluu <khluu000@gmail.com>
Signed-off-by: Tsvika Shapira <tsvika@moonmath.ai>
Signed-off-by: charlifu <charlifu@amd.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Signed-off-by: deng451e <838677410@qq.com>
Signed-off-by: Teemu Virolainen <teemu.virolainen@amd.com>
Signed-off-by: Zhewen Li <zhewenli@inferact.ai>
Signed-off-by: zihaomu <zmu@amd.com>
Signed-off-by: Micah Williamson <micah.williamson@amd.com>
Signed-off-by: gnovack <novackgm@gmail.com>
Signed-off-by: pei.zhang <pei.zhang@amd.com>
Signed-off-by: augusto.yjh <augusto.yjh@antgroup.com>
Signed-off-by: Change72 <changg@nvidia.com>
Signed-off-by: YanXu <yancey.yx@alibaba-inc.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
Signed-off-by: Lei Gong <gonglei25@huawei.com>
Signed-off-by: Alberto Perdomo <aperdomo@redhat.com>
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: Mingfei Guo <1800012773@pku.edu.cn>
Signed-off-by: Viranjan Pagar <vpagar@nvidia.com>
Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Signed-off-by: reidliu41 <reid201711@gmail.com>
Signed-off-by: Yuchen Fan <functionhx@gmail.com>
Signed-off-by: Xu Zhou <xuzhou9417@163.com>
Signed-off-by: Yipeng Hu <i26268@metax-tech.com>
Signed-off-by: HuYiPeng <144002351+MynameFelix@users.noreply.github.com>
Signed-off-by: larryli2-amd <larryli2@amd.com>
Signed-off-by: larryli2-amd <Larry.Li@amd.com>
Signed-off-by: Ashwin Giridharan <girida@amazon.com>
Signed-off-by: Michael Goin <mgoin64@gmail.com>
Signed-off-by: Joe Rowell <joerowell4@gmail.com>
Signed-off-by: Jimmy Lee <hirejimmylee@gmail.com>
Signed-off-by: Robert Shaw <robertgshaw2-redhat@ip-172-31-18-125.us-east-2.compute.internal>
Signed-off-by: mayuyuace <qiming1.zhang@intel.com>
Signed-off-by: peiyuanz <peiyuanz@inferact.ai>
Signed-off-by: Lai, Yejing <yejing.lai@intel.com>
Signed-off-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com>
Signed-off-by: ErenAta16 <erena6466@gmail.com>
Signed-off-by: Ievgen Bondarenko <ibondarenko@student.sierracollege.edu>
Signed-off-by: Ievgen (Jack) Bondarenko <ibondarenko@student.sierracollege.edu>
Signed-off-by: AlejandroParedesLT <alejandroparedeslatorre@gmail.com>
Signed-off-by: Zhenhui Zhao <zhenhui.zhao@intel.com>
Signed-off-by: vx120 <893600387@qq.com>
Signed-off-by: vx120 <57470515+vx120@users.noreply.github.com>
Signed-off-by: aoshen02 <aoshen@inferact.ai>
Signed-off-by: hcenteno <hugo.centeno@estudiantat.upc.edu>
Signed-off-by: Liran Schour <lirans@il.ibm.com>
Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
Signed-off-by: Alex <alex.tech.lab@outlook.com>
Signed-off-by: Tanish Malekar <tanishmalekar32@gmail.com>
Signed-off-by: Tan Pin Siang <tanpinsiang@gmail.com>
Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com>
Signed-off-by: zzt <zengzetang.zzt@antgroup.com>
Signed-off-by: fai <fangzhouai@gmail.com>
Signed-off-by: guybd <guy.boudoukh@intel.com>
Signed-off-by: Guy Boudoukh <guy.boudoukh@intel.com>
Signed-off-by: Rehan Khan <Rehan.Khan7@ibm.com>
Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>
Signed-off-by: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com>
Signed-off-by: Karthik Kothuri <karthikkothuri2009@gmail.com>
Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>
Signed-off-by: Omer Ullman Argov <118735753+omera-nv@users.noreply.github.com>
Signed-off-by: Yan Ma <yan.ma@intel.com>
Signed-off-by: zengxian <xiangdong.zeng@intel.com>
Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com>
Signed-off-by: Stefan Koncarevic <stefan.koncarevic@amd.com>
Signed-off-by: Gavin Morris <gmorriscs@gmail.com>
Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com>
Signed-off-by: Pavani Majety <pmajety@nvidia.com>
Signed-off-by: shanjiaz <zsjwpianpian@gmail.com>
Signed-off-by: wenpengw-nv <wenpengw@nvidia.com>
Signed-off-by: atalhens <sneh.lata@nutanix.com>
Signed-off-by: Lang Zhao <lang.zhao@galileo.ai>
Signed-off-by: emricksini-h <emrick.birivoutin@hcompany.ai>
Signed-off-by: Zhenzhong1 <zhenzhong.xu@intel.com>
Signed-off-by: Zhenzhong Xu <zhenzhong.xu@intel.com>
Signed-off-by: Matthias Gehre <matthias.gehre@amd.com>
Signed-off-by: Saeid Rostami <srostami@amd.com>
Signed-off-by: nemanjaudovic <nudovic@amd.com>
Signed-off-by: Matt Woodson <mwoodson@redhat.com>
Signed-off-by: wangxiyuan <wangxiyuan1007@gmail.com>
Signed-off-by: Bartosz Stefaniak <bstefaniak@nvidia.com>
Signed-off-by: Hoang Nguyen Tien <hoang.nguyentien.2601@gmail.com>
Signed-off-by: xiangdong <40376367+zxd1997066@users.noreply.github.com>
Signed-off-by: Artur Fierka <artur.fierka@intel.com>
Signed-off-by: gangula-karthik <gkarthik923@gmail.com>
Signed-off-by: Tomer Gilad <tgilad@nvidia.com>
Signed-off-by: Itay Etelis <Itay.etelis@gmail.com>
Signed-off-by: Adhithya Balakrishnan <adhithya.b2004@gmail.com>
Signed-off-by: Daoyuan Li <94409450+DaoyuanLi2816@users.noreply.github.com>
Signed-off-by: srinivas_oo7 <sklinkedin0120@gmail.com>
Signed-off-by: rishitdholakia13 <rishit+github@cohere.com>
Signed-off-by: Divakar Verma <divakar.verma@amd.com>
Signed-off-by: Chaemin Lim <chaemin.lim@mangoboost.io>
Signed-off-by: Edwin Lim <edwinlim0919@gmail.com>
Signed-off-by: Joan Velja <joan.velja22@gmail.com>
Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com>
Signed-off-by: HDCharles <charlesdavidhernandez@gmail.com>
Signed-off-by: Russell Bryant <rbryant@redhat.com>
Signed-off-by: Jane Xu <janeyx@meta.com>
Signed-off-by: Shengqi Chen <i@harrychen.xyz>
Signed-off-by: Mikhail Kostryukov <mike@triptrack.net>
Signed-off-by: Giuseppe Grossi <ggrossi@amd.com>
Signed-off-by: aarushjain29 <aarushi.jain2@amd.com>
Signed-off-by: Turner Jabbour <doubleujabbour@gmail.com>
Co-authored-by: Simon Mo <simon.mo@hey.com>
Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: Benjamin Chislett <chislett.ben@gmail.com>
Co-authored-by: Ameen Patel <ameenp360@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>
Co-authored-by: liuzhenwei <zhenweiliu@habana.ai>
Co-authored-by: Robin <jdmjdm1998@163.com>
Co-authored-by: Roberto L. Castro <38211239+LopezCastroRoberto@users.noreply.github.com>
Co-authored-by: Sting Lin <sting.lin@cienet.com>
Co-authored-by: Nils Matteson <nilsmatteson@icloud.com>
Co-authored-by: Nils Matteson <nils@thaw.sh>
Co-authored-by: Rishabh Saini <rishabhsaini01@gmail.com>
Co-authored-by: Benjamin Chislett <bchislett@nvidia.com>
Co-authored-by: Matthew Bonanni <mbonanni@redhat.com>
Co-authored-by: Nick Cao <ncao@redhat.com>
Co-authored-by: Sahil Kadadekar <147995121+Sahil170595@users.noreply.github.com>
Co-authored-by: vanshbhatia-amd <vansh.bhatia@amd.com>
Co-authored-by: vanshbhatia-amd <210711135+vanshbhatia-amd@users.noreply.github.com>
Co-authored-by: Jason Li <jasonlizhengjian@gmail.com>
Co-authored-by: Juan Pérez de Algaba <124347725+jperezdealgaba@users.noreply.github.com>
Co-authored-by: yzong-rh <yzong@redhat.com>
Co-authored-by: AmeenP <ameen@primeintellect.ai>
Co-authored-by: stefankoncarevic <skoncare@amd.com>
Co-authored-by: Rohan Potdar <rohan.potdar@amd.com>
Co-authored-by: Rahul Vishwakarma <168823860+rahulssv-ibm@users.noreply.github.com>
Co-authored-by: Martin Hickey <martin.hickey@ie.ibm.com>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: Posedge_Lin <jacklin78911@gmail.com>
Co-authored-by: Or Ozeri <oro@il.ibm.com>
Co-authored-by: Walter Beller-Morales <walterbm@users.noreply.github.com>
Co-authored-by: Ting SUN <suntcrick@gmail.com>
Co-authored-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Ethan Yang <ethany@amd.com>
Co-authored-by: amd-ethany <amd-ethany@users.noreply.github.com>
Co-authored-by: Jeff (Junze) Ma <93145857+majunze2001@users.noreply.github.com>
Co-authored-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com>
Co-authored-by: Wei Zhao <51183510+wzhao18@users.noreply.github.com>
Co-authored-by: Zach Zhu <zzqshu@126.com>
Co-authored-by: Nicklas Frahm <nicklas.frahm@gmail.com>
Co-authored-by: Muhammad Fawaz <mfawaz182@gmail.com>
Co-authored-by: Mahad Durrani <114791389+mahadrehmann@users.noreply.github.com>
Co-authored-by: Hongxia Yang <62075498+hongxiayang@users.noreply.github.com>
Co-authored-by: Ace Eldeib <aeldeib@coreweave.com>
Co-authored-by: manayang <jackmanayang@gmail.com>
Co-authored-by: manayang <manayang@tencent.com>
Co-authored-by: Nicolò Lucchesi <nlucches@redhat.com>
Co-authored-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
Co-authored-by: Chaojun Zhang <chaojun.zhang@intel.com>
Co-authored-by: sungbin1015 <sbin@solbox.com>
Co-authored-by: Canlin Guo <canlinguosdu@gmail.com>
Co-authored-by: Thien Tran <gau.nernst@yahoo.com.sg>
Co-authored-by: aoright <102943475+aoright@users.noreply.github.com>
Co-authored-by: 辰言 <oncwnuIWp30GguOyJ615Fqj8H-yc@git.weixin.qq.com>
Co-authored-by: Jee Jee Li <jeejeelee@inferact.ai>
Co-authored-by: Jee Jee Li <pandaleefree@gmail.com>
Co-authored-by: Bugen Zhao <i@bugenzhao.com>
Co-authored-by: rasmith <Randall.Smith@amd.com>
Co-authored-by: Saddss <108515797+Saddss@users.noreply.github.com>
Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai>
Co-authored-by: Saddss <28726669061@qq.com>
Co-authored-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: Vivek Sharma <Vivek.Sharma20@ibm.com>
Co-authored-by: vivek sharma <vivsharm@redhat.com>
Co-authored-by: Puneet Sharma <puneet.sharma21@ibm.com>
Co-authored-by: depthfirst-app[bot] <184448029+depthfirst-app[bot]@users.noreply.github.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
Co-authored-by: Robert Shaw <robertgshaw2-redhat@dgx-b200-02.mgmt.accl-001.lab.rdu2.dc.redhat.com>
Co-authored-by: Michael Goin <mgoin64@gmail.com>
Co-authored-by: Martin Vit <martin@voipmonitor.org>
Co-authored-by: shawn <shawnyht@gmail.com>
Co-authored-by: almayne <anna.mayne@arm.com>
Co-authored-by: Ilya Markov <markovilya197@gmail.com>
Co-authored-by: Markov Ilya <markovilya19@gmail.com>
Co-authored-by: Chris Leonard <chleonar@redhat.com>
Co-authored-by: Kaihang Jiang <88449510+kjiang249@users.noreply.github.com>
Co-authored-by: djramic <djoramic@amd.com>
Co-authored-by: Itay Etelis <92247226+Etelis@users.noreply.github.com>
Co-authored-by: Itay Etelis <itay.etelis@ibm.com>
Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: qizixi <22851944+zixi-qi@users.noreply.github.com>
Co-authored-by: music-dino <111048524+music-dino@users.noreply.github.com>
Co-authored-by: Chauncey <chaunceyjiang@gmail.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
Co-authored-by: JooHo Lee <96564470+BWAAEEEK@users.noreply.github.com>
Co-authored-by: JooHo Lee <BWAAEEEK@users.noreply.github.com>
Co-authored-by: Tiezhen WANG <38108242+xianbaoqian@users.noreply.github.com>
Co-authored-by: wenjun liu <wenjun.liu@intel.com>
Co-authored-by: jun,du <jun.du@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk>
Co-authored-by: Woosuk Kwon <woosuk@inferact.ai>
Co-authored-by: Kevin H. Luu <khluu000@gmail.com>
Co-authored-by: Brandon Pelfrey <bpelfrey@nvidia.com>
Co-authored-by: Tsvika Shapira <tsvikas@gmail.com>
Co-authored-by: Charlie Fu <charlifu@amd.com>
Co-authored-by: Codex <codex@openai.com>
Co-authored-by: weishu <838677410@qq.com>
Co-authored-by: Teemu Virolainen <teemu.virolainen@amd.com>
Co-authored-by: TJian <tunjian.tan@embeddedllm.com>
Co-authored-by: Zhewen Li <zhewenli@meta.com>
Co-authored-by: Zhewen Li <zhewenli@inferact.ai>
Co-authored-by: ZihaoMu <zmu@amd.com>
Co-authored-by: Micah Williamson <micah.williamson@amd.com>
Co-authored-by: ap9272 <ap9272@users.noreply.github.com>
Co-authored-by: gnovack <novackgm@gmail.com>
Co-authored-by: peizhang56 <pei.zhang@amd.com>
Co-authored-by: Augusto Yao <augusto.yjh@antgroup.com>
Co-authored-by: Chang Guo <changg@nvidia.com>
Co-authored-by: Yan Xu <yancey.yx@alibaba-inc.com>
Co-authored-by: Jingyi Yang <girasoleyang@gmail.com>
Co-authored-by: GongLei-HW <1327185943@qq.com>
Co-authored-by: Lei Gong <gonglei25@huawei.com>
Co-authored-by: alberto <aperdomo@redhat.com>
Co-authored-by: Jiangyun Zhu <riverclouds.zhu@qq.com>
Co-authored-by: Mingfei Guo <52491257+guoriyue@users.noreply.github.com>
Co-authored-by: Andrey Talman <atalman@fb.com>
Co-authored-by: Andrey Talman <atalman@users.noreply.github.com>
Co-authored-by: ViranjanPagar <vpagar@nvidia.com>
Co-authored-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Co-authored-by: Roger Wang <hey@rogerw.io>
Co-authored-by: Reid <61492567+reidliu41@users.noreply.github.com>
Co-authored-by: FAN YUCHEN <2994114386@qq.com>
Co-authored-by: XuZhou <17717803682@163.com>
Co-authored-by: Xu Zhou <xuzhou9417@163.com>
Co-authored-by: Hoseung Kim <ghyutjik123@gmail.com>
Co-authored-by: HuYiPeng <144002351+MynameFelix@users.noreply.github.com>
Co-authored-by: Yipeng Hu <i26268@metax-tech.com>
Co-authored-by: larryli2-amd <Larry.Li@amd.com>
Co-authored-by: Ashwin Giridharan <ashwing@users.noreply.github.com>
Co-authored-by: Joe Rowell <joerowell4@gmail.com>
Co-authored-by: Jimmy Lee <58957694+thisisjimmyfb@users.noreply.github.com>
Co-authored-by: Robert Shaw <robertgshaw2-redhat@ip-172-31-18-125.us-east-2.compute.internal>
Co-authored-by: Qiming Zhang <qiming1.zhang@intel.com>
Co-authored-by: peiyuanz <peiyuanz@inferact.ai>
Co-authored-by: zhouzhou <zhouzhou@zhouzhoudeMacBook-Pro.local>
Co-authored-by: Yejing Lai <yejing.lai@intel.com>
Co-authored-by: ErenAta16 <erena6466@gmail.com>
Co-authored-by: Ievgen Bondarenko <ibondarenko@student.sierracollege.edu>
Co-authored-by: Alejandro Paredes La Torre <99832715+AlejandroParedesLT@users.noreply.github.com>
Co-authored-by: zhao, zhenhui <zhenhui.zhao@intel.com>
Co-authored-by: vx120 <57470515+vx120@users.noreply.github.com>
Co-authored-by: crp0128 <191679376@qq.com>
Co-authored-by: aoshen02 <aoshen@inferact.ai>
Co-authored-by: yewentao256 <zhyanwentao@126.com>
Co-authored-by: Hugo Centeno <133872718+hugo-cen@users.noreply.github.com>
Co-authored-by: liranschour <liranschour@users.noreply.github.com>
Co-authored-by: omerpaz95 <73347585+omerpaz95@users.noreply.github.com>
Co-authored-by: AlexHuang <alex.tech.lab@outlook.com>
Co-authored-by: Tanish Malekar <60835372+tanish-malekar@users.noreply.github.com>
Co-authored-by: Tan Pin Siang <tanpinsiang@gmail.com>
Co-authored-by: vllmellm <vllm.ellm@embeddedllm.com>
Co-authored-by: Hongxia Yang <hongxia.yang@amd.com>
Co-authored-by: Jun Kang Chow <junkangchow@gmail.com>
Co-authored-by: zzt <zengzetang.zzt@antgroup.com>
Co-authored-by: Fangzhou Ai <31551580+Fangzhou-Ai@users.noreply.github.com>
Co-authored-by: guybd <guy.boudoukh@intel.com>
Co-authored-by: Rehan Khan <Rehan.Khan7@ibm.com>
Co-authored-by: Matthew Wong <Matthew.Wong2@amd.com>
Co-authored-by: Mohammad Miadh Angkad …
plasticchris added a commit to plasticchris/vllm that referenced this pull request Jul 20, 2026
…1100

Native HIP skinny-decode GEMM (wvSplitK_int4_g) + Triton prefill for W4A16.
Supports compressed-tensors asymmetric uint4 (group_size 32/64/128), which the
symmetric-only RDNA3 kernel (PR vllm-project#41394) rejects. Registered ahead of
TritonW4A16LinearKernel in the ROCm mixed-precision priority list.

Adds on_gfx1151() to platforms/rocm.py (required by the hybrid prefill path).

Decode throughput on 2x RX 7900 XTX, Qwen3.6-27B W4A16-G32 asym, TP=2:
  ctx    50:  16.4 -> 28.1 tok/s (1.71x)
  ctx  8000:  14.5 -> 23.5 tok/s (1.62x)
  ctx 24000:  12.2 -> 18.6 tok/s (1.52x)
  ctx 48000:   9.9 -> 14.1 tok/s (1.42x)

Correctness: tests/kernels/quantization/test_rdna_hybrid_w4a16.py 83/83 pass
(forward parity sym+asym, all shapes/group sizes; process_weights repack).
BlivionIaG added a commit to opengfx1030/vllm-rdna that referenced this pull request Sep 18, 2026
…t#40977)

The cherry-pick of 923672d (rdna2 w4a16 multi-config prefill) deleted
this test file because rdna2_extras source was on v0.27.1 base and did
not yet have PR vllm-project#40977 (HybridW4A16LinearKernel). On v0.28.0 base the
file exists and tests HybridW4A16LinearKernel — unrelated to RDNA2's
dispatcher. Restore it from the v0.28.0 base.

Verified on gfx1030 4xV620:
- _rocm_C.abi3.so builds clean (22 _rocm_C:: ops registered, 15 RDNA2)
- pytest tests/kernels/quantization/test_rdna2_w4a16.py: 18/18 PASSED
- vllm.LLM eager-mode smoke Qwen3.5-0.8B: PASSED on GPU 1, generated
  'Paris, France.\nTrue or is' from 'The capital of France is'

Co-authored-by: MiniMax-M3 <agent@minimax.local>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ci/build performance Performance-related issues ready ONLY add when PR is ready to merge/full CI is needed rocm Related to AMD ROCm

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

8 participants