perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix by penglu12-quark · Pull Request #1264 · ROCm/hip-ep

penglu12-quark · 2026-10-20T04:09:44Z

Summary

Speeds up the speculative-decoding verify pass (2-8 tokens per target call) and greedy argmax on gfx1151, and fixes TopK largest=0 / sorted=0 being dropped during conversion. Four areas:

QMoE: verify bursts of up to 8 tokens take the fused decode kernels instead of the grouped path; the dp4a kernels serve all token*k slots in one set of launches.
LinearAttention: decode and bursts of up to 16 steps run in one launch with the state in registers, split into 32-column tiles across blocks, reading past_state directly (no past->present copy).
fix(onnx-to-hip): onnx.TopK declares largest/sorted as SI64Attr, but the conversion read them as BoolAttr, so they were always true and every smallest-k TopK returned the largest values.
TopK: block-parallel last-axis kernels for k=1 (argmax) and 2 <= k <= 32, with a chunked two-pass path for rows longer than 8192. The two-pass path also takes 32 < k <= 256 when no thread is handed more than 32 candidates in either pass, which covers the k = 128 candidate head of EAGLE-4 drafts over a 248320 vocab.

Related issue or design

None. The QMoE burst dispatch ports #1067.

Why

Speculative verify on MoE / hybrid targets (Qwen3.6-35B-A3B int4) is dominated by fixed per-call costs at small widths. The grouped QMoE path pads every expert to WMMA tiles, the LinearAttention decode path copies state and launches per step, and the generic TopK scans a 248k vocab row with one thread (~31 ms per greedy argmax).

What

New knobs, each documented at its definition in lib/Runtime/debug_log.h. Each one defaults to the new behaviour and can restore the old path for A/B:
- HIPDNN_EP_QMOE_FUSED_MAX_TOKENS (default 8; 1 restores decode-only fused)
- HIPDNN_EP_QMOE_BURST (default on; 0 restores the per-token dp4a loop)
- HIPDNN_EP_LA_BURST_MAX_TOKENS (default 16; 0 disables)
- HIPDNN_EP_LA_COLS (0 disables the column-split kernel)
hip_qmoe_decode_fused_dp4a gains a num_tokens argument; new hip_linear_attention_burst entry point.
kernel_resources_*.json: the qmoe_decode_fc1_dp4a_kernel baseline and allow-list entries move to the new mangled name, with values unchanged.

Test plan

On gfx1151 (Strix Halo, Radeon 8060S), Windows, Release, branch on ffeed9de (since rebased onto 991b9859, a CI-only ORT update touching none of these files):

pytest tests/test_qmoe.py: 11 passed. Adds 2, 5, 8 (fused) and 9 (grouped) tokens. The fused cases, including the 1-token case that takes the same dp4a path on main, exceeded the fixed atol=0.1. The default W4A8 path quantizes activations to int8, an error that scales with |output| (~100 at these shapes; max diff 0.49-0.75, cosine >= 0.99996). Fused cases are now bounded at 2% of max |output| with cosine >= 0.9999 (tightened from 0.99); grouped cases keep atol=0.1.
pytest tests/test_linear_attention.py: 21 passed (adds seq_len 2, 5, 16).
pytest tests/test_topk.py: 13 passed (on dd010ba3). Adds k=1 over 32000, k=16 over 1000, k=32 over 20000 (two-pass), and k=128 over 248320 and k=256 over 100000 (two-pass with k > 32), in both directions. test_topk_axis1[largest=0] fails on main because of the conversion bug.
lit: Conversion/onnx-to-hip (149 tests) plus the TopK hip-to-llvm and e2e tests and LinearAttention hip-to-llvm, all passed. New lit case for largest = 0, sorted = 0.
pre-commit run --from-ref origin/main --to-ref HEAD: all hooks pass.
check_kernel_resources.py --sources top_k_kernel,qmoe_kernel,linear_attention_kernel: the local SDK is not the baseline compiler, so main itself reports 11 errors / 14 warnings. Run on main and on this branch with the same SDK, the reports are identical except for the renamed kernel (same numbers) and one warning: qmoe_decode_reduce_kernel +15% instructions (105 -> 121) for the token loop. The new TopK and LinearAttention kernels use no scratch.

Performance

Strix Halo, Qwen3.6-35B-A3B int4, measured per commit (details in each commit message):

