Repository navigation
test(miopen): cover the shipped Navi system DBs in the GPU-free dbsync runner - #13411
Open
yassinsolim wants to merge 1 commit into
Open
yassinsolim wants to merge 1 commit into
yassinsolim wants to merge 1 commit into
Conversation
1 task done
✅ 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. |
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
Effective coverage depends on the still-open ROCm/TheRock#8906 scheduling change and architecture-emulation behavior requiring human validation.
0 open findings
What changed in this PR
Extends the GPU-free MIOpen DBSync runner to validate five shipped Navi system databases.
Changes:
- Supports multiple targets per GPU family.
- Adds gfx1100, gfx1102, gfx1151, gfx1200, and gfx1201 configurations.
- Adds retargeting for architectures lacking upstream rocjitsu configurations.
| File | Description |
|---|---|
projects/miopen/test/gtest/dbsync/run_dbsync_rocjitsu.py |
Adds per-target Navi DBSync execution and rocjitsu configuration generation. |
🧠 Review effort: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
This was referenced Oct 9, 2026
…c runner StaticFDBSync has a param for each of the five shipped Navi system DB pairs (gfx110060, gfx110220, gfx115128, gfx120020, gfx120140), and every one of them is skipped on every runner the lane currently lands on. Develop run 37944725308 shows params 7, 8, 9, 10 and 11 all SKIPPED, because all four Test miopen shards of Windows::release / Test gfx110X-all ran on gfx1101 runners. Across the last 40 completed develop runs of Multi-Arch CI, 90 of 90 gfx110X-all MIOpen jobs landed on gfx1101. So nothing gates those databases, which is how #13394 reached develop. Extend the existing rocjitsu runner instead of asking for new hardware. FAMILY_MAP now maps a family to a list of targets, because gfx110X-all ships two system DBs (gfx1100, gfx1102) and gfx120X-all ships two (gfx1200, gfx1201). cu_correct_config becomes write_config and gains an optional arch retarget for gfx1102 and gfx1200, which have no upstream rocjitsu KMD config; each is derived from its same-generation sibling by rewriting gfx_target_version. The CK grouped-conv plugin assertion now runs per target arch. The cus values are physical CU counts, which is what the DB basename encodes in hex, because Handle::GetDbBasename takes the GetMaxHardwareComputeUnits branch for gfx11xx and gfx12xx. A wrong value cannot validate the wrong file: SetupPaths asserts the basename exists. Validated against e56d656. gfx1100 at 96 CU resolves gfx110060, gfx1102 at 32 CU resolves gfx110220, gfx1151 at 40 CU resolves gfx115128; all three pass with zero MIOPEN_DBSYNC_CLEAN=1 removals, and a planted corrupt perf config fails and is removed on exactly the planted entry. The emulated gfx1100 run agrees with the same database on a real RX 7900 XTX, down to the identical post-run file hashes. Refs: #12013, #13394
yassinsolim
force-pushed
the
users/ysoliman/dbsync-rocjitsu-navi-families
branch
from
October 22, 2026 00:17
b19d81a to
c712767
Compare
yassinsolim
added a commit
to ROCm/TheRock
that referenced
this pull request
Oct 22, 2026
The miopen-dbsync component is opted in to gfx942 and gfx950, so the five Navi system find/perf DB pairs that rocm-libraries ships (gfx110060, gfx110220, gfx115128, gfx120020, gfx120140) have no gate at all. StaticFDBSync has a param for each, but they are skipped everywhere the lane currently lands: in rocm-libraries develop run 37944725308 all four Test miopen shards of Windows::release / Test gfx110X-all ran on gfx1101 runners and params 7, 8, 9, 10 and 11 all reported SKIPPED. Across the last 40 completed develop runs, 90 of 90 gfx110X-all MIOpen jobs landed on gfx1101. That is how ROCm/rocm-libraries#13394 reached develop. Add gfx1100, gfx1102, gfx1151, gfx1200 and gfx1201 to the linux include_family list. gfx1151 is declared advisory through the expect_failure plumbing added by #8813, with a TODO on ROCm/rocm-libraries#13429: its dbsync passes under the rocjitsu interposer but has never run on real gfx1151 silicon, where CK is_applicable aborts with invalid device function, and the interposer CU correction is load-bearing there (the config declares 32 CU while gfx115128 is a 40 CU database). gfx1100, gfx1102, gfx1200 and gfx1201 are blocking. This costs no new builds. amdgpu_family_matrix.py already builds gfx110x, gfx1151 and gfx120x on linux for presubmit, postsubmit, submodule_bump and nightly, and the component is linux_cpu_runner, so it takes no GPU test capacity. Those families carry tests_on_trigger ["nightly"], so this lands as a nightly gate. Unit tests cover selection by family group string and by individual gfx target, the linux-only scoping, CPU-runner scheduling when the GPU runner is gated, and the emitted expect_failure boolean for every included target. Requires ROCm/rocm-libraries#13411, which adds the matching FAMILY_MAP entries to run_dbsync_rocjitsu.py.
This branch has not been deployed
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.
The five Navi system find/perf DB pairs shipped in
projects/miopen/src/kernels/have no CI gate today.StaticFDBSynchas a param for each of them, but every param is skipped on every runner the lane currently lands on.Evidence, from develop run 37944725308 (2026-10-09T16:3xZ), all four
Test miopenshards ofWindows::release / Test gfx110X-all:Params 7, 8, 9, 10, 11 are gfx1100, gfx1102, gfx1151, gfx1200, gfx1201. All four shards ran on
azure-windows-11-gfx1101-1gpu-rocm-scale-runner-*, and gfx1101 matches none of those params. Across the last 40 completed develop runs of TheRock Multi-Arch CI there were 90gfx110X-allTest miopenjobs and 100% of them ran on gfx1101 runners.That is the gap behind #13394:
gfx120020shipped with CK perf-db entries MIOpen could no longer resolve, and nothing in presubmit or postsubmit checked it. #13420 removed 20 of them, #13697 fixed the remaining 11 (and removed 835 more rows); #13394 is now closed. This is also the concrete answer to #12013.This change extends the existing GPU-free dbsync runner, which already does exactly the right thing for gfx942 and gfx950, to the Navi families. No new hardware, no new runner pool.
What changed
FAMILY_MAPnow maps a family to a list of targets rather than a single arch, becausegfx110X-allships two system DBs (gfx1100 and gfx1102) andgfx120X-allships two (gfx1200 and gfx1201). Each target gets its own rocjitsu run.New entries and the DB each one validates:
Two details are worth a reviewer's attention.
The CU values are physical CU counts, which is what the DB basename encodes in hex. For gfx11xx and gfx12xx,
Handle::GetDbBasenametakes theGetMaxHardwareComputeUnits()branch, which isGetMaxComputeUnits() * 2because HIP reports WGPs on RDNA. So gfx110060 is 0x60 = 96 physical CU, gfx115128 is 0x28 = 40, gfx120020 is 0x20 = 32, gfx120140 is 0x40 = 64. A wrong CU value cannot silently validate the wrong file:SetupPathsasserts the basename exists.gfx1151.jsonupstream declaressimd_count64, which is a 32 CU part. The shipped DB is gfx115128, tuned for 40 CU. The existing CU-correction handles it, but it means the correction is load-bearing for gfx1151 rather than just an occupancy fix-up.The upstream rocjitsu configs directory has no entry for gfx1102 or gfx1200, so those two targets are marked
retargetand derived from the same-generation sibling by rewritinggfx_target_version(gfx1100_w7900.json for RDNA3, gfx1201_r9700.json for RDNA4). The pairs share family_id, wave size and LDS size; the target version is what HIP needs to select the right code object, and a mismatch aborts loudly with "invalid device function" rather than passing silently. Deriving them here avoids a rocm-systems change plus aROCJITSU_REFbump for two JSON files.Validation
Every number below comes from a re-run on this branch rebased onto develop
9b13db87ae0c9e12967af48560ea3565c4f344cd. The rebase is clean and touches one file. The earlier revision of this description reported runs againste56d6561and covered gfx1100, gfx1102 and gfx1151 only; gfx1200, gfx1201 and the gfx1200 retarget path had never been run.Rebasing matters most for gfx120X, where both DBs moved after
e56d6561:The three gfx11xx pairs are byte-identical to the earlier round, so for those the re-run is about the 216 commits of MIOpen and CK that landed in between. The two gfx12xx pairs are new content, and gfx120140 changed as recently as 2026-10-20 (#13415).
Method. The driver is this PR's own
run_dbsync_rocjitsu.py, invoked unmodified withAMDGPU_FAMILIESset to each family in turn (file sha256de6eaedb148903891d6c2c8d31b33adfcd44953d1a63194dc5385b300510288b, verified in the container against the branch blob). CK and MIOpen were built from the rebased tree withMIOPEN_TEST_DBSYNC=ON, rocjitsu from the pinnedROCJITSU_REFb2099a6b8703c089134ad526253752d26bbbf43a,MIOPEN_DBSYNC_MAX_THREADS=1, Ubuntu 24.04. One substitution:THEROCK_BIN_DIR/miopen_gtestis a symlink totest_db_syncfrom the same build, which carries the sameSmoke/CPU_DBSync_NONE.StaticFDBSyncparams, so the whole gtest suite did not have to be built. Family lookup, config retarget, CU correction, the rocjitsu build and the launch all ran as written.Per target: baseline (no
MIOPEN_DBSYNC_CLEAN), then oracle (MIOPEN_DBSYNC_CLEAN=1), then a re-verify after pruning whatever the oracle removed.[ OK ] .../7 (58845 ms)[ OK ] (58757 ms)[ OK ] .../8 (28902 ms)[ OK ] (28862 ms)[ OK ] .../9 (3336 ms)[ OK ] (3618 ms)[ OK ] .../10 (20557 ms)[ OK ] (21098 ms)[ OK ] .../11 (113796 ms)[ OK ] (112928 ms)Handle CUis half the configured CU because HIP reports WGPs on RDNA; MIOpen doubles it back. Every run resolved the intended basename, which is the check that thecusvalues are right, and every pruned file came back byte-identical to the shipped one (0 removed lines, 0 added lines in all ten files).The wall times are host dependent and should not be compared across rows: gfx110X ran on a 192 core node and gfx1151 and gfx120X on 384 core nodes, and the earlier revision's gfx1100 number (11778 ms) came from a third machine. Within a family the three passes are comparable.
The RDNA4 retarget path now has direct evidence. gfx1200 has no upstream rocjitsu config, and the config derived from the gfx1201 one executes real gfx1200 code objects instead of aborting:
The RDNA3 retarget reproduced the same way in this round:
config gfx1100_w7900.json -> rj_config_gfx1102_32.json (arch=gfx1102, gfx_target_version=110002, cu=32, simd_count=64, num_cu_per_sh=32), resolvinggfx110220withHandle CU count: 16.Negative control, so the zeros above are not vacuous. For each of the five targets, one CK perf-db config in a scratch copy of the pair was rewritten to the literal string
notaperfconfig. The check then fails on exactly that entry and the oracle removes exactly that entry. gfx1201, for example:The gfx1200 control hit
1000-1-1-1x1-1024-1-1-128-0x0-1x1-1x1-0-NCHW-FP16-Band the gfx1151 one hit1000-13-13-1x1-512-13-13-2-0x0-1x1-1x1-0-NCHW-BF16-B, same shape of result. In every case the pristine pair was restored afterwards and the test passed again. No file in the repo was touched; the corruption only ever existed in the installed copy undershare/miopen/dbinside a throwaway container.Real silicon cross-check, gfx1200. The same rebased tree, built the same way, was run natively (no interposer) on an RX 9060 XT (gfx1200, 32 CU). It agrees with the interposer run on every observable: same resolved
gfx120020,Handle CU count: 16,[ OK ] Smoke/CPU_DBSync_NONE.StaticFDBSync/10 (19227 ms), 0 oracle removals, and the installed files byte-identical afterwards (fdb sha256294a5eb5232f43dcbd79b1e6a77a6e4464c48cab610c2ef4744ee114cc1e6836, pdb sha256b159206a86b139b45d5becde551cb12b96f6c916d6936a1161cfa07b789a5f3a).What was not run. gfx1201 was not cross-checked on silicon: every gfx1201 node in the pool was held by 14 day jobs, so a GPU slot was not obtainable. gfx1151 remains interposer-only, because on real gfx1151 silicon CK grouped-conv
is_applicableraisesHIP runtime error: invalid device functionand aborts the test before it reads a single entry (#13429, still open). gfx1151 is also the one target where the CU correction is load-bearing, so the companion PR marks it advisory until #13429 is fixed. The gfx1100 silicon comparison in the previous revision (RX 7900 XTX) was againste56d6561; it was not repeated here, and the gfx110060 pair has not changed since.One environment note for whoever runs this outside TheRock's image: rocjitsu needs a C++20
<format>. On Ubuntu 22.04 the build dies withlib/util/include/util/log.h:8:10: fatal error: 'format' file not found. Ubuntu 24.04 is fine.ensure_build_toolsinstalls cmake, build-essential, libdrm-dev and git but does not check the standard library version.Not in this change
Making a rocm-libraries PR that touches
projects/miopen/src/kernels/*.fdb.txt.bz2or*.db.txt.bz2run this check in presubmit needs more than a rocm-libraries edit. The earlier version of this description said the families only had to be added tolinux_amdgpu_familiesin.github/workflows/therock-multi-arch-ci.ymlbehind a path condition. That is necessary and not sufficient, and TheRock #8813 is where it came up.On a rocm-libraries PR event the set of test components comes from TheRock's
test_tools/determine_rocm_test_dependencies.pyplustest_tools/test_policies.toml. A MIOpen change resolves to themiopengraph node, and[component.miopen]hastest_include = ["hipdnn_integration_tests", "hipdnn_samples"]. The stringmiopen-dbsyncappears nowhere in that path: on TheRock main it exists only inbuild_tools/github_actions/fetch_test_configurations.py, its unit test, and the label list inbuild_tools/github_actions/configure_ci_labels.py. So the component is only ever requested by atest:miopen-dbsynclabel or by an all-components (*) sweep such as the nightly. Adding the families tolinux_amdgpu_familieswould build the artifacts and still schedule no dbsync job. A TheRock-side selection change is needed as well. That is a decision for the CI owners, with a real build cost attached, so it is not bundled here. Tracked in #12013.The Navi families have
tests_on_trigger: ["nightly"]inamdgpu_family_matrix.py, so with the companion PR this lands as a nightly gate first.Also note
test_typesexcludesquick, so the check does not run on the quick tier. The suite that surfaced #13394 ismiopen_gtest_quick_gfx120X_suite. Leaving the tier gate alone is deliberate: dbsync is a 15 minute check and the nightly gate is the right place for it, but a reviewer may want to revisit that.Companion change, ROCm/TheRock #8906
File:
build_tools/github_actions/fetch_test_configurations.py, the"miopen-dbsync"component.include_familymatching accepts either the family group string or an individual gfx target within the family (_family_matches, resolved fromfetch-gfx-targets), so listing individual targets is unambiguous.gfx950 is already on TheRock main: #8813 (merged 2026-10-20) added it to
include_familyand made it advisory through the platform-keyedexpect_failuredict. #8906 no longer touches it. What#8906 adds is the five Navi targets:
"include_family": { - "linux": ["gfx942", "gfx950"], + "linux": [ + "gfx942", + "gfx950", + # Navi system DBs shipped in rocm-libraries projects/miopen/src/kernels/: + # gfx110060, gfx110220, gfx115128, gfx120020, gfx120140. Nothing gated them + # before: the Windows gfx110X lane runs on gfx1101 runners, where every Navi + # StaticFDBSync param is SKIPPED. See rocm-libraries#12013 and #13394. + "gfx1100", + "gfx1102", + "gfx1151", + "gfx1200", + "gfx1201", + ], }, "expect_failure": { - "linux": ["gfx950"], + # TODO(rocm-libraries#13429): gfx1151 is validated under the interposer only. On real + # gfx1151 silicon CK grouped-conv is_applicable aborts with "invalid device function" + # before the test reads an entry, and gfx1151 is also the target where the CU + # correction is load-bearing. Drop this once #13429 is fixed and a silicon run is green. + "linux": ["gfx950", "gfx1151"], },gfx1100, gfx1102, gfx1200 and gfx1201 go in as must-pass. The runs in the table above are at the same base the companion PR will be validated against, so there is no open question left on those four. gfx1151 is advisory for the reason in the TODO, and moving any other target to advisory is a one-line edit on the same list.
The comment block above the component also needs updating: it currently names only gfx942 and gfx950.
Cross-repo ordering
run_dbsync_rocjitsu.pyships in the MIOpen dist that TheRock pulls from the rocm-libraries submodule, so the FAMILY_MAP in this PR is what a labelled #8906 run executes. Until the submodule pin contains this commit, every Navi family hits the runner's "no rocjitsu KMD config, skipping" path and the job reports green without testing anything. The order is: merge this PR, then bump TheRock's rocm-libraries submodule pin (or run #8906 with a submodule override), then run #8906 withci:gfx110X-all,ci:gfx1151,ci:gfx120X-all,test:miopen-dbsyncandtest_filter:standard. Thetest_filterlabel is required:test:labels only narrow what runs, and the component declarestest_typeswithoutquick, so a quick-tier PR schedules no job at all.