ci(windows): target gfx11-7-generic for gfx117x and fix the WMMA conv there by zz002 · Pull Request #1329 · ROCm/hip-ep

zz002 · 2026-10-26T08:50:16Z

Summary

Build the gfx117x custom kernels for gfx11-7-generic, so one DLL covers gfx1170 / gfx1171 / gfx1172. Switch to an SDK that ships gfx1171 hipBLASLt data, and fix the fp16 WMMA conv, which compiled to a trap on gfx117x.

Related issue or design

None

Why

The pinned SDK (ROCm/llvm-project a2529029) defines gfx11-7-generic, so one generic target covers the whole gfx117x family.
The previous SDK has no hipblaslt\library\gfx1171, so hipBLASLt cannot find its Tensile library on gfx1171.
gfx117x is in the gfx11 family but implements only the gfx12-style WMMA. conv_kernel.hip called the gfx11 builtin directly, which falls back to __builtin_trap() on gfx117x, so every fp16 WMMA conv faulted on the GPU. This is what broke the gemma3 vision encoder on gfx1171.

What

HIP_ARCHITECTURES: gfx11-generic;gfx1170 -> gfx11-generic;gfx11-7-generic. The four generic-target tables (genericTargetFor(), kGenericTargets, _lut_arch_members(), _GENERIC_MEMBERS) map gfx1170 / gfx1171 / gfx1172 to gfx11-7-generic.
cmake/deps.txt: SDK -> v0.5.2 rocm-dist-windows-gfx115x-gfx117x-7.16.0-hipblaslt-gfx117x-trimmed, which adds gfx1171 hipBLASLt data; gfx1170 / gfx1171 keep the same five fp16 Tensile stems (built by ci(windows): build gfx1171 in the SDK selfbuild zz002/hip-ep#7).
conv_kernel.hip: conv_wmma_kernel and conv_window_kernel use the WMMA helpers from hip_arch_compat.h (fragment width, K offset, accumulator row, builtin), as gemm_wmma_kernel.hip does. The gfx11 encoding and memory access are unchanged.

Test plan

build-windows green; custom_kernels_gfx11-7-generic.dll and kernel-tests/gfx11-7-generic/ are staged
gpu-test-package contains bin\hipblaslt\library\gfx1171

Notes for reviewers

The DLL size gate reports custom_kernels_gfx11-7-generic.dll +5.58 MB FAIL because it compares by file name and main has no file by that name, so the baseline counts as 0. The DLL replaces main's custom_kernels_gfx1170.dll; hipgpu.dll and custom_kernels_gfx11-generic.dll are unchanged (Δ 0.00).

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.

github-actions · 2026-10-26T08:50:36Z

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.

github-actions · 2026-10-26T13:27:55Z

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: 8143 - Commit: 463fe41

github-actions · 2026-10-26T13:27:58Z

Binary sizes

Baseline: run 40994390384 - 5e4585c (dll-sizes)

File	This (MB)	Main (MB)	Delta (MB)
hipgpu.dll	59.38	59.38	0.00
custom_kernels_gfx11-7-generic.dll	5.65	-	new
custom_kernels_gfx11-generic.dll	5.87	5.87	0.00

MorphiZen EP Performance Results

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	8.92	7.15	248	3	1158
GroupQueryAttention_seq128	11047.60	2.57102	0	5	168
matmul_down_seq128	733.83	3.28	120	3	268

EPContext Export Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	8.98	47.33	253	3	15505

EPContext Import Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	8.97	10.02	248	3	15544

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	89.3	81.3	1.25	12.34
Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	1	5	128	128	118.5	49.5	1.14	5.40

OGA Wheel Smoke (Python benchmark_e2e.py)

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

VLM Benchmark Results

Model	TTFT (ms)	TPS	Output
gemma-4-E4B-it-fp16-int4-kquant-prune-logits	3827.2	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	4855.8	14.6	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 40994390384 - 5e4585c | Gate: TTFT +20% / TPS -20%

Suite	Model	TTFT (ms)	Main	Delta	TPS	Main	Delta
oga	gpt-oss-20b-webgpu-int4-rtn-block-32	89.3	91.0	-1.9%	81.3	82.7	-1.7%
oga	Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	118.5	116.4	+1.9%	49.5	49.6	-0.2%
vlm	gemma-4-E4B-it-fp16-int4-kquant-prune-logits	3827.2	3775.7	+1.4%	7.3	8.0	-7.9%
vlm	Qwen3.8-27B-fp16-ve-fp16-int4-k_quant-gs128-text-dml-onnx	4855.8	4880.8	-0.5%	14.6	14.6	0.0%

Run: 8143 - Commit: 463fe41

github-actions · 2026-10-26T13:35:26Z

Kernel Unit Test Results

Op	Leaf	Arch	Cases	Suite (ms)	Verdict
matmul_nbits	Xfp16_Wu2_Yfp16	gfx1151	12	6480.1	PASS
matmul_nbits	Xfp16_Wu3_Yfp16	gfx1151	12	6788.1	PASS
matmul_nbits	Xfp16_Wu4_Yfp16	gfx1151	12	4197.8	PASS
matmul_nbits	Xfp16_Wi8_Yfp16	gfx1151	12	5616.3	PASS
matmul_nbits	Xfp32_Wu2_Yfp32	gfx1151	12	6503.6	PASS
matmul_nbits	Xfp32_Wu3_Yfp32	gfx1151	12	6651	PASS
matmul_nbits	Xfp32_Wu4_Yfp32	gfx1151	12	4266.1	PASS
matmul_nbits	Xfp32_Wi8_Yfp32	gfx1151	12	5638.4	PASS
gqa/decode	Xfp16_Wfp16_Yfp16	gfx1151	17	3306.1	PASS
gqa/decode	Xfp16_Wi8_Yfp16	gfx1151	13	2358.5	PASS
gqa/prefill	Xfp16_Wfp16_Yfp16	gfx1151	26	5217.4	PASS
gqa/prefill	Xfp16_Wi8_Yfp16	gfx1151	9	3705.8	PASS
qmoe_ragged_mm	Xfp16_Wu4_Yfp16	gfx1151	40	1643.2	PASS
gather_block_quantized	Xi64_Wu4_Yfp16	gfx1151	32	391	PASS
gather_block_quantized	Xi64_Wu8_Yfp16	gfx1151	32	364.5	PASS
gather_block_quantized	Xi64_Wfp8_Yfp16	gfx1151	24	398	PASS

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

Mode: lookup | Coverage: 1 | Run: 8143 - Commit: 463fe41

zz002 · 2026-10-28T08:01:58Z

Test	PR 1329	main	Compare
LLM + VLM, gfx1170	multi_board 558	multi_board 552	compare

Co-Authored-By: Copilot <copilot@github.com>

Point therock_windows at rocm-dist-windows-gfx115x-gfx117x-7.16.0-hipblaslt-gfx117x-trimmed (v0.5.2 asset, built by zz002#7). It adds bin\hipblaslt\library\gfx1171, trimmed to the same HHS/HSS stems as gfx1170; without it hipBLASLt finds no Tensile library on gfx1171. Co-Authored-By: Copilot <copilot@github.com>

gfx117x is numbered in the gfx11 family but implements only the gfx12-style WMMA encoding (see hip_arch_compat.h), so __builtin_amdgcn_wmma_f32_16x16x16_f16_w32 does not exist for gfx11-7-generic. conv_kernel.hip called that builtin directly and fell back to __builtin_trap() without it, so every WMMA conv kernel in custom_kernels_gfx11-7-generic.dll compiled to s_trap and no v_wmma. Every fp16 Conv reaches those kernels, and on Windows the trap hangs the wave until the driver resets the GPU (TDR). gemma3's vision encoder fails this way on gfx1171, with hip_transpose_2d_tiled reporting the first error. Route conv_wmma_kernel and conv_window_kernel through the hip_arch_compat.h helpers, as gemm_wmma_kernel.hip already does: the fragment width, the per-lane K offset, the accumulator row and the builtin follow the target. gfx11 keeps its encoding (K offset 0, row e*2+pair), so its WMMA and memory instructions are unchanged. Tested on gfx1171 (Radeon 940M) with the v0.5.2 SDK: fp16 convs and the gemma3 vision encoder match the CPU EP (relL2 <= 5e-3), the gemma3 VLM benchmark passes with 1100 hipBLASLt calls and no TDR, and the kernel unit tests pass. Co-authored-by: Cursor <cursoragent@cursor.com>

The conv fragment-layout comments still described the gfx11 encoding only (K replicated across lane halves, accumulator row e*2+pair). Point them at loadWmmaFrag / hipdnn_wmma_acc_row instead, and drop the trap comment's restatement of hip_arch_compat.h. Co-Authored-By: Copilot <copilot@github.com>

A DLL with no same-name entry in main's baseline was compared against 0, so renaming or adding a kernel DLL failed the gate on its full size. List it as "new" instead; the next PR after it lands on main compares as usual. Co-Authored-By: Copilot <copilot@github.com>

zz002 mentioned this pull request Oct 27, 2026

ci(windows): build gfx1171 in the SDK selfbuild zz002/hip-ep#7

Open

2 tasks

zz002 changed the title ~~ci(windows): build custom kernels for gfx11-7-generic instead of gfx1170~~ ci(windows): target gfx11-7-generic for gfx117x and fix the WMMA conv there Oct 28, 2026

zz002 marked this pull request as ready for review October 28, 2026 07:21

zz002 requested review from CaymanYang, amd-bqiao, amd-mingw, asparkhi, edelaye, fhanuman, qianglin-amd, sambanerxilinxamd and wcy123 as code owners October 28, 2026 07:21

zz002 mentioned this pull request Oct 28, 2026

feat(kernels): replace hipBLASLt with Composable Kernel for all GEMMs #960

Open

9 of 12 tasks

zz002 and others added 4 commits October 28, 2026 05:51

ci(windows): build custom kernels for gfx11-7-generic instead of gfx1170

51cea15

Co-Authored-By: Copilot <copilot@github.com>

zz002 force-pushed the ci/gfx11-7-generic-probe branch from e7deb5c to 83bad03 Compare October 28, 2026 09:51

amd-mingw approved these changes Oct 30, 2026

View reviewed changes

zz002 merged commit 789b424 into main Oct 30, 2026
11 checks passed

zz002 deleted the ci/gfx11-7-generic-probe branch October 30, 2026 12:40

Navigation Menu

Repository navigation

ci(windows): target gfx11-7-generic for gfx117x and fix the WMMA conv there - #1329

ci(windows): target gfx11-7-generic for gfx117x and fix the WMMA conv there#1329
zz002 merged 5 commits into
mainfrom
ci/gfx11-7-generic-probe

zz002 commented Oct 26, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Oct 26, 2026

Uh oh!

github-actions Bot commented Oct 26, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Oct 26, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Oct 26, 2026 •

edited

Loading

Uh oh!

zz002 commented Oct 28, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

zz002 commented Oct 26, 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

Notes for reviewers

Checklist

Uh oh!

github-actions Bot commented Oct 26, 2026

Uh oh!

github-actions Bot commented Oct 26, 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 26, 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!

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

Uh oh!

Kernel Unit Test Results

Uh oh!

zz002 commented Oct 28, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

zz002 commented Oct 26, 2026 •

edited

Loading

github-actions Bot commented Oct 26, 2026 •

edited

Loading

github-actions Bot commented Oct 26, 2026 •

edited

Loading

github-actions Bot commented Oct 26, 2026 •

edited

Loading