[JartX]

Motivation

This work was driven by a concrete production problem: on RDNA3 the only fast W4A16 path was ExLlama, which is fp16-only, and casting bf16-trained models down to fp16 introduces numerical instability. Modern model families (Qwen2/Qwen3, Qwen3-Coder, Llama 3.x, Mistral, etc.) are trained and released in bf16; their weight ranges and activation magnitudes routinely exceed fp16's ±65504 dynamic range, producing inf/NaN propagation, silently degraded outputs, and garbled tool-calling / structured-output generations under load. The Triton W4A16 kernel does support bf16, but at ~82 tk/s it's roughly 3× slower than ExLlama fp16, forcing users into a "fast but numerically wrong" vs "correct but unusably slow" trade-off.

This PR removes that trade-off: native bf16 W4A16 at 205–345 tk/s (2.5–4.2× faster than Triton bf16), plus a fp16 path that beats ExLlama by 6–17% for users who don't need bf16. Models stay in their native bf16 dtype, no manual casts, no quality regression.

Summary

Adds a native HIP W4A16 kernel for AMD RDNA3 (gfx1100/1101/1102 — RX 7900 XT/XTX class), registered as RDNA3W4A16LinearKernel ahead of the Triton kernel on the ROCm priority list. Falls through gracefully on non-RDNA3 ROCm devices (CDNA / MI300) via can_implement() gating on vllm.platforms.rocm.on_gfx11().

Two compute paths share a single torch.ops._C.gptq_gemm_rdna3 entry (all dispatch in C++ to keep apply_weights torch.compile-friendly):

• Scalar dot-product kernel for decode and any M < 16. fp16 uses v_dot2_f32_f16 (__builtin_amdgcn_fdot2); bf16 widens to fp32 to sidestep the missing v_pk_fma_bf16 on gfx11.
• v_wmma_f32_16x16x16 kernel for bf16 prefill / batched (M ≥ 16). Two variants share a launcher: v1 (1 wave/block, 16M×16N tile) for M < 32, v2 (2 waves/block, 32M×16N tile, double-buffered LDS) for M ≥ 32.

The full architecture write-up, fragment-layout discussion, lessons learned (including reverted experiments) and bench commands live in csrc/quantization/gptq/README_RDNA3.md.

No open PR adds a native HIP W4A16 kernel for RDNA3. ExLlama (fp16-only) and Triton W4A16 are the existing options on ROCm; this kernel beats both at every concurrency tested and adds bf16 support that ExLlama lacks.

Note:
Also can load AWQ Like cyankiwi/Qwen3.5-9B-AWQ-BF16-INT4 or cyankiwi/Qwen3.5-9B-AWQ-INT4

End-to-end performance

All benchmarks run with evalscope perf openqa against vLLM serving Qwen3.6-27B-GPTQ-W4A16-G32 on RX 7900 XTX, tp=2, 50 requests, concurrency 100, output ~1800 tok.

Summary table (Qwen3.6-27B-GPTQ-W4A16-G32)

Kernel	dtype	max-num-seqs=8	max-num-seqs=32
Triton W4A16	bf16	82.4 tk/s	—
Triton W4A16	fp16	83.2 tk/s	—
ExLlama (no bf16)	fp16	255.0 tk/s	382.5 tk/s
RDNA3 W4A16 (this PR)	bf16	205.3 tk/s	345.6 tk/s
RDNA3 W4A16 (this PR)	fp16	270.2 tk/s	445.7 tk/s

Common evalscope client command

evalscope perf \
  --url "http://127.0.0.1/v1/chat/completions" \
  --parallel 100 \
  --model INCCODER \
  --number 50 \
  --api openai \
  --dataset openqa \
  --stream

12 files changed, +2820 / 0 lines.

Test plan

• Build the C++ extension in a ROCm container (gfx1100 target):

  uv pip install -e . --torch-backend=auto
  

• Correctness — gptq_gemm_rdna3_wmma_probe op confirms mode-1 fragment layout with random-A.
• Correctness — v1 vs v2 kernel parity check across multiple shapes / M / dtype:

  # max abs diff between v1 and v2 outputs is at numerical-rounding level# across {M=32..2048, K=N ∈ {4096, 5120, 11008}, fp16 / bf16}
  

