feat(miopen): refresh Navi system databases (gfx1100, gfx1102, and gfx1201) by yassinsolim · Pull Request #9529 · ROCm/rocm-libraries

yassinsolim · 2026-07-16T22:30:41Z

Supersedes #6736.

Motivation

Refresh the packaged Navi system databases and their validation coverage.

Technical Details

Refresh the perf/find databases for gfx1100, gfx1102, and gfx1201.
Update runtime database lookup and installation registration for these assets.
Add structural validation that each find DB has a matching perf DB.
KDB assets are unchanged.

Test Plan

Validate database integrity with StaticFDBSync.
Run immediate-mode lookup smoke tests on target hardware using a clean user DB/cache.
Confirm the packaged system databases are selected at runtime.

Submission Checklist

Reviewed the contributing guidelines.

Refs ALMIOPEN-1195

therock-pr-bot · 2026-07-16T22:35:25Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
🌿 Branch Name	✅ Pass	—
📝 PR Title/Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 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.

therock-pr-bot · 2026-07-16T22:35:27Z

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

0xDELUXA · 2026-07-22T11:30:32Z

Looking forward to a follow-up PR adding the system DB for gfx1200.

mgehre-amd · 2026-07-22T11:51:55Z

@yassinsolim, are you going to do gfx1150 and gfx1151 next?

yassinsolim · 2026-07-22T16:27:32Z

Another PR will be made for the systemDBs on the gfx1200, gfx1150, and gfx1151.

yassinsolim · 2026-07-24T04:25:03Z

CI triage update:

All Jenkins/Math CI, DBsync, clang-tidy, component, and pre-commit checks passed.
The gfx94X shard 3 failure was a softmax GPU hang; its failed-only retry passed unchanged.
The repeated Windows gfx1151 failures are runner/runtime failures: pre-job cleanup timeout, hipInfo GPU-sanity timeout, and GPU launch/COMGR crashes on different tests and hosts.
This PR does not change gfx942 or gfx1151 basename behavior; the special path is gated to gfx1100, gfx1102, and gfx1201.

A third failed-only retry is queued for the three Windows jobs. No unrelated product-code change is justified by these logs.

yassinsolim · 2026-07-24T09:13:34Z

TheRock retry update:

The unchanged gfx94X retry passed. The three Windows gfx1151 jobs failed again on attempt 3, but with different runner/runtime symptoms:

shard 1: unspecified GPU launch failure followed by COMGR abort in a 3D backward-weight group-conv test;
shard 2: COMGR abort in a 2D backward-data group-conv test;
shard 4: one-hour timeout in a 2D forward group-conv test.

Across attempts these failures moved between tests and hosts, while all Jenkins/Math CI, DBsync, clang-tidy, component, and Linux checks passed. This PR does not change gfx1151 basename behavior. Further blind reruns are unlikely to help; this needs Windows gfx1151 runner/runtime triage or maintainer guidance on the CI gate.

Add dbsync-validated system perf/find databases for gfx1100, gfx1102, and gfx1201. Register the architectures for installation and StaticFDBSync, and add a structural find/perf pairing check. Supersedes #6736 for these architectures. Refs #6736, ALMIOPEN-1195

