Skip to content

ci(windows): build gfx1171 in the SDK selfbuild - #7

Open
zz002 wants to merge 2 commits into
ci/sdk-rc7-11dfrom
ci/sdk-rc7-11d-gfx1171
Open

ci(windows): build gfx1171 in the SDK selfbuild#7
zz002 wants to merge 2 commits into
ci/sdk-rc7-11dfrom
ci/sdk-rc7-11d-gfx1171

Conversation

@zz002

@zz002 zz002 commented Oct 27, 2026

Copy link
Copy Markdown
Owner

Summary

Builds gfx1171 in the SDK selfbuild, with its Tensile logic trimmed to the same four HHS/HSS GridBased stems as gfx1170. Stacked on ci/sdk-rc7-11d (#5).

Why

The current SDK ships no bin\hipblaslt\library\gfx1171, so on gfx1171 hipBLASLt finds no Tensile library, every hipBLASLt matmul fails and VLM output is garbage. ROCm#1329 moves the gfx1171 custom kernels to gfx11-7-generic; hipBLASLt needs matching data.

What

  • GPU_ARCHS: add gfx1171.
  • Step 6 keep-set and the trim check treat gfx1171 as a trimmed arch, like gfx1170; bin\hipblaslt\library\gfx1171 is a required path.
  • Tarball: rocm-dist-windows-gfx115x-gfx117x-<ver>-hipblaslt-gfx117x-trimmed.

Test plan

  • Windows SDK Selfbuild completes
  • hipBLASLt logic: gfx1171 reports more than 0 and at most 20 .co; gfx1150-gfx1153 counts unchanged

@zz002
ci(windows): build gfx1171 in the SDK selfbuild
8d6db0e 
Add gfx1171 to GPU_ARCHS and trim its Tensile logic to the same four
HHS/HSS GridBased stems as gfx1170, so the SDK ships
bin\hipblaslt\library\gfx1171. Without it hipBLASLt finds no Tensile
library on gfx1171 and every matmul fails.

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

github-actions Bot commented Oct 27, 2026

Copy link
Copy Markdown

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.

zz002 added a commit to ROCm/hip-ep that referenced this pull request Oct 27, 2026
@zz002
ci(windows): use the SDK that ships gfx1171 hipBLASLt data
ab1acb6 
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>
@zz002 zz002 mentioned this pull request Oct 28, 2026
Merged
2 of 6 tasks
@zz002 @cursoragent
ci(windows): keep the transB=T fp16 Tensile stem for gfx117x
fe4b301 
gemma-4's vision pooler multiplies by a transposed operand, which reaches
hipBLASLt as transA=N, transB=T, fp16 in/out, f32 compute. The gfx1170 and
gfx1171 keep-set has only the Bljk (transB=N) stems, so the heuristic
returns no algorithm, the matmul fails and the image features are garbage.

Add Cijk_Ailk_Bjlk_HHS to the keep-set: one more .co per gfx117x arch
(about 1 MB), well inside the trimmed-set limit. Fail the step when a
keep-set stem is missing, instead of skipping it, so the trimmed library
cannot silently lose a stem the workloads need.

Co-authored-by: Cursor <cursoragent@cursor.com>
zz002 added a commit to ROCm/hip-ep that referenced this pull request Oct 28, 2026
@zz002
ci(windows): use the SDK that ships gfx1171 hipBLASLt data
bd212ec 
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>
zz002 added a commit to ROCm/hip-ep that referenced this pull request Oct 30, 2026
@zz002 @cursoragent
ci(windows): target gfx11-7-generic for gfx117x and fix the WMMA conv…
789b424 
… there (#1329)

## 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 zz002#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

- [x] 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.
- [x] Substantial AI assistance is disclosed, and I reviewed and
understand the result.

---------

Co-authored-by: Copilot <copilot@github.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

1 participant

@zz002