Repository navigation
fix(miopen): drop CK perf configs invalidated by the large-tensor gate on gfx1100 - #12117
Merged
Merged
Conversation
✅ All Policy Checks Passed
📖 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. |
…e on gfx1100 MIOpen gates CK off for convolutions whose tensor byte span exceeds INT_MAX, because CK forms byte offsets in 32-bit arithmetic there and returns silently wrong results. #10043 introduced that for grouped wrw on element count, and #11585 extended it to byte span across all directions. #11585 also cleaned the system databases it invalidated, but only the CDNA ones (gfx942130, gfx942e4, gfx950100). The Navi databases were not part of that sweep. These four gfx1100 entries are the leftovers, and Smoke/CPU_DBSync_NONE.StaticFDBSync/7 fails on them. Each names a CK instance that is not a large-tensor instance for a problem whose input byte span is over INT_MAX: 3.06 GiB, 3.06 GiB, 2.01 GiB and 2.01 GiB. The discriminator is byte span, not the instance. The three instance strings involved appear in 285 entries in this database and only these four fail. For one of them, 52 entries use it, exactly one has a byte span over INT_MAX, and that one is the failing entry. The removal list was produced by running the test with MIOPEN_DBSYNC_CLEAN=1 on gfx1100 hardware, so it comes from MIOpen's own predicate. A clean re-run afterwards passes with zero invalid configs.
yassinsolim
force-pushed
the
users/ysoliman/fix-gfx1100-invalid-ck
branch
from
September 15, 2026 17:00
0f746a8 to
e56bd74
Compare
yassinsolim
changed the title
jdcampbe
approved these changes
Sep 16, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
evedovelli
pushed a commit
that referenced
this pull request
Sep 23, 2026
…e on gfx1100 (#12117) JIRA ID: ROCM-29026 ## Motivation MIOpen gates CK off for convolutions whose tensor byte span exceeds INT_MAX. CK forms byte offsets in 32-bit arithmetic there, so those shapes return silently wrong results. `#10043` introduced that guard for grouped wrw on element count, and `#11585` extended it to byte span across all directions. `#11585` also cleaned the system databases it invalidated, but only the CDNA ones (`gfx942130`, `gfx942e4`, `gfx950100`). The Navi databases were not part of that sweep. This PR cleans gfx1100. The rest of the Navi family is still affected, see Remaining scope below. `Smoke/CPU_DBSync_NONE.StaticFDBSync/7` fails on these four entries. ## Technical Details Each removed entry names a CK instance that is not a large-tensor instance, for a problem whose input byte span is over INT_MAX. | shape | dir | input byte span | |---|---|---| | `128-112-112-3x3-64-112-112-1024` | bwd | 3.06 GiB | | `64-224-224-3x3-64-224-224-512` | bwd | 3.06 GiB | | `128-513-513-3x3-128-256-256-32` | fwd | 2.01 GiB | | `336-56-56-1x1-192-56-56-1024` | fwd | 2.01 GiB | The discriminator is byte span, not the instance. The three instance strings involved appear in 285 entries in this database and only these four fail. Taking one of them, `DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_WmmaPorted<128, 128, 64, 32, Filter1x1Stride1Pad0, 32, 32, 2, 2, 8, 8, 8, 1, 1, 1>`: 52 entries use it, exactly one has a byte span over INT_MAX, and that one is the failing entry. The other 51 resolve normally. ### File deltas Find database: 104253 lines before and after, no keys lost. Four lines have the gated solver stripped and keep their remaining winners. Perf database: 96165 to 96163 lines. Two lines held only the CK entry and are deleted. Two lines lose the CK entry and keep `ConvBinWinogradRxSf3x2`. So four find-db entries and four perf-db entries are removed, spread over four and four lines respectively, and two of those perf-db lines survive with their other solver intact. Two of the rewritten find-db lines come back with their remaining winners in descending time order rather than ascending. That is an artifact of the database being rewritten from an unordered container. It does not affect selection: `DbRecord` stores values in a `std::unordered_map`, so file order is not preserved on load, and both the Find path and immediate mode sort by time before use. ## Remaining scope This PR fixes gfx1100 only. Under the same predicate, four other shipped databases still contain entries that will fail the same way on matching hardware: | database | arch | entries | |---|---|---| | `gfx110220` | gfx1102 | 4 | | `gfx115128` | gfx1151 | 1 | | `gfx120020` | gfx1200 | 107 | | `gfx120140` | gfx1201 | 229 | `#12088` covers the gfx1201 rows. The other three need hardware that is currently unavailable. `StaticFDBSync` parameters 8, 9 and 10 should still fail on their respective hardware after this merge. ## Test Plan Build CK and MIOpen from this branch for gfx1100 and run the DBSync test on matching hardware. ``` miopen_gtest --gtest_filter='Smoke/CPU_DBSync_NONE.StaticFDBSync/*' ``` ## Test Result Run on a Radeon RX 7900 XTX against a develop build of CK and MIOpen. Before this change `StaticFDBSync/7` fails with 4 invalid perf configs. After it: ``` [ OK ] Smoke/CPU_DBSync_NONE.StaticFDBSync/7 (1004 ms) [ PASSED ] 1 test. invalid perf configs: 0 ``` The absolute count here depends on the CK revision the build binds against. A build that loads a CK library from a different revision than the MIOpen it is testing can report a large number of unrelated invalid configs. The reproducible part of this result is the delta: this change removes exactly the four entries that MIOpen's own predicate rejects, and nothing else. Content checks on the committed files: - pre-change tree content is byte identical to develop (pdb `b31c42ac05dcd6c7`, fdb `b3b288712cc3e3b7`, first 16 hex of sha256 of the decompressed files) - the removal list was generated by running the test with `MIOPEN_DBSYNC_CLEAN=1`, so it comes from MIOpen's own predicate - removals confined to the affected solvers, zero outside them - zero find-db keys lost - bz2 round trip verified against the cleaned files ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-libraries/blob/develop/CONTRIBUTING.md#pull-requests
yassinsolim
added a commit
that referenced
this pull request
Sep 28, 2026
…e on gfx1102
Four gfx110220 rows name non-large-tensor CK grouped-conv instances on shapes
whose tensor byte span exceeds INT_MAX. MIOpen rejects these configs at
runtime, so StaticFDBSync("gfx1102", 0) fails on gfx1102 runners. This is the
same class as the gfx1100 cleanup in #12117.
Removes exactly those four (find-db key, solver) pairs from the find-db and
their perf-db entries. The rest of both files is byte-identical.
This was referenced Sep 30, 2026
yassinsolim
added a commit
that referenced
this pull request
Oct 2, 2026
JIRA ID: ALMIOPEN-1195 ## Motivation MIOpen's packaged Navi system databases cover gfx1100, gfx1102, gfx1151, gfx1200, and gfx1201. This adds the validated gfx1101 database for AMD Radeon Pro V710 so immediate-mode lookup can use tuned solutions instead of falling back to runtime search or defaults. ## Technical Details - Add the bzip2-compressed `gfx110136.HIP.fdb.txt` and `gfx110136.db.txt` system database pair. - Add gfx1101 to `MIOPEN_INSTALL_GPU_DATABASES` defaults. - Add gfx1101 to the architecture-specific database basename path. The runtime-selected basename is `gfx110136` (Radeon Pro V710, 54 CU). - Add gfx1101 to `StaticFDBSync` coverage. - Generate from only the recorded Radeon Pro V710 data (`gpu_id=29`, `gfx1101`) for MIOpen commit `16a17c9056`. - Remove the two WinoRage solver families using strict find/perf pairing. Neither solver was present, so this pass changed zero entries. - Run the standard no-orphan cleanup before native validation. - Native `StaticFDBSync` rejected 500 CK (find key, solver) pairs as invalid perf configs. Their serialized configs are still rejected on the current build. Those pairs were first removed. - Re-tune the 500 rejected pairs on the exact PR build with exhaustive search (`MIOPEN_FIND_ENFORCE=4`). Search produced a new valid config for all 500. - Restore 393 of them, each with its newly searched perf config and find time. A pair is restored only when the re-tuned CK solver is more than 5% faster in steady state than the solver the database selects without it, and immediate-mode lookup selects it. The other 107 stay out: 99 where CK is slower, 5 within 5%, and 3 where the find ranking would not select CK. No existing entry is removed or modified by the restore. - Also drop 4 gfx1102 (`gfx110220`) CK perf configs invalidated by the large-tensor gate. They made `StaticFDBSync` fail on gfx1102 runners in this PR's CI. This is the same class as the gfx1100 cleanup in #12117. ## Test Plan - Run the SystemDB structural check from the repository. - Run `StaticFDBSync` against an empty isolated system DB as a negative control. - Run `StaticFDBSync` against the generated pair on matching Radeon Pro V710 hardware using a source-pinned MIOpen build. - Rerun `StaticFDBSync` after pruning only reported invalid solver pairs. - Run an immediate-mode `MIOpenDriver` convolution with isolated user/cache paths against the final packaged database pair. - On a MIOpen and CK build of this PR's code, on Radeon Pro V710 (54 CU): - Compare immediate-mode lookup with exhaustive search on a sample of 1,200 tuned problems, stratified by direction, data type and 2D/3D, plus 10 intentional find-DB misses. Layout was not a sampling stratum because the tuning export has no layout field. Parsing the sampled keys afterwards gives 991 NCHW, 163 NCDHW and 46 mixed NCHW-input/NHWC-output problems. - Measure all 500 re-tuned pairs in steady state (3 rounds, 50 iterations, with sequential single-GPU repeats of borderline cases). - Gate the final pair with the structural checker, `StaticFDBSync`, and lookup of the 393 restored keys and the full sample. ## Test Result - Generator completed with 10,953 find entries and 10,236 performance entries. - Cleanup completed with zero parse failures and zero missing perf lines. - Initial native `StaticFDBSync` reported 500 invalid solver pairs. Exact pruning removed 500 find pairs and 500 perf pairs and preserved every unflagged entry. - The final database contains 10,953 find lines and 10,161 performance lines. 393 find lines each gained one restored CK entry, and 381 performance lines were added. - Structural checker: 0 violations, 0 CK rows without a perf entry. `python3 projects/miopen/test/test_system_db_kernels.py` passed. - Native `StaticFDBSync` on Radeon Pro V710 passes with the final pair: 0 invalid perf configs, 0 missing perf entries, 0 inapplicable solvers. - Restored keys (393): immediate mode selects the restored CK solver on all 393, all verified, 1.06x to 39x faster than the previous database (median 8.8x). None slower. - Sample lookup (1,200 + 10 controls): all 1,200 keys hit the find DB and all 10 controls miss, for both the previous and final DB. The 1,159 unchanged keys keep the same selection. 34 mixed-layout problems cannot be checked by the reference implementation. 1 problem fails verification identically with both databases (pre-existing, same solver in exhaustive search). - Immediate mode loaded `gfx110136.HIP.fdb.txt` and `gfx110136.db.txt` from the installed database path. - The 1,200-problem sample covers part of the 10,953 keys. Keys outside the sample are validated structurally and by `StaticFDBSync`, and their selections were not individually benchmarked. **Residual quality audit.** On the exact `bf4db207` build on Radeon Pro V710, built in the ROCm 7.14 development image (`rocm/dev-ubuntu-22.04`), the final database was compared against exhaustive search (`MIOPEN_FIND_ENFORCE=4`) on the same 1,200-shape sample. Of the 1,165 sampled shapes that verify in both runs, 125 are repeat-confirmed more than 5% slower than the exhaustive winner, and 15 of those are more than 1.5x slower. Over the 1,011 keys with repeat-confirmed timing or an identical selection, total kernel time is 1.24% above the exhaustive winners, counting one call per shape with no invocation weights. In 108 of the 125 the faster solver is already in the find DB but ranks lower because its recorded time is stale. In the other 17 the faster solver is absent from the find DB, and 14 of those are naive reference solvers. These numbers come from that development stack and are not a performance result for a shipping ROCm release. This database therefore does not select the fastest solver for every shape. Closing that gap needs the misranked entries re-timed on the shipping ROCm stack, plus a decision from the DB owners on whether naive solvers belong in system DBs. The 34 mixed-layout shapes the reference verifier does not support and 1 pre-existing verification failure are outside this audit. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/rocm-libraries/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: d5b <d5b@amd.com>
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>
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: ROCM-29026
Motivation
MIOpen gates CK off for convolutions whose tensor byte span exceeds INT_MAX. CK forms byte offsets in 32-bit arithmetic there, so those shapes return silently wrong results.
#10043introduced that guard for grouped wrw on element count, and#11585extended it to byte span across all directions.#11585also cleaned the system databases it invalidated, but only the CDNA ones (gfx942130,gfx942e4,gfx950100). The Navi databases were not part of that sweep. This PR cleans gfx1100. The rest of the Navi family is still affected, see Remaining scope below.Smoke/CPU_DBSync_NONE.StaticFDBSync/7fails on these four entries.Technical Details
Each removed entry names a CK instance that is not a large-tensor instance, for a problem whose input byte span is over INT_MAX.
128-112-112-3x3-64-112-112-102464-224-224-3x3-64-224-224-512128-513-513-3x3-128-256-256-32336-56-56-1x1-192-56-56-1024The discriminator is byte span, not the instance. The three instance strings involved appear in 285 entries in this database and only these four fail. Taking one of them,
DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_WmmaPorted<128, 128, 64, 32, Filter1x1Stride1Pad0, 32, 32, 2, 2, 8, 8, 8, 1, 1, 1>: 52 entries use it, exactly one has a byte span over INT_MAX, and that one is the failing entry. The other 51 resolve normally.File deltas
Find database: 104253 lines before and after, no keys lost. Four lines have the gated solver stripped and keep their remaining winners.
Perf database: 96165 to 96163 lines. Two lines held only the CK entry and are deleted. Two lines lose the CK entry and keep
ConvBinWinogradRxSf3x2.So four find-db entries and four perf-db entries are removed, spread over four and four lines respectively, and two of those perf-db lines survive with their other solver intact.
Two of the rewritten find-db lines come back with their remaining winners in descending time order rather than ascending. That is an artifact of the database being rewritten from an unordered container. It does not affect selection:
DbRecordstores values in astd::unordered_map, so file order is not preserved on load, and both the Find path and immediate mode sort by time before use.Remaining scope
This PR fixes gfx1100 only. Under the same predicate, four other shipped databases still contain entries that will fail the same way on matching hardware:
gfx110220gfx115128gfx120020gfx120140#12088covers the gfx1201 rows. The other three need hardware that is currently unavailable.StaticFDBSyncparameters 8, 9 and 10 should still fail on their respective hardware after this merge.Test Plan
Build CK and MIOpen from this branch for gfx1100 and run the DBSync test on matching hardware.
Test Result
Run on a Radeon RX 7900 XTX against a develop build of CK and MIOpen.
Before this change
StaticFDBSync/7fails with 4 invalid perf configs. After it:The absolute count here depends on the CK revision the build binds against. A build that loads a CK library from a different revision than the MIOpen it is testing can report a large number of unrelated invalid configs. The reproducible part of this result is the delta: this change removes exactly the four entries that MIOpen's own predicate rejects, and nothing else.
Content checks on the committed files:
b31c42ac05dcd6c7, fdbb3b288712cc3e3b7, first 16 hex of sha256 of the decompressed files)MIOPEN_DBSYNC_CLEAN=1, so it comes from MIOpen's own predicateSubmission Checklist