JIRA ID: ALMIOPEN-2364 ## Motivation ROCm 7.15 has been renamed to ROCm 10.0, and MIOpen 3.6.0 is the corresponding release. `projects/miopen/CHANGELOG.md` needs a new version section documenting what changed for downstream users, release-notes tooling, and packagers between MIOpen 3.5.2 (ROCm 7.14.0) and MIOpen 3.6.0 (ROCm 10.0.0), consistent with the changelog's existing per-release format. ## Technical Details - Inserted a new `## MIOpen 3.6.0 for ROCm 10.0.0` section at the top of `projects/miopen/CHANGELOG.md`, directly above the existing `## MIOpen 3.5.2 for ROCm 7.14.0` section, using the same `### Added` / `### Changed` / `### Resolved Issues` / `### Removed` subheading structure as every prior release section in the file. - Entries were compiled by diffing `origin/release/therock-10.0` against `origin/release/therock-7.14` for `projects/miopen` and hand-selecting changelog-worthy commits, excluding: test-only commits, CI-only commits, build/version-bump commits, doc-only commits, reverted commit/revert pairs, and commits already surfaced in the 3.5.2/7.14.0 section via cherry-pick. - Resulting entries added: - **Added:** gfx950 (MI350X/MI355X) 7x7 depthwise forward and backward-data convolution support, fp16/bf16 (#9168) — fixes a slow fallback-to-naive-kernel regression in ConvNeXt-style depthwise convolutions. - **Changed:** - Restored gfx12x support in the Winograd Rage solver (#6439), recovering performance lost when earlier gfx12 support was reverted. - Refreshed the gfx1100, gfx1102, and gfx1201 (Navi) SystemDBs (#9529) with updated tuned find/perf-database entries. - Refreshed the gfx950 SystemDB (#9111) with additional tuned entries. - **Resolved Issues:** - Fixed an off-by-stride indexing bug in BatchNorm backward `CalcStats` mean/variance remainder loop causing a ~1% systematic bias in NCHW batchnorm backward results (#9031). - Fixed a tensor-op integer overflow on gfx1250 for large allocations that could cause memory access faults (#9049). - Fixed naive convolution solvers silently failing to launch (and leaving a stale HIP error visible after Find returned success) when a global work size reached 2^32 or more work-items; such launches are now rejected up front (#10188). - **Removed:** the OpenCL (OCL) backend (#8433); MIOpen now supports the HIP backend only. - No code, build, or test files were touched — this is a pure documentation addition; the underlying functional changes were already merged to `release/therock-10.0` in their own PRs (referenced above) and are only being documented here. ## Risk Assessment Single file (`projects/miopen/CHANGELOG.md`), 17 lines added, 0 removed, no code or build-system changes. **Risk Level: 🟢 Low** ### Impacted Components - `projects/miopen/CHANGELOG.md` only — no source, kernel, solver, build, or test files are touched. ### Potential Side Effects - None functionally. The only observable effect is the new changelog text itself; no build artifacts, APIs, or runtime behavior change. - Cosmetic risk only: if any release-notes/packaging tooling parses this file programmatically, it should be checked for compatibility with the new section, though the added section strictly follows the existing heading/bullet conventions used by all prior sections in the file. ### Mitigation Steps - Followed the exact existing section format (`##` version header, `###` category subheadings, `* ` bullet entries, `[Component]` prefix tags) used by the adjacent 3.5.2/7.14.0 and 3.5.1/7.13.0 sections, minimizing formatting risk. - Each entry was cross-referenced against a specific merged PR number (#9168, #6439, #9529, #9111, #9031, #9049, #10188, #8433) rather than paraphrased from memory, so reviewers can verify each line against its source commit. ## Test Plan No functional tests apply — this is a documentation-only change to a Markdown changelog file. Verification consisted of manually reviewing the diff against the existing file's structure and confirming the new section's heading hierarchy and bullet formatting match every other version section in `CHANGELOG.md`. ## Test Result - `git diff origin/develop...HEAD --stat` confirms the change is scoped to exactly one file (`projects/miopen/CHANGELOG.md`, +17/-0). - Visual/manual review confirms the new `## MIOpen 3.6.0 for ROCm 10.0.0` section renders as valid Markdown consistent with surrounding sections. No build or CI run is applicable to a docs-only change. --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>