Metric	Before	After
QMoE target call T(w=2)	24.46 ms (grouped)	20.73 ms
QMoE target call T(w=4)	29.95 ms	25.94 ms
QMoE target call T(w=8)	35.97 ms	33.15 ms
30 gated-delta LA layers, w=1 / 4 / 8	1.77 / 2.53 / 2.87 ms	1.13 / 1.55 / 1.91 ms
Greedy decode T(1), GPU argmax graph	49.1 ms	16.9 ms
TopK k=16 over 7x248320	2.4 s	0.80 ms

End-to-end speculative decoding: best speedup per drafter

W4A16 ONNX drafters on Strix Halo: Radeon 8060S iGPU (gfx1151), Windows, ONNX Runtime + hip-ep, strict iGPU mode (HIPDNN_EP_STRICT=1, CPU fallback disabled), HIPDNN_EP_MATMUL_DP4A_MAX_M=8. Spec-Bench, 240 prompts (first 40 of each of 6 subtasks), greedy, 128 new tokens, decode time only. Speedup = pooled tok/s over plain decode measured in the same process.

Directory	Speedup	tok/s (draft / plain)
gemma-4-26B-A4B-it/mtp	1.37×	72.60 / 52.98
gemma-4-26B-A4B-it/eagle3	1.17×	61.00 / 51.95
gemma-4-26B-A4B-it/dspark	1.18×	61.41 / 52.10
gemma-4-26B-A4B-it/dflash	1.29×	67.81 / 52.46
gpt-oss-20b/eagle3	1.21×	102.67 / 84.86
gpt-oss-20b/dflash	1.11×	94.68 / 84.96
qwen3.6-35B-A3B/eagle3	1.58×	99.99 / 63.41
qwen3.6-35B-A3B/mtp	1.60×	102.26 / 63.73
qwen3.6-35B-A3B/dspark	1.57×	98.97 / 62.93
qwen3.6-35B-A3B/dflash2	1.94×	122.53 / 63.19
nemotron-3.5-lightning-30B-A3B/mtp	2.37×	15.91 / 6.73
nemotron-3.5-lightning-30B-A3B/eagle3	1.70×	11.40 / 6.73
nemotron-3.5-lightning-30B-A3B/dspark	2.78×	18.70 / 6.73
nemotron-3.5-lightning-30B-A3B/dflash	2.40×	16.11 / 6.73
qwen3.8-27B/eagle4	2.31×	33.28 / 14.38

Bold rows: best drafter for each target. GPT-OSS-20B runs use HIPDNN_EP_QMOE_FUSED_MAX_TOKENS=1 (its large experts favour the grouped path). The Nemotron target export runs all 128 experts per token (no QMoE), so its speedups are not comparable with the QMoE targets. Qwen3.8-27B is a dense hybrid target (48 Gated DeltaNet + 16 GQA layers, AMD int4 ONNX export) with the EAGLE-Corp/Qwen3.8-27B-EAGLE4 drafter (eagle4:5:c:t0.5: k <= 5, cap, τ = 0.5); it needs the k = 128 two-pass TopK above. Its 240-prompt run used this branch at 7316a9da plus that TopK change; on the current head (dd010ba3) a one-prompt check runs end to end in strict mode. Drafter, target graph and harness: amd/Qwen3.8-27B-EAGLE4-hip-ep.

Notes for reviewers

Commit 3 (the TopK conversion fix) is independent and can be split into its own PR if preferred. Commit 4's new largest=0 tests depend on it.
Numeric tests were run with ORT 1.30, which rejects IR 14 models, through a local pytest plugin that caps generated models at IR 13. It is not part of this PR.

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 implemented the kernels, runtime dispatch, conversion fix and tests, ran the measurements and validation listed above, and is credited in the Co-authored-by trailers.