• Microbench M-sweep (csrc/quantization/gptq/tools/bench_prefill_msweep.py) over the 5 Qwen-class shapes confirms expected behaviour: WMMA bf16 wins at M ≥ 16, scalar fp16 stays competitive across all M, v2 saves 5–7% vs v1 at M ≥ 64.
• End-to-end evalscope perf openqa against vllm serve on RX 7900 XTX, tp=2 — numbers in the table above.
• Verified the kernel does NOT regress decode-only configs: pre-PR ExLlama fp16 baselines reproduced when the kernel is disabled via VLLM_DISABLED_KERNELS=RDNA3W4A16LinearKernel.
• VGPR / spill audit (gfx1100 .note.AMDGPU.metadata static analysis): bf16 M_COUNT=8 went from 192-VGPR cap with 144 spilled to scratch → 130 VGPR / 0 spill. fp16 path register-clean throughout. WMMA v1/v2 both 43–44 VGPR.

Hardware tested

• AMD Radeon RX 7900 XTX (gfx1100, 96 CUs, 24 GB)
• ROCm runtime: rocm-X series (hipcc 7.2.1)
• Model: Qwen3.6-27B-GPTQ-W4A16-G32 (group_size=32)
• vLLM serving config: --attention-backend TRITON_ATTN --kv-cache-dtype fp8 --enable-prefix-caching (full flags in README)

The kernel is gated to gfx11 only; CDNA / MI3xx / RDNA4 fall through to TritonW4A16LinearKernel automatically.

Out of scope (tracked in README "Known limitations")

• fp16 → WMMA dispatch: tested end-to-end, was a wash with scalar+fdot2 (445.7 vs 440.7 tk/s, within run-to-run variance) despite a 47% kernel microbench advantage. Dispatch stays bf16-only; the standalone WMMA op remains callable for fp16 if a future workload benefits.
• Multi-wave WMMA with LDS-shared A (not just B). Realistic 10–15% more on bf16 prefill but a significantly bigger rewrite.
• fp16 → fp32-output-buffer experiment: untested.
• rocprof returns 0 for most counters on gfx1100 (only SQ_BUSY_CYCLES and SQ_WAVES work in practice on rocprofv3) — limits diagnostic depth.

## Motivation

This work was driven by a concrete production problem: on RDNA3 the only fast W4A16 path was ExLlama, which is fp16-only, and casting bf16-trained models down to fp16 introduces numerical instability. Modern model families (Qwen2/Qwen3, Qwen3-Coder, Llama 3.x, Mistral, etc.) are trained and released in bf16; their weight ranges and activation magnitudes routinely exceed fp16's ±65504 dynamic range, producing inf/NaN propagation, silently degraded outputs, and garbled tool-calling / structured-output generations under load. The Triton W4A16 kernel does support bf16, but at ~82 tk/s it's roughly 3× slower than ExLlama fp16, forcing users into a "fast but numerically wrong" vs "correct but unusably slow" trade-off.

This PR removes that trade-off: native bf16 W4A16 at 205–345 tk/s (2.5–4.2× faster than Triton bf16), plus a fp16 path that beats ExLlama by 6–17% for users who don't need bf16. Models stay in their native bf16 dtype, no manual casts, no quality regression.

Summary

Adds a native HIP W4A16 kernel for AMD RDNA3 (gfx1100/1101/1102 — RX 7900 XT/XTX class), registered as RDNA3W4A16LinearKernel ahead of the Triton kernel on the ROCm priority list. Falls through gracefully on non-RDNA3 ROCm devices (CDNA / MI300) via can_implement() gating on vllm.platforms.rocm.on_gfx11().

Two compute paths share a single torch.ops._C.gptq_gemm_rdna3 entry (all dispatch in C++ to keep apply_weights torch.compile-friendly):

• Scalar dot-product kernel for decode and any M < 16. fp16 uses v_dot2_f32_f16 (__builtin_amdgcn_fdot2); bf16 widens to fp32 to sidestep the missing v_pk_fma_bf16 on gfx11.
• v_wmma_f32_16x16x16 kernel for bf16 prefill / batched (M ≥ 16). Two variants share a launcher: v1 (1 wave/block, 16M×16N tile) for M < 32, v2 (2 waves/block, 32M×16N tile, double-buffered LDS) for M ≥ 32.