JIRA ID: ALMIOPEN-1195 ## Motivation MIOpen ships pre-tuned system find/performance databases so that users get good kernel selection out of the box, without paying for an exhaustive search at runtime. Navi (RDNA) coverage has been incomplete: #9529 landed gfx1100, gfx1102, and gfx1201, leaving the rest of the family to fall back on generic heuristics. This PR extends that coverage to **gfx1151** and **gfx1200**, which are the next two Navi parts with healthy validation hardware available. Both were tuned on matching silicon and taken through the full validation gate before packaging, so immediate-mode lookups now resolve against a curated database instead of defaults. Follow-up to #9529. The two remaining Navi arches (gfx1101, gfx1150) are blocked on hardware faults and are intentionally out of scope here. ## Technical Details **Database assets added** (bzip2-compressed, installed with the other packaged DBs): | File | Rows | |------|------| | `projects/miopen/src/kernels/gfx115128.HIP.fdb.txt.bz2` | 4,998 find | | `projects/miopen/src/kernels/gfx115128.db.txt.bz2` | 4,382 perf | | `projects/miopen/src/kernels/gfx120020.HIP.fdb.txt.bz2` | 90,688 find | | `projects/miopen/src/kernels/gfx120020.db.txt.bz2` | 77,761 perf | KDB assets are unchanged. **Supporting changes:** - `projects/miopen/CMakeLists.txt` — add `gfx1151` and `gfx1200` to `ALL_GPU_DATABASES` so the new assets are globbed and installed. - `projects/miopen/src/include/miopen/handle.hpp` — register gfx1151/gfx1200 in the existing RDNA basename special-case list used for runtime database lookup, following the same pattern already established for gfx1100, gfx1102, and gfx1201 in #9529. - `projects/miopen/test/gtest/db_sync.cpp` — add `StaticFDBSync` parameter rows so the new databases are structurally validated in CI. Packaged database basenames are `gfx115128` and `gfx120020`, derived by the existing architecture-specific naming logic in `handle.hpp`. That special-case list must be extended whenever a new RDNA architecture is packaged, otherwise the runtime will look for a basename that does not match the installed asset. **gfx1151 winner updates.** Three entries were replaced with benchmark-confirmed faster solvers rather than the raw exhaustive-search pick: - FP32 Bwd → `GemmBwd1x1_stride1` (~22% faster than Winograd) - FP16 WrW → `GemmWrwUniversal` (~35% faster than CK) - BF16 WrW → `ConvHipImplicitGemmGroupWrwXdlops` (CK ~3x faster than Gemm) ## Test Plan Each architecture was taken through the same gate, in order, on matching hardware: 1. Tune on matching silicon. 2. Generate the FDB + PDB pair. 3. Clean dead/orphan entries using the architecture-aware policy (forward `ConvWinoRageRxS` is deliberately retained on gfx120x). 4. `StaticFDBSync` (`test_db_sync`) — required to reach zero failures. 5. Prune only exact (find key, solver) pairs surfaced by any failures; re-sync until clean. 6. Runtime lookup with `MIOPEN_FIND_MODE=2`, cleared user DB and cache, verbose logs asserting the packaged system database is the one loaded. 7. Mini-exhaustive overlap with `MIOPEN_FIND_MODE=1` plus search, comparing cached winners against exhaustive winners on a representative sample. 8. Local structural check via `projects/miopen/test/test_system_db_kernels.py` (find/perf pairing). ## Test Result **gfx1151** — `StaticFDBSync` clean (0 failures); runtime lookup 9/9; mini-exhaustive 9/9 numerical. Cached-vs-exhaustive exact agreement 6/9; of the three differences, two are the intentional Gemm winner updates listed above (benchmarked faster than the exhaustive pick) and one is a WinoFury-vs-f3x2 divergence on an FP16 forward case that is performance-neutral in measurement. **gfx1200** — `StaticFDBSync` clean; runtime lookup 9/9; mini-exhaustive 9/9, with cached and exhaustive winners identical on every sampled case. **Known CI limitation (pre-existing, not introduced here).** The Windows gfx1151 `StaticFDBSync` shard reports `Db file does not exist "gfx115128.HIP.fdb.txt"` with no directory prefix. This is not a problem with the packaged databases: `MIOPEN_SYSTEM_DB_PATH` is only populated when `BUILD_DEV=On`, while `db_path.cpp.in` consumes that configure-time value unconditionally on Windows, so the path is empty in a `BUILD_DEV=Off` CI build and lookup resolves relative to the ctest working directory instead of the directory the databases were generated into. Linux is unaffected because it falls back to `<lib>/share/miopen/db`. gfx1151 is the only Windows runner and this parameter row is new, so the path had never executed on Windows before — the gap is newly exposed, not newly created. Being addressed separately on the CI side; a follow-up issue tracks the underlying CMake fix. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

yassinsolim requested a review from a team as a code owner July 16, 2026 22:30

github-actions Bot added the project: miopen label Jul 16, 2026

therock-pr-bot Bot added the Not ready to Review label Jul 16, 2026