Speculative verify presents 2-8 tokens per MoE layer, which used to take the grouped path (one pass over the union of experts in padded WMMA tiles). Bursts up to HIPDNN_EP_QMOE_FUSED_MAX_TOKENS (default 8) now take the fused decode path instead. The dp4a kernels take num_tokens and run all token*k slots in the same launches (quantize inputs, FC1+SwiGLU, quantize slots, FC2, plus the per-token reduce); bursts autotune separately from decode. The fp16 kernel is iterated a token at a time. HIPDNN_EP_QMOE_BURST=0 restores the per-token dp4a loop and HIPDNN_EP_QMOE_FUSED_MAX_TOKENS=1 the decode-only behaviour, for A/B. Ports the burst dispatch of #1067. Strix Halo, Qwen3.6-35B-A3B int4, target call T(w) in ms: w 2 3 4 5 6 8 loop 21.97 26.09 29.03 32.48 - - burst 20.73 24.04 25.94 28.52 30.53 33.15 grouped 24.46 27.98 29.95 31.69 33.37 35.97 test_qmoe adds 2, 5, 8 (fused) and 9 (grouped) tokens. The fused cases, 1 token included (the same dp4a decode path as on main), exceeded the fixed atol=0.1: the default W4A8 dp4a path quantizes activations to int8, an error that scales with |output| (~100 at the GPT-OSS shapes; max diff 0.49-0.75, cosine >= 0.99996). They are now bounded at 2% of max |output| with cosine >= 0.9999 instead of 0.99; grouped cases keep atol=0.1. Burst and loop produce the same error stats on every case. 11/11 pass on gfx1151. qmoe_decode_fc1_dp4a_kernel gains a num_tokens parameter, so its entries in kernel_resources_*.json (baseline and allow list) move to the new mangled name unchanged; with the same SDK check_kernel_resources.py reports the same resources for it as on main. Co-authored-by: Cursor <cursoragent@cursor.com>

