[donghanbj]

Summary

Single-token GQA decode gets a split-K kernel for the cases the fused fp16 flash-decode path does not template (head_dim 512 and KV-sharing layers), and decode over an in-place KV cache no longer reads seqlens_k back to the host. Gemma-4 E4B decode is 5-6% faster on gfx1151.

Related issue or design

None.

Why

• The fused fp16 flash-decode path does not template head_dim 512 or zero-length new KV (KV-sharing layers that re-read an earlier layer's cache), so those single-token decodes fell back to the multi-kernel hipBLASLt pipeline (Steps 8-10).
• Every single-token decode read seqlens_k back to the host (D2H + stream sync) in the first GQA layer of each forward pass, although with a shared past/present buffer the append, RoPE and attention kernels can take the length from the device.

What

• hip_gqa_decode_f32: single-query split-K decode for head_dim 128/256/512 (1, 2, 4 or 8 query heads per KV head; fused Q RoPE only for head_dim >= 256), fp32 or fp16 Q/KV/out, fused non-interleaved Q RoPE, last-block merge. gqa_forward_hipblaslt routes eligible decodes here instead of Steps 8-10; other geometries keep the existing pipeline.
• Device-length mode: each batch derives its key range, sliding-window start and RoPE position from seqlens_k on the device. The host only plans the split grid from an estimate (RuntimeState::seqlens_hint_total: set by every real read, advanced by sq once per forward pass). The last split runs to the end of the range, so any estimate is correct and an accurate one is fastest. Empty splits skip the q setup.
• The fused path skips the seqlens_k readback for in-place decode and gives the tuning lookup the same estimate.

Test plan

• Built hipgpu.dll and the custom kernels, rebuilt hip-mlir-opt / hip-compiler: MorphizenMLIRLitTests 651 passed, 9 unsupported, 0 failed.
• Gemma-4 E4B AWQ uint4 (split fp16 export) end to end through OGA on gfx1151, 128 generated tokens, 3 iterations per prompt, 0 subgraphs falling back to CPU.
• Mistral-7B and Llama-3.1-8B (CI configs): greedy output tokens identical to main.

Performance

Decode TPS, Gemma-4 E4B AWQ uint4, gfx1151 (Windows), 128 generated tokens, mean of 3 iterations:

Prompt tokens	main	this PR	delta	delta with HIPDNN_EP_MATMUL_DP4A=0
388	51.63	54.44	+5.4%	+5.1%
2279	50.08	52.99	+5.8%	+5.7%
16375	43.43	46.17	+6.3%	+6.7%

Prefill / TTFT unchanged (decode-only change).

Notes for reviewers

• Measured on main at d5979f7 plus a local fix for the Gemma-4 shape-inference regression reported in fix(compiler): refine incomplete shapes, unused Split outputs, and unit-slice copies #1306; without it 8 Gemma-4 subgraphs fall back to the CPU EP on current main.
• With the default DP4A GEMV path the 388- and 2279-token outputs diverge late (about word 75 and word 29) and stay coherent; with HIPDNN_EP_MATMUL_DP4A=0 the output is identical to main. The split-K kernel accumulates the attention in a different order than the hipBLASLt pipeline.
• HIPDNN_EP_GQA_F32_DECODE=0 disables the new kernel.

Checklist

• The change is focused, or links a design/series explaining its scope.
• Relevant tests were added or updated and the results are documented.
• User-facing or design documentation was updated when needed.
• Substantial AI assistance is disclosed, and I reviewed and understand the result.

AI assistance (Cursor, Claude Opus 5.5) was used for profiling analysis, the kernel and dispatch implementation, and the benchmark runs. I reviewed the changes and validated them with the lit suite and the end-to-end runs above.

[github-actions]

Thanks for opening a PR!

This project follows LLVM's incremental-development and AI-tool-use
guidance. See CONTRIBUTING.md
for the project workflow.

Before requesting review, please check that:

1. The change is focused. Substantial work links the relevant issue
   or design discussion.
2. The PR documents relevant test results and updates affected
   documentation.
3. If AI tools provided substantial assistance, the description
   explains what was assisted and how it was validated, and commit
   trailers identify the tool. The contributor has reviewed and
   understands the result.

Reviewers are assigned through
CODEOWNERS where ownership
is configured.

[github-actions]

L2 Accuracy Results (EP vs CPU)

Model	Combined L2	Total Elems	Skipped NaN/Inf
conv_test_hybrid	1.12256E-07	64	0
GroupQueryAttention_seq256	59.2561	2621440	0
MatMulNBits_o_seq128	259.905	368640	0
QMoE_seq128	24.2583	368640	0

Threshold: 0.01 | Run: 8055 - Commit: 354c96d

[github-actions]

Binary sizes

Baseline: run 40617608668 - cf8ff09 (dll-sizes)

File	This (MB)	Main (MB)	Delta (MB)
hipgpu.dll	59.34	59.34	0.00
custom_kernels_gfx1170.dll	5.67	5.55	+0.12
custom_kernels_gfx11-generic.dll	5.91	5.79	+0.12

MorphiZen EP Performance Results

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	8.77	7.46	277	3	1158
GroupQueryAttention_seq128	11422.40	2.53303	1	5	229
matmul_down_seq128	510.24	3.28	119	3	274

EPContext Export Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	8.83	48.45	252	3	15505

EPContext Import Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	8.83	9.97	253	3	15640

OGA Benchmark Results

Model	Warmup	Reps	Prompt Len	Gen Tokens	TTFT (ms)	TPS	Peak Mem (GB)	GPU Mem (GB)
gpt-oss-20b-webgpu-int4-rtn-block-32	1	5	128	128	94.1	75.3	1.25	12.24
Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	1	5	128	128	118.1	47.5	1.15	5.23

OGA Wheel Smoke (Python benchmark_e2e.py)

Model	TTFT (ms)	TPS
Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	117	48.1

VLM Benchmark Results

Model	TTFT (ms)	TPS	Output
gemma-4-E4B-it-fp16-int4-kquant-prune-logits	4029.3	7.3	This is a high-quality, outdoor photograph featuring a beautiful, fluffy white dog, likely a West Highland White Terrier or a si
Qwen3.8-27B-fp16-ve-fp16-int4-k_quant-gs128-text-dml-onnx	4906.7	14.1	The user wants a detailed description of the image. Let me carefully observe all elements before writing. Subject: A white do

TTFT / TPS vs main

Baseline: run 40617608668 - cf8ff09 | Gate: TTFT +20% / TPS -20%

Suite	Model	TTFT (ms)	Main	Delta	TPS	Main	Delta
oga	gpt-oss-20b-webgpu-int4-rtn-block-32	94.1	92.4	+1.8%	75.3	75.0	+0.4%
oga	Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	118.1	119.4	-1.1%	47.5	47.7	-0.4%
vlm	gemma-4-E4B-it-fp16-int4-kquant-prune-logits	4029.3	4096.2	-1.6%	7.3	7.3	-0.3%
vlm	Qwen3.8-27B-fp16-ve-fp16-int4-k_quant-gs128-text-dml-onnx	4906.7	4928.4	-0.4%	14.1	14.0	+0.4%

Run: 8055 - Commit: 354c96d

[github-actions]

Kernel Unit Test Results

Op	Leaf	Arch	Cases	Failed	Suite (ms)	Verdict
matmul_nbits	Xfp16_Wu2_Yfp16	gfx1151	12	0	6510	PASS
matmul_nbits	Xfp16_Wu3_Yfp16	gfx1151	12	0	6858.4	PASS
matmul_nbits	Xfp16_Wu4_Yfp16	gfx1151	12	0	4379.8	PASS
matmul_nbits	Xfp16_Wi8_Yfp16	gfx1151	12	0	5810.7	PASS
matmul_nbits	Xfp32_Wu2_Yfp32	gfx1151	12	0	6508.9	PASS
matmul_nbits	Xfp32_Wu3_Yfp32	gfx1151	12	0	6883	PASS
matmul_nbits	Xfp32_Wu4_Yfp32	gfx1151	12	0	4329.4	PASS
matmul_nbits	Xfp32_Wi8_Yfp32	gfx1151	12	0	5798.3	PASS
gqa/decode	Xfp16_Wfp16_Yfp16	gfx1151	17	0	3323.6	PASS
gqa/decode	Xfp16_Wi8_Yfp16	gfx1151	13	0	2483.1	PASS
gqa/prefill	Xfp16_Wfp16_Yfp16	gfx1151	26	0	5213.5	PASS
gqa/prefill	Xfp16_Wi8_Yfp16	gfx1151	9	0	3753.3	PASS
qmoe_ragged_mm	Xfp16_Wu4_Yfp16	gfx1151	40	0	1640	PASS
gather_block_quantized	Xi64_Wu4_Yfp16	gfx1151	32	0	323.7	PASS
gather_block_quantized	Xi64_Wu8_Yfp16	gfx1151	32	0	379.3	PASS
gather_block_quantized	Xi64_Wfp8_Yfp16	gfx1151	24	0	270.6	PASS

Leaves: 16 (0 failed) | Cases: 289 (0 failed)

Mode: lookup | Coverage: 1 | Run: 8055 - Commit: 354c96d

[donghanbj]

/run-test models=all node=xsjnucstrhalo09

[amd-yanjunz]

ROCm GenAI CI finished with result: FAILURE.

• PR: perf(gqa): single-query split-K decode with device-side sequence length #1311
• Source comment: models=all node=xsjnucstrhalo09
• TEST_MODE: performance,accuracy
• MODELZOO: LLM_ROCm_dynamic
• TEST_MODELS: (all)
• FORBID_MODELS: (none)
• NODE: xsjnucstrhalo09
• CI_OPTIONS: (none)
• Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2815/
• Main baseline: unavailable (the package base commit could not be resolved: the selected package run ID was not found in Jenkins)
• NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6ae21f5b5f5e4a2dc2f6b2ec
• Compare with latest daily baseline (fallback): http://xsjnts.xilinx.com/ROCm/compare?base_url=http%3A%2F%2Fxsjnts.xilinx.com%2FROCm%2Fhipdnn-detail%3Fids%3D6ae15c276ac3b9cf4901cc52%26collection%3DHIP-EP-Dynamic&base_job=&base=&curr_url=&curr_job=daily_hip_ep_genai&curr=2815