assistant-librarian Bot added the external contribution Code contribution from users community.. label Jul 16, 2026

yassinsolim added project: hipdnn organization: ROCm and removed external contribution Code contribution from users community.. project: hipdnn labels Jul 16, 2026

yassinsolim force-pushed the users/ysoliman/systemdb-navi-refresh branch from 7f24aa8 to 27e9525 Compare July 16, 2026 22:55

yassinsolim changed the title ~~feat(miopen): refresh Navi system DBs (gfx1100/1102/1201) and fix gfx1102 CU basename~~ feat(miopen): refresh Navi DBs (gfx1100/1102/1201) + gfx1102 CU fix Jul 16, 2026

yassinsolim force-pushed the users/ysoliman/systemdb-navi-refresh branch from 27e9525 to 6a3dec2 Compare July 16, 2026 23:09

therock-pr-bot Bot removed the Not ready to Review label Jul 16, 2026

yassinsolim force-pushed the users/ysoliman/systemdb-navi-refresh branch 3 times, most recently from a498f13 to 180ef6f Compare July 21, 2026 23:11

yassinsolim changed the title ~~feat(miopen): refresh Navi DBs (gfx1100/1102/1201) + gfx1102 CU fix~~ feat(miopen): refresh Navi system databases Jul 21, 2026

jdcampbe mentioned this pull request Jul 21, 2026

feat(miopen): Update system DBs and add gfx1100/1101/1102/1151/1201 #6736

Closed

jdcampbe approved these changes Jul 21, 2026

View reviewed changes

yassinsolim changed the title ~~feat(miopen): refresh Navi system databases~~ feat(miopen): refresh Navi system databases (gfx1100, gfx1102, and gfx1201) Jul 22, 2026

yassinsolim force-pushed the users/ysoliman/systemdb-navi-refresh branch 3 times, most recently from d1a9b14 to d219e59 Compare July 23, 2026 08:17

yassinsolim force-pushed the users/ysoliman/systemdb-navi-refresh branch from d219e59 to a8d09c8 Compare July 24, 2026 10:55

yassinsolim merged commit b41371a into develop Jul 27, 2026
81 of 83 checks passed

yassinsolim deleted the users/ysoliman/systemdb-navi-refresh branch July 27, 2026 08:24

This was referenced Jul 28, 2026

feat(miopen): add validated gfx1151 and gfx1200 Navi system DBs #10050

Merged

[MIOpen] Windows: GetSystemDbPath() is empty for BUILD_DEV=Off builds, packaged system DBs are not found #10077

Closed

JonathanLichtnerAMD mentioned this pull request Jul 31, 2026

docs(miopen): add ROCm 10.0.0 changelog section #10208

Merged

yassinsolim mentioned this pull request Sep 15, 2026

fix(miopen): drop invalid CK perf configs from the gfx1201 system database #12088

Merged

1 task done

Navigation Menu

Repository navigation

feat(miopen): refresh Navi system databases (gfx1100, gfx1102, and gfx1201) - #9529

feat(miopen): refresh Navi system databases (gfx1100, gfx1102, and gfx1201)#9529
yassinsolim merged 1 commit into
developfrom
users/ysoliman/systemdb-navi-refresh

yassinsolim commented Jul 16, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Jul 16, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Jul 16, 2026 •

edited

Loading

Uh oh!

0xDELUXA commented Jul 22, 2026

Uh oh!

mgehre-amd commented Jul 22, 2026

Uh oh!

yassinsolim commented Jul 22, 2026

Uh oh!

yassinsolim commented Jul 24, 2026

Uh oh!

yassinsolim commented Jul 24, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

yassinsolim commented Jul 16, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Jul 16, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Jul 16, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

0xDELUXA commented Jul 22, 2026

Uh oh!

mgehre-amd commented Jul 22, 2026

Uh oh!

yassinsolim commented Jul 22, 2026

Uh oh!

yassinsolim commented Jul 24, 2026

Uh oh!

yassinsolim commented Jul 24, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

yassinsolim commented Jul 16, 2026 •

edited

Loading

therock-pr-bot Bot commented Jul 16, 2026 •

edited

Loading

therock-pr-bot Bot commented Jul 16, 2026 •

edited

Loading