… short bursts Every S column evolves independently under all four update rules, so the (batch, kv_head) state is split into 32-column tiles across blockIdx.y (128 blocks instead of 32 for Qwen3.6's 32 kv heads) and each thread keeps its rows of S in registers: one read and one write of S per launch instead of a past->present copy plus two reads and a write per step. seq_len <= HIPDNN_EP_LA_BURST_MAX_TOKENS (default 16) runs all steps in that launch with S in fp32 between steps (operands staged in LDS, two barriers per step) and reads past_state directly, so the separate state copy is skipped for decode and speculative verify. HIPDNN_EP_LA_COLS=0 or HIPDNN_EP_LA_BURST_MAX_TOKENS=0 restores the old paths. Standalone, 30 gated-delta layers (Hq=16, Hkv=32, d=128, ping-pong state): w 1 2 4 8 old 1.77 2.34 2.53 2.87 ms new 1.13 1.32 1.55 1.91 ms In the full 35B-A3B target the saving is 0.2-0.7 ms per call at w<=4. test_linear_attention passes, now also at seq_len 2, 5 and 16. Co-authored-by: Cursor <cursoragent@cursor.com>

onnx.TopK declares largest and sorted as SI64Attr, but the conversion read them with getAttrOfType<BoolAttr>, which never matches, so both stayed at their default of true: every smallest-k TopK returned the largest values (test_topk_axis1 with largest=0 failed on main). Read them as IntegerAttr, as axis already is, and add a lit case that checks largest = false and sorted = false reach hip.top_k. Co-authored-by: Cursor <cursoragent@cursor.com>

The generic TopK assigns one thread per row and scans the axis serially. For a greedy-decode argmax over a 248k vocab that is ~31 ms per call on gfx1151, more than the rest of the 35B-A3B forward, and 2 <= k <= 32 was as slow (2.4 s for k=16 over 7x248320). - k=1: one 1024-thread block per row with a shared-memory (value, index) reduction; ties keep the lowest index, matching numpy argmax. - 2 <= k <= 32: one 256-thread block per row with per-thread register top-K lists; rows longer than 8192 take a chunked first pass into a cached scratch buffer plus a per-row merge. Values and lowest-index tie order match a stable sort. Measured on Strix Halo (Qwen3.6-35B-A3B int4, GPU-argmax graph): T(1) 49.1 -> 16.9 ms, vs 18.1 ms for host argmax. Token-identical to host argmax on 10 Spec-Bench prompts x 128 tokens. k=16 over 7x248320: 0.80 ms; k=32 over 8x32000: 0.37 ms. test_topk covers each new path in both directions: k=1 over 32000, k=16 over 1000 (one pass) and k=32 over 20000 (two pass). 9/9 pass on gfx1151. Co-authored-by: Cursor <cursoragent@cursor.com>

github-actions · 2026-10-20T04:10:13Z

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:

The change is focused. Substantial work links the relevant issue
or design discussion.
The PR documents relevant test results and updates affected
documentation.
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.

penglu12-quark · 2026-10-20T04:55:32Z

/run-test models=all

amd-yanjunz · 2026-10-20T04:57:24Z

ROCm GenAI CI finished with result: ABORTED.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
Source comment: models=all
TEST_MODE: performance,accuracy
MODELZOO: LLM_ROCm_dynamic
TEST_MODELS: (all)
FORBID_MODELS: (none)
NODE: (GenAI job default)
CI_OPTIONS: (none)
Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2507/
Main baseline commit: 991b985
Main baseline build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2504/
Main baseline: not usable, the baseline build has not finished yet
Package build run: https://github.com/ROCm/hip-ep/actions/runs/39368091230
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6ad732350776acf7f561f071
Compare with latest daily baseline (fallback): http://xsjnts.xilinx.com/ROCm/compare?base_url=http%3A%2F%2Fxsjnts.xilinx.com%2FROCm%2Fhipdnn-detail%3Fids%3D6ad59f43c497c7d6b761eea8%26collection%3DHIP-EP-Dynamic&base_job=&base=&curr_url=&curr_job=daily_hip_ep_genai&curr=2507

github-actions · 2026-10-20T05:52:15Z

L2 Accuracy Results (EP vs CPU)

Model	Combined L2	Total Elems
conv_test_hybrid	1.12256E-07	64
GroupQueryAttention_seq256	59.2561	2621440
MatMulNBits_o_seq128	259.905	368640
QMoE_seq128	24.2583	368640

Threshold: 0.01 | Run: 8092 - Commit: cf7e1fd

github-actions · 2026-10-20T05:52:18Z

Binary sizes

Baseline: run 40801212116 - d3df783 (dll-sizes)

File	This (MB)	Main (MB)	Delta (MB)
hipgpu.dll	59.34	59.34	+0.01
custom_kernels_gfx1170.dll	5.60	5.55	+0.05
custom_kernels_gfx11-generic.dll	5.85	5.79	+0.05

MorphiZen EP Performance Results

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	8.78	6.91	247	3	1158
GroupQueryAttention_seq128	11650.90	2.44631	1	5	229
matmul_down_seq128	501.35	3.18	125	3	275

EPContext Export Performance

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

EPContext Import Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	8.84	9.91	248	3	15639

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	92.2	75.0	1.25	12.27
Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	1	5	128	128	118.7	47.6	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	118	48.1

VLM Benchmark Results

Model	TTFT (ms)	TPS	Output
gemma-4-E4B-it-fp16-int4-kquant-prune-logits	3927.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	4928.3	13.9	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 40801212116 - d3df783 | Gate: TTFT +20% / TPS -20%

Suite	Model	TTFT (ms)	Main	Delta	TPS	Main	Delta
oga	gpt-oss-20b-webgpu-int4-rtn-block-32	92.2	93.7	-1.7%	75.0	74.9	+0.1%
oga	Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	118.7	118.8	-0.1%	47.6	47.7	-0.2%
vlm	gemma-4-E4B-it-fp16-int4-kquant-prune-logits	3927.3	4005.3	-1.9%	7.3	8.0	-8.8%
vlm	Qwen3.8-27B-fp16-ve-fp16-int4-k_quant-gs128-text-dml-onnx	4928.3	4880.6	+1.0%	13.9	14.0	-0.9%

Run: 8092 - Commit: cf7e1fd

Resolve test_qmoe.py: keep the fused-path tolerance helper for the skewed-routing test and add main's router_weights test unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

penglu12-quark · 2026-10-21T23:48:04Z

/run-test models=all

amd-yanjunz · 2026-10-21T23:52:33Z

ROCm GenAI CI finished with result: UNSTABLE.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
Source comment: models=all
TEST_MODE: performance,accuracy
MODELZOO: LLM_ROCm_dynamic
TEST_MODELS: (all)
FORBID_MODELS: (none)
NODE: (GenAI job default)
CI_OPTIONS: (none)
Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2573/
Main baseline commit: 9d38afb
Main baseline build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2575/
Main baseline: not usable, the baseline build has not finished yet
Package build run: https://github.com/ROCm/hip-ep/actions/runs/39695507121
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6ad9a1a4f27f47880d1b66a0
Compare with latest daily baseline (fallback): http://xsjnts.xilinx.com/ROCm/compare?base_url=http%3A%2F%2Fxsjnts.xilinx.com%2FROCm%2Fhipdnn-detail%3Fids%3D6ad7d74ee9ec6e3e0461ed9d%26collection%3DHIP-EP-Dynamic&base_job=&base=&curr_url=&curr_job=daily_hip_ep_genai&curr=2573

github-actions · 2026-10-22T03:05:41Z

Kernel Unit Test Results

Op	Leaf	Arch	Cases	Suite (ms)	Verdict
matmul_nbits	Xfp16_Wu2_Yfp16	gfx1151	12	6580.8	PASS
matmul_nbits	Xfp16_Wu3_Yfp16	gfx1151	12	6792.2	PASS
matmul_nbits	Xfp16_Wu4_Yfp16	gfx1151	12	4217.5	PASS
matmul_nbits	Xfp16_Wi8_Yfp16	gfx1151	12	5641.3	PASS
matmul_nbits	Xfp32_Wu2_Yfp32	gfx1151	12	6637.5	PASS
matmul_nbits	Xfp32_Wu3_Yfp32	gfx1151	12	6772.5	PASS
matmul_nbits	Xfp32_Wu4_Yfp32	gfx1151	12	4297.2	PASS
matmul_nbits	Xfp32_Wi8_Yfp32	gfx1151	12	5753.4	PASS
gqa/decode	Xfp16_Wfp16_Yfp16	gfx1151	17	3288.7	PASS
gqa/decode	Xfp16_Wi8_Yfp16	gfx1151	13	2550.9	PASS
gqa/prefill	Xfp16_Wfp16_Yfp16	gfx1151	26	5198	PASS
gqa/prefill	Xfp16_Wi8_Yfp16	gfx1151	9	3775.2	PASS
qmoe_ragged_mm	Xfp16_Wu4_Yfp16	gfx1151	40	1625	PASS
gather_block_quantized	Xi64_Wu4_Yfp16	gfx1151	32	382.2	PASS
gather_block_quantized	Xi64_Wu8_Yfp16	gfx1151	32	286.2	PASS
gather_block_quantized	Xi64_Wfp8_Yfp16	gfx1151	24	256.3	PASS

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

Mode: lookup | Coverage: 1 | Run: 8092 - Commit: cf7e1fd

Replace the global mutex-guarded scratch buffer with the RuntimeState workspace, following the qgemm pattern. Add hip_top_k_workspace_bytes(), pass workspace/workspace_bytes into hip_top_k, size it in wrap_top_k after K is read, and fall back to the one-pass topk_small_last_axis_kernel when no workspace is available. Co-authored-by: Cursor <cursoragent@cursor.com>

Clear stale HIP error state before launching, then log hipGetErrorString and return -1 on launch failure. Co-authored-by: Cursor <cursoragent@cursor.com>

Update TN=4/TN=8 baseline keys to the new mangled signature (extra int parameter). Co-authored-by: Cursor <cursoragent@cursor.com>

amd-yanjunz · 2026-10-22T13:42:37Z

ROCm GenAI CI finished with result: UNSTABLE.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
Source comment: models=all
TEST_MODE: performance,accuracy
MODELZOO: LLM_ROCm_dynamic
TEST_MODELS: (all)
FORBID_MODELS: (none)
NODE: (GenAI job default)
CI_OPTIONS: (none)
Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2591/
Main baseline commit: 76c6664
Main baseline: not usable, no baseline result for this commit yet
Package build run: https://github.com/ROCm/hip-ep/actions/runs/39788343808
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6ada63d490a7a5eec37aa868
Compare with latest daily baseline (fallback): http://xsjnts.xilinx.com/ROCm/compare?base_url=http%3A%2F%2Fxsjnts.xilinx.com%2FROCm%2Fhipdnn-detail%3Fids%3D6ad9baedf27f47880d1b66a8%26collection%3DHIP-EP-Dynamic&base_job=&base=&curr_url=&curr_job=daily_hip_ep_genai&curr=2591

qianglin-amd · 2026-10-26T04:07:55Z

/run-test models=all

amd-yanjunz · 2026-10-26T04:12:32Z

ROCm GenAI CI finished with result: UNSTABLE.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
Source comment: models=all
TEST_MODE: performance,accuracy
MODELZOO: LLM_ROCm_dynamic
TEST_MODELS: (all)
FORBID_MODELS: (none)
NODE: (GenAI job default)
CI_OPTIONS: (none)
Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2707/
Main baseline commit: a8398e2
Main baseline: not usable, the baseline build is queued
Package build run: https://github.com/ROCm/hip-ep/actions/runs/40269413951
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6adf07f56cb9d0c74b7ab37d
Compare with latest daily baseline (fallback): http://xsjnts.xilinx.com/ROCm/compare?base_url=http%3A%2F%2Fxsjnts.xilinx.com%2FROCm%2Fhipdnn-detail%3Fids%3D6ade45bfa9af1b06b77aacf3%26collection%3DHIP-EP-Dynamic&base_job=&base=&curr_url=&curr_job=daily_hip_ep_genai&curr=2707

Resolve kernel_resources_{gfx11-generic,gfx1170}.json: take main's regenerated baselines and reapply this branch's rename of the qmoe_decode_fc1_dp4a_kernel keys to the new signature (extra int parameter); apply the same rename to main's new kernel_resources_gfx1151.json. Co-authored-by: Cursor <cursoragent@cursor.com>

Resolve qmoe.cpp: keep this branch's fused burst gate (num_tokens <= hipdnn_ep_qmoe_fused_max_tokens()) and add main's swiglu requirement, so relu2 experts skip the SwiGLU fused path. Resolve test_qmoe.py imports: keep both os and dataclass. Co-authored-by: Cursor <cursoragent@cursor.com>

penglu12-quark · 2026-10-26T12:49:22Z

/run-test models=all

amd-yanjunz · 2026-10-26T12:52:40Z

ROCm GenAI CI finished with result: UNSTABLE.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
Source comment: models=all
TEST_MODE: performance,accuracy
MODELZOO: LLM_ROCm_dynamic
TEST_MODELS: (all)
FORBID_MODELS: (none)
NODE: (GenAI job default)
CI_OPTIONS: (none)
Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2721/
Main baseline commit: 7579ebf
Main baseline build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2715/
Package build run: https://github.com/ROCm/hip-ep/actions/runs/40325379734
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6adf983ba9af1b06b77aaee9
Compare with main baseline 7579ebf: http://xsjnts.xilinx.com/ROCm/compare?base_url=http%3A%2F%2Fxsjnts.xilinx.com%2FROCm%2Fhipdnn-detail%3Fdatabase%3Ddevelopment%26collection%3Ddaily_hip_ep_genai%26id%3D6adf7ab8848fa63c277aafa3&base_job=&base=&curr_url=&curr_job=daily_hip_ep_genai&curr=2721

penglu12-quark · 2026-10-27T09:47:37Z

/run-test models=all

amd-yanjunz · 2026-10-27T09:52:42Z

ROCm GenAI CI finished with result: UNSTABLE.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
Source comment: models=all
TEST_MODE: performance,accuracy
MODELZOO: LLM_ROCm_dynamic
TEST_MODELS: (all)
FORBID_MODELS: (none)
NODE: (GenAI job default)
CI_OPTIONS: (none)
Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2771/
Main baseline commit: 84c0443
Main baseline build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2773/
Main baseline: not usable, the baseline build has not finished yet
Package build run: https://github.com/ROCm/hip-ep/actions/runs/40468083919
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6ae093e310cc69816601cba3
Compare with latest daily baseline (fallback): http://xsjnts.xilinx.com/ROCm/compare?base_url=http%3A%2F%2Fxsjnts.xilinx.com%2FROCm%2Fhipdnn-detail%3Fids%3D6adffc53a9af1b06b77aaf3b%26collection%3DHIP-EP-Dynamic&base_job=&base=&curr_url=&curr_job=daily_hip_ep_genai&curr=2771

penglu12-quark · 2026-10-27T13:30:45Z

/run-test models=all

amd-yanjunz · 2026-10-27T13:32:47Z

ROCm GenAI CI finished with result: UNSTABLE.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
Source comment: models=all
TEST_MODE: performance,accuracy
MODELZOO: LLM_ROCm_dynamic
TEST_MODELS: (all)
FORBID_MODELS: (none)
NODE: (GenAI job default)
CI_OPTIONS: (none)
Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2778/
Main baseline commit: 3c569ca
Main baseline build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2791/
Main baseline: not usable, the baseline build has not finished yet
Package build run: https://github.com/ROCm/hip-ep/actions/runs/40512303359
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6ae1078f34dd0c11b501cb64
Compare with latest daily baseline (fallback): http://xsjnts.xilinx.com/ROCm/compare?base_url=http%3A%2F%2Fxsjnts.xilinx.com%2FROCm%2Fhipdnn-detail%3Fids%3D6adffc53a9af1b06b77aaf3b%26collection%3DHIP-EP-Dynamic&base_job=&base=&curr_url=&curr_job=daily_hip_ep_genai&curr=2778

…uncated The two-pass TopK accepted only k <= 32, so larger k over a vocab row fell back to the one-thread-per-row kernel. The two-pass kernels stay exact for larger k as long as no thread is handed more than kTopKSmallMax candidates in either pass (chunk pass: kTopKChunk / block, merge pass: chunks * k / block). Allow 32 < k <= 256 under that bound. This covers the low-rank candidate head of EAGLE-4 drafts (k = 128 over a 248320 vocab: 61 chunks, 31 candidates per thread in the merge). test_topk: adds k=128 over 248320 and k=256 over 100000, both directions. Co-authored-by: Cursor <cursoragent@cursor.com>

penglu12-quark · 2026-10-28T00:26:13Z

/run-test models=all

amd-yanjunz · 2026-10-28T00:27:48Z

ROCm GenAI CI finished with result: UNSTABLE.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
Source comment: models=all
TEST_MODE: performance,accuracy
MODELZOO: LLM_ROCm_dynamic
TEST_MODELS: (all)
FORBID_MODELS: (none)
NODE: (GenAI job default)
CI_OPTIONS: (none)
Jenkins build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2802/
Main baseline commit: 3c569ca
Main baseline build: http://xsjvitisaijenkins:8080/job/daily_hip_ep_genai/2791/
Main baseline: not usable, the baseline build has not finished yet
Package build run: https://github.com/ROCm/hip-ep/actions/runs/40605968670
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6ae1893002a637697001cd2a
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=2802

penglu12-quark · 2026-10-28T06:22:55Z

/run-test models=all node=xsjnucstrhalo09

amd-yanjunz · 2026-10-28T06:27:43Z

ROCm GenAI CI finished with result: UNSTABLE.

PR: perf: speculative-decode verify on gfx1151 (QMoE bursts, LinearAttention burst, TopK) + TopK largest=0 fix #1264
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/2819/
Main baseline commit: cf8ff09
Main baseline: not usable, the baseline uses a different NODE
Package build run: https://github.com/ROCm/hip-ep/actions/runs/40642524770
NTS report: http://xsjnts.xilinx.com/ROCm/hipdnn-detail?database=development&collection=daily_hip_ep_genai&id=6ae275ac59b3b76f90f6b40f
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=2819

Conflict in lib/Runtime/real/qmoe.cpp with #1264, which split the fused decode into a whole-burst dp4a launch (HIPDNN_EP_QMOE_BURST) and a per-token loop. Kept main's code and placed the step marks on the new paths: "decode.dp4a" on the burst launch, and one mark before the per-token loop, "decode.dp4a.per_token" or "decode.fp", since a mark names every launch until the next one. Co-Authored-By: Cursor Claude Opus 5.5 <cursoragent@cursor.com> Made-with: Cursor

penglu12-quark and others added 4 commits October 19, 2026 21:55

penglu12-quark requested review from CaymanYang, amd-bqiao, amd-mingw, asparkhi, edelaye, fhanuman, qianglin-amd, sambanerxilinxamd and wcy123 as code owners October 20, 2026 04:09

penglu12-quark self-assigned this Oct 20, 2026

penglu12-quark requested a review from ChaoLi-AMD October 20, 2026 04:53

penglu12-quark added the enhancement New feature or request label Oct 20, 2026

Merge origin/main into perf/moe-spec-decode-strix

a626f98

Resolve test_qmoe.py: keep the fused-path tolerance helper for the skewed-routing test and add main's router_weights test unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

Merge branch 'main' into perf/moe-spec-decode-strix

09f5802

mingyueliuh reviewed Oct 22, 2026

View reviewed changes

Comment thread lib/Runtime/Kernels/hip/top_k_kernel.hip Resolved

Comment thread lib/Runtime/Kernels/tools/kernel_resources_gfx11-generic.json Resolved

Comment thread lib/Runtime/Kernels/hip/linear_attention_kernel.hip Resolved

penglu12-quark and others added 3 commits October 22, 2026 07:32

LinearAttention: check launch errors in launch_la_cols

ad8eaf4

Clear stale HIP error state before launching, then log hipGetErrorString and return -1 on launch failure. Co-authored-by: Cursor <cursoragent@cursor.com>

Kernels: rename qmoe_decode_fc1_dp4a_kernel resource baselines

7316a9d

Update TN=4/TN=8 baseline keys to the new mangled signature (extra int parameter). Co-authored-by: Cursor <cursoragent@cursor.com>

penglu12-quark and others added 2 commits October 25, 2026 22:30

masteryi-0018 mentioned this pull request Oct 26, 2026

fix(qmoe): choose the decode GEMV tiling before computing, not on real calls #1328

Open

4 tasks done

Merge branch 'main' into perf/moe-spec-decode-strix

82786af

penglu12-quark added 3 commits October 26, 2026 20:51

Merge branch 'main' into perf/moe-spec-decode-strix

7056de0

Merge branch 'main' into perf/moe-spec-decode-strix

9d5214e

Merge branch 'main' into perf/moe-spec-decode-strix

6380d1c

Merge branch 'main' into perf/moe-spec-decode-strix

4b68d2a

Merge branch 'main' into perf/moe-spec-decode-strix

49ff975

Merge branch 'main' into perf/moe-spec-decode-strix

cf7e1fd

qianglin-amd approved these changes Oct 29, 2026

View reviewed changes

qianglin-amd merged commit 7926971 into main Oct 29, 2026
11 checks passed

qianglin-amd deleted the perf/moe-spec-decode-strix branch October 29, 2026 13:27

penglu12-quark mentioned this pull request Oct 30, 2026

perf: 2-bit MatMulNBits models end to end on Strix Halo (u2 fp32 I/O, fp32-KV flash prefill, Q/DQ chain fusions) #1318

Draft

3 of 4 tasks

donghanbj mentioned this pull request Oct 30, 2026

feat(profiling): add optional ROCTx operator ranges #949

Draft

Conversation

penglu12-quark commented Oct 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Related issue or design

Why

What

Test plan

Performance

End-to-end speculative decoding: best speedup per drafter

Notes for reviewers

Checklist

Uh oh!

github-actions Bot commented Oct 20, 2026

Uh oh!

penglu12-quark commented Oct 20, 2026

Uh oh!

amd-yanjunz commented Oct 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

github-actions Bot commented Oct 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

L2 Accuracy Results (EP vs CPU)

Uh oh!

github-actions Bot commented Oct 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Binary sizes

MorphiZen EP Performance Results

EPContext Export Performance

EPContext Import Performance

OGA Benchmark Results

OGA Wheel Smoke (Python benchmark_e2e.py)

VLM Benchmark Results

TTFT / TPS vs main

Uh oh!

penglu12-quark commented Oct 21, 2026

Uh oh!

amd-yanjunz commented Oct 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

github-actions Bot commented Oct 22, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Kernel Unit Test Results

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-yanjunz commented Oct 22, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

qianglin-amd commented Oct 26, 2026

Uh oh!

amd-yanjunz commented Oct 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

penglu12-quark commented Oct 26, 2026

Uh oh!

amd-yanjunz commented Oct 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

penglu12-quark commented Oct 27, 2026

Uh oh!

amd-yanjunz commented Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

penglu12-quark commented Oct 27, 2026

Uh oh!

amd-yanjunz commented Oct 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

penglu12-quark commented Oct 28, 2026

Uh oh!

amd-yanjunz commented Oct 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

penglu12-quark commented Oct 28, 2026

Uh oh!

amd-yanjunz commented Oct 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

penglu12-quark commented Oct 20, 2026 •

edited

Loading

amd-yanjunz commented Oct 20, 2026 •

edited

Loading

github-actions Bot commented Oct 20, 2026 •

edited

Loading

github-actions Bot commented Oct 20, 2026 •

edited

Loading

amd-yanjunz commented Oct 21, 2026 •

edited

Loading

github-actions Bot commented Oct 22, 2026 •

edited

Loading

amd-yanjunz commented Oct 22, 2026 •

edited

Loading

amd-yanjunz commented Oct 26, 2026 •

edited

Loading

amd-yanjunz commented Oct 26, 2026 •

edited

Loading

amd-yanjunz commented Oct 27, 2026 •

edited

Loading

amd-yanjunz commented Oct 27, 2026 •

edited

Loading

amd-yanjunz commented Oct 28, 2026 •

edited

Loading

amd-yanjunz commented Oct 28, 2026 •

edited

Loading