Skip to content

feat(miopen): add gfx1250 build to the MIOpen dev image - #12496

Merged
JonathanLichtnerAMD merged 1 commit into
developfrom
users/jlichtne/almiopen-2644-gfx1250-dev-image
Sep 29, 2026
Merged

JonathanLichtnerAMD merged 1 commit into
developfrom
users/jlichtne/almiopen-2644-gfx1250-dev-image

Conversation

@JonathanLichtnerAMD

Copy link
Copy Markdown
Contributor

feat(miopen): add gfx1250 build to the MIOpen dev image

JIRA ID: ALMIOPEN-2644

Description

Adds gfx1250 and gfx1250-strict to the two GPU-arch lists in
projects/miopen/vars/utils.groovy that drive MIOpen's dev Docker images:
getDockerImage()'s ci branch (the arch list published as
rocm/miopen-dev via publishDevDockerImage()) and
buildTheRockDockerImage()'s :therock base-image arch list. These lists
are threaded automatically into the Dockerfile via --build-arg THEROCK_ASIC, which flows into CK's -D GPU_TARGETS and rbuild's
-DGPU_ARCHS, so no other files need changes. gfx1250/gfx1250-strict were
already supported as build targets elsewhere (CK, MIOpen's cdna5 arch libs,
hiptensor, rocWMMA); this closes the gap in the two dev-image arch lists
that had been missed.

Risk Assessment

Risk Level: Low

Two-line, additive-only change to a Jenkins/CI build-config list (no
existing archs removed or reordered). Verified inside a
rocm/miopen-dev:latest container that a fresh MIOpen CMake configure with
-DGPU_TARGETS="gfx1250;gfx1250-strict" succeeds, and that
libMIOpenCKGroupedConv_gfx1250-strict.so builds with a real ~318MB
.hip_fatbin section containing compiled device kernels (confirmed via
clang-offload-bundler/objcopy, not just symbol-name acceptance).

Test Result

Build-only enablement — no gfx1250 hardware exists yet in TheRock's CI
matrix, so there is no on-hardware test to run. Validated via the manual
CMake configure/build described above.

Add gfx1250 and gfx1250-strict to the arch lists that produce the
published rocm/miopen-dev image (getDockerImage's "ci" branch) and the
:therock base image (buildTheRockDockerImage), so CK and MIOpen build
kernels for gfx1250 hardware. Verified in the current dev image
toolchain that both offload-archs compile cleanly and produce distinct
device codeobjects.

JIRA: ALMIOPEN-2644
@therock-pr-bot

therock-pr-bot Bot commented Sep 23, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass PR does not contain code files — Unit Test auto-passed
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@iq136boy iq136boy left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@JonathanLichtnerAMD
JonathanLichtnerAMD merged commit 320f7a8 into develop Sep 29, 2026
176 of 186 checks passed
@JonathanLichtnerAMD
JonathanLichtnerAMD deleted the users/jlichtne/almiopen-2644-gfx1250-dev-image branch September 29, 2026 16:22
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

JonathanLichtnerAMD added a commit that referenced this pull request Oct 2, 2026
…12800)

JIRA ID: ALMIOPEN-2820

## Motivation

Reverts #12496 (320f7a8), which added `gfx1250;gfx1250-strict` to
the MIOpen dev image arch lists. With those archs, the CK stage of the
Docker image build (`build_ck`, Dockerfile:375) fails:

```
ld.lld: error: local memory (131552) exceeds limit (65536)
```