The full architecture write-up, fragment-layout discussion, lessons learned (including reverted experiments) and bench commands live in [csrc/quantization/gptq/README_RDNA3.md](https://claude.ai/chat/csrc/quantization/gptq/README_RDNA3.md).

No open PR adds a native HIP W4A16 kernel for RDNA3. ExLlama (fp16-only) and Triton W4A16 are the existing options on ROCm; this kernel beats both at every concurrency tested and adds bf16 support that ExLlama lacks.

End-to-end performance

All benchmarks run with evalscope perf openqa against vLLM serving Qwen3.6-27B-GPTQ-W4A16-G32 on RX 7900 XTX, tp=2, 50 requests, concurrency 100, output ~1800 tok.

Summary table (Qwen3.6-27B-GPTQ-W4A16-G32)

Kernel	dtype	max-num-seqs=8	max-num-seqs=32
Triton W4A16	bf16	82.4 tk/s	—
Triton W4A16	fp16	83.2 tk/s	—
ExLlama (no bf16)	fp16	255.0 tk/s	382.5 tk/s
RDNA3 W4A16 (this PR)	bf16	205.3 tk/s	345.6 tk/s
RDNA3 W4A16 (this PR)	fp16	270.2 tk/s	445.7 tk/s

Full comparative table (Qwen3.6-27B-GPTQ-W4A16-G32)

All runs: RX 7900 XTX, tp=2, 50 requests, concurrency 100, ~1800 output tokens.

Kernel	dtype	max-num-seqs	Output rate (tk/s)	Total time (s)	Avg lat. (s)	P99 lat. (s)	Avg TTFT (s)	Avg TPOT (s)	Speedup vs Triton*
Triton W4A16	bf16	8	82.37	1068.21	591.27	1068.20	429.72	0.092	1.00× (baseline)
Triton W4A16	fp16	8	83.23	1066.78	564.48	1066.77	408.99	0.088	1.01×
ExLlama	fp16	8	255.00	346.19	195.14	346.18	142.82	0.030	3.10×
ExLlama	fp16	32	382.52	229.91	154.43	229.90	42.83	0.064	4.65×
RDNA3 W4A16 (this PR)	bf16	8	205.34	430.76	236.44	430.75	172.70	0.036	2.49×
RDNA3 W4A16 (this PR)	bf16	32	345.64	259.83	172.64	259.83	48.81	0.069	4.20×
RDNA3 W4A16 (this PR)	fp16	8	270.19	332.03	188.50	332.02	137.84	0.028	3.28×
RDNA3 W4A16 (this PR)	fp16	32	445.74	203.65	137.94	203.65	39.61	0.054	5.41×

^* Speedup vs Triton W4A16 bf16 @ max-num-seqs=8 (slowest baseline).

Head-to-head deltas

Comparison	Output rate	Δ
RDNA3 fp16 vs Triton fp16 (seqs=8)	270.19 vs 83.23 tk/s	+224.6% (3.25×)
RDNA3 bf16 vs Triton bf16 (seqs=8)	205.34 vs 82.37 tk/s	+149.3% (2.49×)
RDNA3 fp16 vs ExLlama fp16 (seqs=8)	270.19 vs 255.00 tk/s	+5.96%
RDNA3 fp16 vs ExLlama fp16 (seqs=32)	445.74 vs 382.52 tk/s	+16.5%
RDNA3 bf16 vs ExLlama fp16 (seqs=8)†	205.34 vs 255.00 tk/s	−19.5%
RDNA3 bf16 vs ExLlama fp16 (seqs=32)†	345.64 vs 382.52 tk/s	−9.6%

^† ExLlama does not support bf16; this comparison is included only to quantify the cost of bf16 vs fp16 on the new kernel against the best fp16-only alternative. For the apples-to-apples bf16 comparison, see the Triton bf16 row (RDNA3 bf16 is 2.49–4.20× faster).

Detailed benchmarks (Qwen3.6-27B-GPTQ-W4A16-G32)

ExLlama — fp16, max-num-seqs=8

VLLM_DISABLED_KERNELS=TritonW4A16LinearKernel,ConchLinearKernel,RDNA3W4A16LinearKernel \
VLLM_TUNED_CONFIG_FOLDER=/app/moe_configs \
vllm serve /models/Qwen3.6-27B-GPTQ-W4A16-G32 \
  --gpu-memory-utilization 0.95 --max_model_len 262144 -tp 2 --port 8000 \
  --attention-backend TRITON_ATTN --enable-prefix-caching \
  --max-num-seqs 8 --kv-cache-dtype fp8 --dtype float16

Metric	Value
Total generated	88,277 tokens
Avg output rate	255.00 tk/s
Total test time	346.19 s
Avg latency	195.140 s
P99 latency	346.182 s
Avg TTFT	142.824 s
Avg TPOT	0.030 s
Success rate	100.0%

ExLlama — fp16, max-num-seqs=32

VLLM_DISABLED_KERNELS=TritonW4A16LinearKernel,ConchLinearKernel,RDNA3W4A16LinearKernel \
VLLM_TUNED_CONFIG_FOLDER=/app/moe_configs \
vllm serve /models/Qwen3.6-27B-GPTQ-W4A16-G32 \
  --gpu-memory-utilization 0.95 --max_model_len 262144 -tp 2 --port 8000 \
  --attention-backend TRITON_ATTN --enable-prefix-caching \
  --max-num-seqs 32 --kv-cache-dtype fp8 --dtype float16

Metric	Value
Total generated	87,947 tokens
Avg output rate	382.52 tk/s
Total test time	229.91 s
Avg latency	154.427 s
P99 latency	229.904 s
Avg TTFT	42.827 s
Avg TPOT	0.064 s
Success rate	100.0%

Triton W4A16 — bf16, max-num-seqs=8

VLLM_DISABLED_KERNELS=ConchLinearKernel,RDNA3W4A16LinearKernel \
VLLM_TUNED_CONFIG_FOLDER=/app/moe_configs \
vllm serve /models/Qwen3.6-27B-GPTQ-W4A16-G32 \
  --gpu-memory-utilization 0.95 --max_model_len 262144 -tp 2 --port 8000 \
  --attention-backend TRITON_ATTN --enable-prefix-caching \
  --max-num-seqs 8 --kv-cache-dtype fp8

Metric	Value
Total generated	87,993 tokens
Avg output rate	82.37 tk/s
Total test time	1068.21 s
Avg latency	591.269 s
P99 latency	1068.199 s
Avg TTFT	429.724 s
Avg TPOT	0.092 s
Success rate	100.0%

Triton W4A16 — fp16, max-num-seqs=8

VLLM_DISABLED_KERNELS=ConchLinearKernel,RDNA3W4A16LinearKernel \
VLLM_TUNED_CONFIG_FOLDER=/app/moe_configs \
vllm serve /models/Qwen3.6-27B-GPTQ-W4A16-G32 \
  --gpu-memory-utilization 0.95 --max_model_len 262144 -tp 2 --port 8000 \
  --attention-backend TRITON_ATTN --enable-prefix-caching \
  --max-num-seqs 8 --kv-cache-dtype fp8 --dtype float16

Metric	Value
Total generated	88,784 tokens
Avg output rate	83.23 tk/s
Total test time	1066.78 s
Avg latency	564.476 s
P99 latency	1066.771 s
Avg TTFT	408.985 s
Avg TPOT	0.088 s
Success rate	100.0%

RDNA3 W4A16 (this PR) — fp16, max-num-seqs=8

VLLM_TUNED_CONFIG_FOLDER=/app/moe_configs \
vllm serve /models/Qwen3.6-27B-GPTQ-W4A16-G32 \
  --gpu-memory-utilization 0.95 --max_model_len 262144 -tp 2 --port 8000 \
  --attention-backend TRITON_ATTN --enable-prefix-caching \
  --max-num-seqs 8 --kv-cache-dtype fp8 --dtype float16

Metric	Value
Total generated	89,711 tokens
Avg output rate	270.19 tk/s
Total test time	332.03 s
Avg latency	188.498 s
P99 latency	332.022 s
Avg TTFT	137.837 s
Avg TPOT	0.028 s
Success rate	100.0%

RDNA3 W4A16 (this PR) — bf16, max-num-seqs=8

VLLM_TUNED_CONFIG_FOLDER=/app/moe_configs \
vllm serve /models/Qwen3.6-27B-GPTQ-W4A16-G32 \
  --gpu-memory-utilization 0.95 --max_model_len 262144 -tp 2 --port 8000 \
  --attention-backend TRITON_ATTN --enable-prefix-caching \
  --max-num-seqs 8 --kv-cache-dtype fp8

Metric	Value
Total generated	88,453 tokens
Avg output rate	205.34 tk/s
Total test time	430.76 s
Avg latency	236.441 s
P99 latency	430.752 s
Avg TTFT	172.696 s
Avg TPOT	0.036 s
Success rate	100.0%

RDNA3 W4A16 (this PR) — bf16, max-num-seqs=32

VLLM_TUNED_CONFIG_FOLDER=/app/moe_configs \
vllm serve /models/Qwen3.6-27B-GPTQ-W4A16-G32 \
  --gpu-memory-utilization 0.95 --max_model_len 262144 -tp 2 --port 8000 \
  --attention-backend TRITON_ATTN --enable-prefix-caching \
  --max-num-seqs 32 --kv-cache-dtype fp8

Metric	Value
Total generated	89,808 tokens
Avg output rate	345.64 tk/s
Total test time	259.83 s
Avg latency	172.637 s
P99 latency	259.827 s
Avg TTFT	48.806 s
Avg TPOT	0.069 s
Success rate	100.0%

RDNA3 W4A16 (this PR) — fp16, max-num-seqs=32

VLLM_TUNED_CONFIG_FOLDER=/app/moe_configs \
vllm serve /models/Qwen3.6-27B-GPTQ-W4A16-G32 \
  --gpu-memory-utilization 0.95 --max_model_len 262144 -tp 2 --port 8000 \
  --attention-backend TRITON_ATTN --enable-prefix-caching \
  --max-num-seqs 32 --kv-cache-dtype fp8 --dtype float16

Metric	Value
Total generated	90,775 tokens
Avg output rate	445.74 tk/s
Total test time	203.65 s
Avg latency	137.941 s
P99 latency	203.645 s
Avg TTFT	39.607 s
Avg TPOT	0.054 s
Success rate	100.0%

Common evalscope client command

evalscope perf \
  --url "http://127.0.0.1/v1/chat/completions" \
  --parallel 100 \
  --model INCCODER \
  --number 50 \
  --api openai \
  --dataset openqa \
  --stream

12 files changed, +2820 / 0 lines.

Test plan

• Build the C++ extension in a ROCm container (gfx1100 target):

  uv pip install -e . --torch-backend=auto

• Correctness — gptq_gemm_rdna3_wmma_probe op confirms mode-1 fragment layout with random-A.
• Correctness — v1 vs v2 kernel parity check across multiple shapes / M / dtype:

  # max abs diff between v1 and v2 outputs is at numerical-rounding level
  # across {M=32..2048, K=N ∈ {4096, 5120, 11008}, fp16 / bf16}

• Microbench M-sweep (csrc/quantization/gptq/tools/bench_prefill_msweep.py) over the 5 Qwen-class shapes confirms expected behaviour: WMMA bf16 wins at M ≥ 16, scalar fp16 stays competitive across all M, v2 saves 5–7% vs v1 at M ≥ 64.
• End-to-end evalscope perf openqa against vllm serve on RX 7900 XTX, tp=2 — numbers in the table above.
• Verified the kernel does NOT regress decode-only configs: pre-PR ExLlama fp16 baselines reproduced when the kernel is disabled via VLLM_DISABLED_KERNELS=RDNA3W4A16LinearKernel.
• VGPR / spill audit (gfx1100 .note.AMDGPU.metadata static analysis): bf16 M_COUNT=8 went from 192-VGPR cap with 144 spilled to scratch → 130 VGPR / 0 spill. fp16 path register-clean throughout. WMMA v1/v2 both 43–44 VGPR.

Hardware tested

• AMD Radeon RX 7900 XTX (gfx1100, 96 CUs, 24 GB)
• ROCm runtime: rocm-X series (hipcc 7.2.1)
• Model: Qwen3.6-27B-GPTQ-W4A16-G32 (group_size=32)
• vLLM serving config: --attention-backend TRITON_ATTN --kv-cache-dtype fp8 --enable-prefix-caching (full flags in README)

The kernel is gated to gfx11 only; CDNA / MI3xx / RDNA4 fall through to TritonW4A16LinearKernel automatically.

Out of scope (tracked in README "Known limitations")

• fp16 → WMMA dispatch: tested end-to-end, was a wash with scalar+fdot2 (445.7 vs 440.7 tk/s, within run-to-run variance) despite a 47% kernel microbench advantage. Dispatch stays bf16-only; the standalone WMMA op remains callable for fp16 if a future workload benefits.
• Multi-wave WMMA with LDS-shared A (not just B). Realistic 10–15% more on bf16 prefill but a significantly bigger rewrite.
• fp16 → fp32-output-buffer experiment: untested.
• rocprof returns 0 for most counters on gfx1100 (only SQ_BUSY_CYCLES and SQ_WAVES work in practice on rocprofv3) — limits diagnostic depth.

IMPORTANT(UPDATE: 2026-05-11):

This kernel is complemented by the following HIP kernels for kvcache int8_per_token_head and int4_per_token_head from the branch:

https://github.com/JartX/vllm/tree/feat/rdna3_int8_int4_hip_kernels

which also leverage the mechanisms of the following pull request:

#40835

Currently, all developments are being compiled into

https://github.com/JartX/vllm/tree/perf/rdna3_full_stack

It can be compiled with:

docker build -f docker/Dockerfile.rocm_fa -t vllm-rocm-rdna3-full-stack .

The Dockerfile includes a temporary fix to enable RCCL, as there's a confirmed bug in the official rccl version on AMD.

ROCm/legacy-rocm-build#6074 (comment)

Regards

[claude]

Claude Code Review

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

[gemini-code-assist]

Code Review

This pull request introduces optimized W4A16 GPTQ kernels specifically designed for AMD RDNA3 (gfx1100) GPUs, providing a high-performance alternative to Triton and ExLlama kernels on the ROCm path. The implementation includes a scalar dot-product path for decode and a WMMA-based matrix-instruction path for prefill, incorporating hardware-specific optimizations such as bf16 widening to fp32 and packed CAS-64 atomic accumulation to overcome RDNA3 ISA limitations. Review feedback suggests replacing #define macros with constexpr constants in the C++ source files to improve type safety and scoping in accordance with modern C++ standards.

[mergify]

Hi @JartX, 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

[AndreasKaratzas]

This looks decent, but before I dive deeper into this I would like to see what other AMD engineers think of this. While this is a good feature to be added, it also forces us to commit into maintaining such features. I am not sure if this should be upstream or in a fork. I'll leave others to comment on that as well.

cc @micah-wil @Rohan138

UPDATE: I feel I owe a bit more explanation regarding this. One of the reasons that I say this is because our CI does not have RDNA3 devices. So one of the pitfalls is that any upcoming regression here would not be detected.

[AndreasKaratzas]

This is getting out of control. There should be a dataclass probably or something else for this one.

[Rohan138]

cc @gshtras this is likely going to get worse in the future with gfx12xx, do we have any threads/thoughts on cleaning this up?

[JartX]

Hello both, the comment is pending removal; it was just explaining why it was approved, but it was a topic to discuss with you, because the list is growing :S

[JartX]

I've also uploaded the readme file as a log, so you can tell me where I could document the architecture, tests, etc., how the kernel is structured, because I had to do it through reverse engineering.

The readme shouldn't just stay there.

[gshtras]

My only concern here is gfx115x which is not a Radeon. Should it also be included?

[JartX]

@gshtras I can find a way to skip it

[mergify]

Hi @JartX, 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

[BowenBao]

could this kernel work for general w4a16 (e.g., awq, etc) or is it specific to gptq?

[JartX]

@BowenBao
Yes, you can also load cyankiwi/Qwen3.5-9B-AWQ-BF16-INT4 or cyankiwi/Qwen3.5-9B-AWQ-INT4

The GPTQ is to follow the convention; I've also changed the references in the PR text :)

[mgehre-amd]

Hey, nice to see those improvements!

I had opened #40977 to address w4a16 for both fp16 and bf16 activations on gfx11/gfx12 (and specifically tuned for Strix Halo) using a new HIP kernel for the skinny GEMM (decode) and a Triton kernel for prefill.

I will check your PR in more detail to see where our approaches differ and how they could be combined.

[JartX]

@mgehre-amd, thanks for taking a look. My implementation is focused solely on the GFX1100; I've gone into the nitty-gritty down to the core. But if we can combine our efforts, even better. I'd also like you to see the downstream version I have.

[JartX]

@mgehre-amd @AndreasKaratzas @Rohan138

https://github.com/JartX/vllm/tree/perf/rdna3_full_stack
https://github.com/JartX/vllm/blob/perf/rdna3_full_stack/csrc/quantization/gptq/README_RDNA3.md
https://github.com/JartX/vllm/blob/perf/rdna3_full_stack/csrc/quantization/gptq/README_RDNA3_CVT_ISA.md
https://github.com/JartX/vllm/blob/perf/rdna3_full_stack/csrc/quantization/gptq/README_RDNA3_FULL_ISA.md
https://github.com/JartX/vllm/blob/perf/rdna3_full_stack/csrc/attention/README_RDNA3_HIP_KERNELS.md

For generate the kvache int8_per_tensor scales

https://github.com/Neurothic-io/generate_int8_scales_for_compressed_tensors

[rasmith]

I'd like to see this broken up into smaller PRs with tests. For example, I'd like to see a test for each custom op.

[JartX]

@rasmith Yes, the last links refer to documentation and how the branch rdna3_full_stack (It's not from this PR) about having right people with the right expertise to have a lookis structured and how many branches it comprises; they were just in case you wanted to take a look. For now, efforts will be focused on working with @mgehre-amd to see how we can combine our efforts. The other branch with everything won't be released all at once, not by a long shot; that would be crazy. But the kernel is the kernel, along with the performance data.

[mgehre-amd]

Thanks again for the detailed write-up — your gfx1100 numbers look great. I gave the kernel a closer look on AMD Strix Halo (gfx1151, RDNA 3.5, 40 CUs, LPDDR5X-8000) since the on_gfx11() gate matches gfx1151 too. Sharing all the numbers in case they're useful — fully understand gfx1151 is outside the PR's stated target.

I compared this PR head-to-head against the hybrid W4A16 kernel I'm working on for gfx1151 (#40977). To make it apples-to-apples I rebased my branch onto the same upstream base as this PR (b4806c8ee1, vllm-project/vllm main from 2026-04-30), and reran the same benchmarks for both.

Setup

• vLLM built locally with uv pip install -e ., torch 2.10.0 + ROCm nightly, gfx1151 venv.
• Both kernels: same upstream base, same venv, same bench command (only --model/--input-len vary):

  vllm bench latency --model <MODEL> \
    --input-len <128|1920> --output-len 128 \
    --max-model-len 4096 --max-num-seqs 1

• Each rep = full server restart + warmup. 3 reps for the main 4 cases (variance), 1 rep for the broader 10-case sweep.
• Each run wrapped in an exclusive GPU lock — no concurrent GPU work.
• Verified the expected kernel name in every startup log (Using RDNA3W4A16LinearKernel or Using HybridW4A16LinearKernel).

Variance (3 reps × 4 cases)

Run-to-run variance is very tight on this rig: TTFT sd ≤ 15 ms, TPOT sd ≤ 0.02 ms — well under the 1% iGPU bench noise floor. Every Δ below is far outside the 5×pooled-stddev threshold.

Case	Kernel	TTFT mean ± sd (ms)	TPOT mean ± sd (ms)
Qwen3-4B-quantized.w4a16 in=128/out=128	RDNA3	141.6 ± 0.7	14.05 ± 0.01
Qwen3-4B-quantized.w4a16 in=128/out=128	Hybrid	68.0 ± 0.0	13.57 ± 0.01
Qwen3-4B-quantized.w4a16 in=1920/out=128	RDNA3	1930.7 ± 14.7	15.62 ± 0.02
Qwen3-4B-quantized.w4a16 in=1920/out=128	Hybrid	759.3 ± 1.9	15.13 ± 0.00
gemma-3-4b-it-quantized.w4a16 in=128/out=128	RDNA3	137.3 ± 0.8	16.72 ± 0.01
gemma-3-4b-it-quantized.w4a16 in=128/out=128	Hybrid	65.7 ± 0.5	15.47 ± 0.00
gemma-3-4b-it-quantized.w4a16 in=1920/out=128	RDNA3	1694.6 ± 1.5	25.35 ± 0.01
gemma-3-4b-it-quantized.w4a16 in=1920/out=128	Hybrid	724.3 ± 0.9	24.07 ± 0.02

(Note: gemma-3-4b's TPOT growth with context length is attention-bound, not GEMM-bound — gemma-3 has head_dim=256 and sliding_window=1024, both of which disable the ROCm custom paged-attn fast path. Both kernels are affected equally; the W4A16 kernel does the same work per decode step regardless of context.)

Full 10-case sweep (5 models × 2 shapes, 1 rep each)

This shows what happens on the larger model suite, including AWQ and asymmetric compressed-tensors models. PR 41394's RDNA3 kernel only supports symmetric uint4b8, so the asymmetric models fall through to TritonW4A16 (correctly, per can_implement()). The hybrid kernel accepts both symmetric and asymmetric.

Model	in/out	RDNA3 TTFT / TPOT (ms)	Kernel	Hybrid TTFT / TPOT (ms)	ΔTTFT	ΔTPOT
RedHatAI/Qwen3-4B-quantized.w4a16 (sym, g=128)	128/128	138 / 13.98	RDNA3 → Hybrid	67 / 13.55	−51%	−3%
RedHatAI/Qwen3-4B-quantized.w4a16 (sym, g=128)	1920/128	1894 / 15.56	RDNA3 → Hybrid	759 / 15.11	−60%	−3%
Qwen/Qwen2.5-7B-Instruct-AWQ (asym, g=128)	128/128	119 / 83.48	Triton (fallback) → Hybrid	98 / 20.92	−18%	−75%
Qwen/Qwen2.5-7B-Instruct-AWQ (asym, g=128)	1920/128	2014 / 85.37	Triton (fallback) → Hybrid	1183 / 21.54	−41%	−75%
RedHatAI/gemma-3-4b-it-quantized.w4a16 (sym, g=128)	128/128	137 / 16.74	RDNA3 → Hybrid	65 / 15.48	−52%	−8%
RedHatAI/gemma-3-4b-it-quantized.w4a16 (sym, g=128)	1920/128	1680 / 25.24	RDNA3 → Hybrid	722 / 24.03	−57%	−5%
RedHatAI/Qwen3-8B-quantized.w4a16 (asym, g=128)	128/128	146 / 112.45	Triton (fallback) → Hybrid	139 / 24.46	−5%	−78%
RedHatAI/Qwen3-8B-quantized.w4a16 (asym, g=128)	1920/128	3204 / 115.20	Triton (fallback) → Hybrid	1533 / 26.08	−52%	−77%
trymirai/SmolLM2-1.7B-Instruct-AWQ (asym, g=32)	128/128	38 / 22.71	Triton (fallback) → Hybrid	47 / 7.29	+24%	−68%
trymirai/SmolLM2-1.7B-Instruct-AWQ (asym, g=32)	1920/128	775 / 25.88	Triton (fallback) → Hybrid	417 / 10.54	−46%	−59%

Notes

• None of this is a knock on the PR — your gfx1100 numbers look great and that's the hardware you're targeting. The WMMA prefill kernel and the M ≥ 32 double-buffered LDS variant are clearly designed around the 96-CU GDDR6 part. On 40-CU LPDDR5X they spend more cycles waiting on memory than on compute, and prefill in particular shows a 2× gap.
• For the asymmetric AWQ / CT models, this PR correctly falls through to TritonW4A16 (per can_implement()) — no regression vs the prior state, just no acceleration. The hybrid kernel covers those too, which is most of the gap on those rows above.

Maybe you want to try running the HybridW4A16 from my PR on your gfx1100 and check how performance compares for you?

[JartX]

@mgehre-amd

Hmm, I'm looking at your PR code, some very interesting things that have given me an idea, and the gap before seeing it is real. Wow, what a pair we've been put together, hahahaha 👿