Repository navigation
feat(miopen): add validated gfx1151 and gfx1200 Navi system DBs - #10050
Merged
yassinsolim merged 7 commits intoAug 26, 2026
Merged
Conversation
Ship dbsync-clean find/perf databases for gfx1151 and gfx1200 with RDNA hex-basename lookup, StaticFDBSync wiring, and CMake install glob support.
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
clang-format 18.1.4 (repo-pinned via .pre-commit-config.yaml) rewraps the gfx1100/gfx1102/gfx1151/gfx1200/gfx1201 special-case condition in Handle::GetDbBasename to the 100-column limit. Formatting only, no behavior change.
MIOpen resolves packaged system databases against the configure-time MIOPEN_SYSTEM_DB_PATH, which projects/miopen/CMakeLists.txt only assigns inside if(BUILD_DEV). db_path.cpp.in consumes that value unconditionally on Windows and has no <lib>/share/miopen/db fallback, so a BUILD_DEV=Off CI build gets an empty path and resolves database lookups relative to the ctest working directory (build/bin/MIOpen) instead of the artifact tree. DATABASE_INSTALL_DIR is bin on Windows, so the databases are installed alongside the test binary in build/bin. Export MIOPEN_SYSTEM_DB_PATH there for Windows test jobs. Left empty on Linux, which already resolves correctly via the share/miopen/db fallback. This unblocks Smoke/CPU_DBSync_NONE.StaticFDBSync on the Windows gfx1151 runner, which is the first StaticFDBSync case to ever execute on Windows.
Member
Author
|
Follow-up issue for the Windows system-DB path gap referenced under Test Result: #10077. Summary of what changed here since the first CI run:
Unrelated red checks observed on the earlier commit, for the record:
|
…ndows" This reverts commit 44358fd.
1 task done
yassinsolim
added a commit
that referenced
this pull request
Aug 12, 2026
…10081) JIRA ID: ALMIOPEN-1195 ## Motivation `Smoke/CPU_DBSync_NONE.StaticFDBSync` fails on the Windows gfx1151 test runner with `Db file does not exist "gfx115128.HIP.fdb.txt"` — note the absence of any directory component. MIOpen resolves its packaged system find/perf databases against `MIOPEN_SYSTEM_DB_PATH`, a value baked in at configure time. `projects/miopen/CMakeLists.txt` assigns that variable only inside `if(BUILD_DEV)`, while `projects/miopen/src/db_path.cpp.in` consumes it unconditionally on Windows and has no install-relative fallback. TheRock CI builds with `BUILD_DEV=Off`, so the compiled-in path is empty and every database lookup resolves relative to the ctest working directory rather than the artifact tree. Linux is unaffected: its non-`BUILD_DEV` branch falls back to `<lib>/share/miopen/db`. This is not a new regression. `StaticFDBSync` skips any parameter whose architecture does not match the runner's GPU, so exactly one parameter executes per runner. gfx1151 is the only Windows runner in this pipeline, and that parameter was only just added, so this assertion had never previously executed on Windows at all. ## Technical Details Adds `MIOPEN_SYSTEM_DB_PATH` to the env block of the Test step in `.github/workflows/therock-test-component.yml`: ```yaml MIOPEN_SYSTEM_DB_PATH: ${{ runner.os == 'Windows' && format('{0}/build/bin', github.workspace) || '' }} ``` - `DATABASE_INSTALL_DIR` is `bin` on Windows, and the test job unpacks artifacts with `THEROCK_BIN_DIR=./build/bin`, so the databases land beside `miopen_gtest.exe` in `build/bin`. - Gated on `runner.os == 'Windows'`. On every other platform the value is the empty string, which MIOpen treats as unset (`db_path.cpp.in` explicitly tests `p.empty()`), so Linux keeps its existing working fallback and behaviour there is unchanged. - No product code is modified. This is a CI-side unblock only. This is deliberately a workaround, not the real fix. The underlying defect — that a `BUILD_DEV=Off` Windows build has no install-relative database fallback at all, which plausibly affects shipped Windows packages and not just CI — is tracked in #10077 and needs a change in MIOpen itself. Split out of #10050 so that the database change and this CI change can be reviewed independently by their respective owners. ## Test Plan - Confirm the workflow parses and the existing pre-commit hook set passes on the changed file. - Confirm the Windows gfx1151 `StaticFDBSync` case locates the packaged databases and passes. - Confirm Linux test jobs are unaffected, since the variable is empty there. ## Test Result - `Check Yaml` pre-commit hook passes on the changed workflow. - Windows `StaticFDBSync` result to be confirmed by CI on this PR. - Linux behaviour is unchanged by construction: the variable is empty outside Windows and MIOpen ignores an empty value. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
9 of 11 tasks
Develop commit 1a14df1 (#10043, ALMIOPEN-2367) tightened ConvHipImplicitGemmGroupWrwXdlops so it requires a large-tensor CK instance once GetElementSpace() exceeds INT_MAX. That gate is solver-level and rank-agnostic, so it applies to every architecture, but #10043 only cleaned the gfx942130 system DB and noted in its own message that other architectures likely carry the same stale rows. The gfx1151 and gfx1200 databases added by this PR were tuned before that gate existed, so they still record a non-large (int32) CK instance for wrw problems above INT_MAX elements. That is precisely the silently-wrong combination the gate now rejects, and StaticFDBSync fails on it. Removed, find DB and paired perf DB: gfx115128 1 entry 64-445-445-7x7-3-889-889-256-...-NCHW-FP16-W gfx120020 5 entries 64-445-445-7x7-3-889-889-208-...-NCHW-FP16-W 64-597-597-7x7-3-1193-1193-{112,114,120,128}-... Only ConvHipImplicitGemmGroupWrwXdlops is stripped from each affected key. Every one retains GemmWrwUniversal, so no problem loses its find/perf entry and none falls back to a naive solver. Confidence differs between the two architectures, deliberately noted: gfx1151 is confirmed by CI. Smoke/CPU_DBSync_NONE.StaticFDBSync/8 with GetParam() = ("gfx1151", 0) failed on the Windows gfx1151 shard of this PR's own build with "Solver is not applicable" and "Invalid perf config found" on exactly that key. gfx1200 is derived by element-count arithmetic over the find DB, not observed. No Windows gfx1200 shard exists in this repo's CI, so those five entries are not exercised by any check. They match the gfx1151 case exactly in shape and solver, but a reviewer should treat them as reasoned rather than measured. Performance note. On gfx1151 the removed entry was also the slower of the two (2158.87 ms versus 1459.64 ms for GemmWrwUniversal), so removal costs nothing. On gfx1200 the removed entries were faster on paper (250-306 ms versus 383-445 ms), so those five shapes get roughly 1.5x slower. Those timings came from the int32 kernel that returns wrong results above INT_MAX, which is the reason the gate rejects it, so the comparison is not like for like. The upstream MIOPEN_DBSYNC_CLEAN pass could not be used to produce this change: the dbsync tooling's default container pins CK at 2026-04-07 and no longer compiles MIOpen's ck_impl wrappers from a develop-based branch (index_t versus long_index_t array width mismatch introduced by the widen_runtime_index_to_long change in #10547). Entries were removed by the same rule the gate applies, with find and perf edits verified 1:1.
jdcampbe
added a commit
that referenced
this pull request
Aug 25, 2026
…dows (#10838) ## Summary On Windows, `miopen::GetSystemDbPath()` returns an empty path for any `BUILD_DEV=Off` build, so every packaged find-db/perf-db filename resolves against the process's current working directory. An installed MIOpen therefore loads its system databases only when the application happens to be started from the install directory, and silently falls back to a full solver search otherwise. This gives Windows the same install-relative resolution Linux already has, by locating the directory of the loaded `MIOpen.dll`. JIRA ID : ALMIOPEN-1195 ## Risk Assessment **Risk 3.** The change is Windows-only, host-side path resolution — no kernel, dispatch, or support-surface change. Its blast radius is nonetheless every Windows user, because the intended effect is that shipped find/perf databases now actually load, which changes solver selection from a search result to the tuned database entry. Those entries come from the validated databases MIOpen already ships, so this is the designed behaviour rather than a new heuristic. `MIOPEN_SYSTEM_DB_PATH` keeps precedence, `BUILD_DEV` still reads from the build tree, Linux is untouched, and if the module lookup ever fails the previous behaviour is retained. One deliberate behaviour change worth reviewer attention: anything relying on the accidental working-directory lookup (placing databases beside an executable and launching from there) must now set `MIOPEN_SYSTEM_DB_PATH=.`. The existing comment shows working-directory resolution was a fallback, not a design goal, but it is user-visible. ## ASIC Coverage Architecture-independent: the diff touches only host-side path resolution and an environment registry entry, with no kernel, dispatch, default-selection, or support-surface change. Standard PR CI is sufficient; no multi-arch sweep is warranted. Verified end-to-end on Windows on a gfx1151 GPU, against the gfx1151 system databases added by #10050. The lane that exercises this for real is the Windows gfx1151 `db_sync` test (`Smoke/CPU_DBSync_NONE.StaticFDBSync`), which only has gfx1151 databases to find once #10050 lands. ## Testing Summary - Unit: three new cases in the existing `db_paths` GTest covering environment-variable precedence, the default being absolute rather than working-directory relative, and the Windows default following `MIOpen.dll`. - Regression guard: the fix was reverted and the tests re-run, to confirm they actually fail without it. - End-to-end: driver configurations decoded from real entries of the shipped system find-db, run from an unrelated working directory with no MIOpen environment variables set and an isolated user database, confirming the solution MIOpen selects is the one recorded in the system database rather than the result of a fresh search. - Repo gate: `pre-commit` (clang-format, MIOpen GTest naming) over all changed files. - Review follow-up: rebuilt MIOpen and the CK grouped-conv plugin on Windows/gfx1151 with the bounded loop (MSVC 14.44 STL, ROCm wheel clang, `BUILD_TESTING=ON`, `MIOPEN_USE_COMPOSABLEKERNEL=ON`), re-ran the `db_paths` GTest, and re-ran the driver from an unrelated working directory to confirm both copies of the loop still resolve the module directory at runtime. Loop termination itself was additionally exercised standalone against stubs, since a misbehaving `GetModuleFileNameW` cannot be produced in-process. ## Testing Checklist - [x] MIOpen db-path unit tests - `build/bin/test_db_paths.exe` - Status: Passed (15/15; 3 new) - [x] Regression guard - reverted the fix, rebuilt, re-ran - Status: Passed (`SystemDbPath_NoEnvVar_IsAbsolute` fails without the fix with `path '' must be absolute`; `SystemDbPath_NoEnvVar_FollowsMIOpenModule` fails with an empty path instead of the module directory) - [x] End-to-end system-DB lookup on Windows, gfx1151 - `MIOpenDriver conv|convfp16|convbfp16 ... -V 1` over 18 configurations taken from the shipped find-db (fp32/fp16/bf16 x forward/backward-data/backward-weights) - Status: Passed (18/18 located the database and used its recorded solution; 0/18 before the fix, which reported `Database directory does not exist`) - [x] Pre-commit - `pre-commit run --files projects/miopen/src/db_path.cpp.in projects/miopen/src/env_debug.cpp projects/miopen/test/gtest/db_paths.cpp` - Status: Passed - [x] Rebuild with the bounded loop - `ninja MIOpen test_db_paths MIOpenDriver` (Windows, gfx1151, CK on) - Status: Passed (both `db_path.cpp.obj` and `solver/ck_impl_lib_loader.cpp.obj` recompiled, no compiler warnings) - [x] MIOpen db-path unit tests after the review fix - `build-dbtest/bin/test_db_paths.exe` - Status: Passed (15/15, including the three new cases) - [x] Runtime module resolution after the review fix - `MIOPEN_LOG_LEVEL=6 MIOpenDriver convfp16 -n 4 -c 32 -H 28 -W 28 -k 64 -y 3 -x 3 -p 1 -q 1 -g 4 -I NHWC -f NHWC -O NHWC -F 1 -V 1 -i 1`, launched from an unrelated directory - Status: Passed (`[GetLibPath] Lib Path: "...\build-dbtest\bin"`, `Loading Composable Kernel dynamic library: ...\build-dbtest\bin\MIOpenCKGroupedConv_gfx1151.dll`, `Forward Convolution Verifies OK`) - [x] Bounded module-path loop, stub-driven - standalone MSVC harness over the final loop body - Status: Passed (real API resolves in 1 iteration; always-"too small" stub terminates at the 32768-character cap after 8 iterations and returns the previous empty-path fallback instead of looping; `MAX_PATH+10` stub grows once and resolves) - [x] Pre-commit (review follow-up) - `pre-commit run --files projects/miopen/src/db_path.cpp.in projects/miopen/src/solver/ck_impl_lib_loader.cpp` - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending - [ ] Windows gfx1151 `Smoke/CPU_DBSync_NONE.StaticFDBSync` - exercises the shipped databases once #10050 provides gfx1151 databases - Status: Pending ## Technical Changes - `src/db_path.cpp.in`: implement `GetLibPath()` for Windows via `GetModuleHandleExW` / `GetModuleFileNameW` on the module owning `miopenCreate`, mirroring `GetOwningModuleDirectory()` in `src/solver/ck_impl_lib_loader.cpp` rather than introducing a second mechanism. - `src/db_path.cpp.in`: `GetSystemDbPath()` now returns that module directory on Windows for non-`BUILD_DEV` builds. `DATABASE_INSTALL_DIR` is `bin` on Windows (`CMakeLists.txt:555-557`), so the databases are installed next to `MIOpen.dll`. Restructured to an early return so the environment override, `BUILD_DEV`, Windows, and Linux branches read independently. - `src/env_debug.cpp`: register `MIOPEN_SYSTEM_DB_PATH` in the debug environment registry, matching the existing `MIOPEN_USER_DB_PATH` / `MIOPEN_CUSTOM_CACHE_DIR` entries, so tests can set it. - `test/gtest/db_paths.cpp`: three cases added to the existing `CPU_DbPaths_NONE` fixture. The Windows case resolves the module directory independently inside the test, so it validates the contract rather than re-running the implementation, and skips under `MIOPEN_BUILD_DEV`. - `src/db_path.cpp.in`, `src/solver/ck_impl_lib_loader.cpp` (review follow-up): cap the `GetModuleFileNameW` retry buffer at 32768 wide characters — the Windows path ceiling — instead of doubling without bound, warn, and fall back to the existing empty-path behaviour. Applied to both copies of the pattern, per review. A ternary is used rather than `std::min` because both translation units include `windows.h` without `NOMINMAX`, whose `min()` macro breaks the `std::` call. ## Related - #10077 — reports this exact defect; its "Proposed fix" section is what this PR implements. It was closed as completed by #10081, which only exports `MIOPEN_SYSTEM_DB_PATH` in the CI workflow and does not change library behaviour for installed users. Its "Impact beyond CI" section asked for confirmation against a real Windows install; the end-to-end testing above is that confirmation. Suggest reopening it, or tracking the user-facing half here. - #10081 — CI-only mitigation that this change makes unnecessary for the library path. - #10050 — Navi gfx1151/gfx1200 system databases; the change that first made this code path execute on Windows, and the source of the databases used for verification here.
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Sep 2, 2026
## Summary Backfills the **MIOpen 3.6.1 / ROCm 10.1.0** changelog section with externally-visible changes that landed since the `release/therock-10.0` branch but weren't yet documented. All 43 MIOpen commits in that range were reviewed; CI/build-infra, test-only, docs, version-bump, and already-documented changes were excluded. 21 entries added: ### Added * [RNN] Fused single-layer LSTM forward-inference kernel, gated behind `MIOPEN_DEBUG_RNN_FUSED_INFERENCE` (#8915) * [Conv] Naive reference kernel enabling 3D INT8 forward convolution (#10051) * [Conv] `ConvHipConv` solver backed by vendored hipconv kernels (#9746) * [Conv] float8 vector type support + deduplicated math kernels (#10130) * [Conv] Validated gfx1151 and gfx1200 (Navi) SystemDBs (#10050) * [Conv] `MIOPEN_DEBUG_DISABLE_SYSTEM_DB` / `MIOPEN_DEBUG_DISABLE_USER_DB` env vars (#10728) * [Conv] CK depthwise convolution on RDNA wave32 GPUs (#9373) ### Changed * [Conv] Winograd Rage RxS f2x3 on gfx950 + gfx942 v4_6_1/v4_9_1 kernels (#9014) * [Conv] Updated 3D AI heuristics for gfx942 & gfx950 (#8954) ### Optimized * [Conv] Single GEMM for point (1x1) output convolution, all directions, 2D/3D (#10912) * [Conv] Naive backward-data kernel work group widened to 1024 (#10684) ### Resolved Issues * [Conv] Batch-split miscalculation on large WrW tensors — ROCM-26011 (#11188) * [Conv] Conv failure on XNACK-unsupported gfx9 APUs (gfx90c) (#11211) * [BatchNorm] Device heap-overflow in fwd spatial `FinalMeanVariance` — ALMIOPEN-2424 (#11093) * [Conv] Deterministic `SetNextValue` polarity for grouped Bwd/Wrw CK solvers (#10240) * [LRN][LayerNorm] Grid alignment + OOB threads on gfx1250 (#10251) * [Conv] Stale gfx950 grouped-wrw SystemDB entries (#11130) * [Conv][Windows] System-DB path resolution from the MIOpen module (#10838) * [Conv] System-DB fallback so MI308X (gfx942) resolves a usable DB (#8802) * JSON perf logs (`MIOPEN_PERFORMANCE_LOGS`) dropping the last solver during Find (#11166) * Build failure against GCC 15 libstdc++ under C++20 (#9642) ## Notes Deliberately excluded: internal/default-off changes (#10219, #10088, #10649, #10231), a near-duplicate of an existing 10.0.0 entry (#10188), and CI/build-infra/test-only/docs commits. Docs-only change; no code impact. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
This was referenced Oct 9, 2026
Open
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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):
projects/miopen/src/kernels/gfx115128.HIP.fdb.txt.bz2projects/miopen/src/kernels/gfx115128.db.txt.bz2projects/miopen/src/kernels/gfx120020.HIP.fdb.txt.bz2projects/miopen/src/kernels/gfx120020.db.txt.bz2KDB assets are unchanged.
Supporting changes:
projects/miopen/CMakeLists.txt— addgfx1151andgfx1200toALL_GPU_DATABASESso the newassets are globbed and installed.
projects/miopen/src/include/miopen/handle.hpp— register gfx1151/gfx1200 in the existing RDNAbasename special-case list used for runtime database lookup, following the same pattern already
established for gfx1100, gfx1102, and gfx1201 in feat(miopen): refresh Navi system databases (gfx1100, gfx1102, and gfx1201) #9529.
projects/miopen/test/gtest/db_sync.cpp— addStaticFDBSyncparameter rows so the new databasesare structurally validated in CI.
Packaged database basenames are
gfx115128andgfx120020, derived by the existingarchitecture-specific naming logic in
handle.hpp. That special-case list must be extended whenever anew 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:
GemmBwd1x1_stride1(~22% faster than Winograd)GemmWrwUniversal(~35% faster than CK)ConvHipImplicitGemmGroupWrwXdlops(CK ~3x faster than Gemm)Test Plan
Each architecture was taken through the same gate, in order, on matching hardware:
ConvWinoRageRxSisdeliberately retained on gfx120x).
StaticFDBSync(test_db_sync) — required to reach zero failures.MIOPEN_FIND_MODE=2, cleared user DB and cache, verbose logs asserting thepackaged system database is the one loaded.
MIOPEN_FIND_MODE=1plus search, comparing cached winners againstexhaustive winners on a representative sample.
projects/miopen/test/test_system_db_kernels.py(find/perf pairing).Test Result
gfx1151 —
StaticFDBSyncclean (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 —
StaticFDBSyncclean; runtime lookup 9/9; mini-exhaustive 9/9, with cached and exhaustivewinners identical on every sampled case.
Known CI limitation (pre-existing, not introduced here). The Windows gfx1151
StaticFDBSyncshardreports
Db file does not exist "gfx115128.HIP.fdb.txt"with no directory prefix. This is not a problemwith the packaged databases:
MIOPEN_SYSTEM_DB_PATHis only populated whenBUILD_DEV=On, whiledb_path.cpp.inconsumes that configure-time value unconditionally on Windows, so the path is empty ina
BUILD_DEV=OffCI build and lookup resolves relative to the ctest working directory instead of thedirectory 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 pathhad 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