CK's `CK_USE_GFX1250`-gated large-LDS wmma grouped-conv instances are
also compiled for gfx1101/gfx1151 (64 KB LDS). This broke the MIOpen
develop nightly (Jenkins `rocm-libraries-folder/MIOpen/develop` #1589
and #1590) at "Build Docker for CI"; the last good run was #1588.

The break was masked for about a day because the CI image tag hash does
not include the arch list, so #1584-#1588 reused a stale image until an
unrelated CK `CMakeLists.txt` change (#12686) forced a rebuild.

The CK-side fix is tracked in AICK-2468. gfx1250 can be re-added to the
dev image once that lands.

## Technical Details

In `projects/miopen/vars/utils.groovy`, remove `gfx1250;gfx1250-strict`
from two arch lists:
- the `ci` list in `getDockerImage()`
- the `:therock` multiarch list in `buildTheRockDockerImage()`

## Risk Assessment

**Risk Level: 🟢 Low**

### Impacted Components

MIOpen Jenkins CI Docker image build
(`projects/miopen/vars/utils.groovy`) only; no library code changes.

### Potential Side Effects

gfx1250 is temporarily absent from the `rocm/miopen-dev` image.

### Mitigation Steps

Re-add gfx1250 to the dev image once the CK fix (AICK-2468) lands.

## Test Plan

Verify on the next develop nightly: the forced CI image rebuild
completes `build_ck`, then Build Docker for CI, Full Tests, Publish Dev
Image and Nightly Tests pass.

## Test Result

Not yet run; verification is the next develop nightly (see Test Plan).
mattwill-amd pushed a commit that referenced this pull request Oct 14, 2026
JIRA ID: ALMIOPEN-2913

## Motivation

Only 5 of the 30 commits that landed against `projects/miopen` since the
3.6.2 version bump (`2d20a8f775e`, #12211) were reflected in
`CHANGELOG.md`. The remaining 25 were reviewed and the 15
externally-visible ones documented, so the ROCm 10.2.0 section
accurately reflects what shipped.

14 entries added (two of the fixes share one entry), plus a
duplicate-heading fix.

### Added
* [Conv] NHWC (channel-last) support in the GEMM convolution solvers,
with new im2col/col2im kernels (#8693)
* [Conv] Cross-architecture AI heuristics for immediate mode
(`MIOPEN_FIND_MODE=2`), extending data-driven selection to Navi (#9680)
* [Conv] Missing 3x3 stride-1 pad-1 tile-28 CK depthwise backward-data
instance — ALMIOPEN-2423 (#11271)

### Changed
* [Conv] `ConvHipConv` held out of the AI perf-config picker until its
model is retrained — ALMIOPEN-2494 (#13337)
* [Conv] CK byte-span gate lifted for forward and backward-data, now
that CK fixes the 32-bit offset overflow itself — ALMIOPEN-2645 (#12396)

### Optimized
* [Conv] `ConvDirectNaiveConv`: per-direction kernel files, float
accumulators, spatial tiling, bfloat16 path — ALMIOPEN-2485/2486 (#7529)
* Kernel cache reuses cached SQLite connections instead of reopening per
load/store — ALMIOPEN-2598 (#11933)

### Resolved Issues
* [Conv] Incorrect backward-weights results from the gfx950 direct
kernel at `waves_q == 1` (missing cross-lane ordering on LDS staging
reuse) — ROCM-31508 (#12406)
* [Conv] Up to 8x slower hipconv `direct`/`direct_l1`/`direct_wgrad`
under ROCm 10.1's full AMDGCN LTO default; fixed with `-fno-offload-lto`
— ROCM-31334 (#12328)
* [Conv] gfx942 multi-node training throughput regression from backward
GEMM on small-batch 1x1 NHWC shapes — ROCM-31155 (#13232)
* [Conv] Invalid CK perf-config entries in the gfx1201 and gfx1100
system databases — ROCM-29026 (#12088, #12117)
* [BatchNorm] JIT compilation failures on non-GFX9 targets; Welford DPP
assembly restricted to GFX9 — ALMIOPEN-2838 (#12889)
* Out-of-bounds write in `miopenLoadSolution` when a serialized solution
carries more than three work dimensions — ROCM-31082 (#12242)
* `uint64_t` typedef conflict between `miopen_cstdint.hpp` and system
`<stdint.h>` on LP64 Linux (#12623)

## Technical Details

Single-file change to `projects/miopen/CHANGELOG.md`.

The 10.2.0 section had **two** `### Resolved Issues` headings. The two
entries under the trailing one (`ConvHipConv` gfx1250 crash/LDS race,
and the perf-db reload config loss) were merged into the first, and the
trailing heading removed — this accounts for the 2 removed bullets in
the diff. Section order is now `Resolved Issues → Added → Changed →
Optimized → Removed`, matching the 3.6.1 section. `Optimized` is new to
this release section.

Already documented, so untouched: #12874 (reduce int32 rejection),
#12000 / #13016 / #13213 (hipconv v0.3.1 / v0.3.2 / v0.4.0), #11761
(gfx803 removal).

## Notes

Deliberately excluded as not externally visible (10 commits):

* **hipDNN convolution forwarding (#12276)** —
`MIOPEN_ENABLE_HIPDNN_WRAPPER` defaults to `OFF`
(`projects/miopen/CMakeLists.txt:162`), so no shipped build compiles the
forwarding code and the runtime `MIOPEN_HIPDNN_FORWARDING` flag has
nothing to enable.
* **Test-only**: #13289 (gfx1250 depthwise test applicability,
ALMIOPEN-2841), #11610 (hipDNN parity harness), #12660 (Jenkins status
post).
* **Internal kernel-math cleanup, generated ISA unchanged**: #13312
(drop direct `__ocml` calls), #13190 (remove broken single-argument
`__half` overloads).
* **Add/revert pairs, net no change**: #12466 + #13066 (gfx1101 Navi
system DBs), #12496 + #12800 (gfx1250 in the dev image arch lists).

## Risk Assessment

Documentation-only change to a single Markdown file; no code, build, or
runtime impact.

**Risk Level: 🟢 Low**

## Test Plan

None required — Markdown-only. Verified that the entry topics are not
already documented elsewhere in the file, that no content was lost when
merging the duplicate heading, and that every new entry maps to a commit
in `2d20a8f775e..origin/develop -- projects/miopen`.

## Test Result

16 bullets added, 2 relocated — 14 net-new entries covering 15 commits.
Duplicate heading removed; section ordering matches 3.6.1.

Co-authored-by: Claude <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants