Repository navigation
fix(miopen): drop the gfx1200 backward-weights entry that returns wrong results - #14037
Merged
yassinsolim merged 5 commits intoOct 26, 2026
Merged
Conversation
…ng results JIRA ID: ALMIOPEN-2090 The gfx1200 system database serves a Composable Kernel backward-weights kernel that computes the wrong weight gradient for one bf16 problem. On gfx1200, MIOpenDriver with the shipped database picks ConvHipImplicitGemmGroupWrwXdlops for 96-18-18-3x3-96-35-35-256-1x1-2x2-1x1-0-NCHW-BF16-W and the result fails verification against the GPU reference with a relative error of about 0.306, on every run, in MIOPEN_FIND_MODE=2 and with the mode unset. The perf config behind it is DeviceGroupedConvBwdWeight_Wmma_CShuffleV3<96, 96, 96, 48, Default, 8, 6, 2, 6, 8, 6, 8, 1, 1, 8, 1, 1>+8 This removes that solver's find-db entry and its perf-db record, so the record's next entry, GemmWrwUniversal, serves the problem instead. Nothing else in either file changes. The underlying Composable Kernel defect is reported separately. All 765 records in this database where ConvHipImplicitGemmGroupWrwXdlops is the fastest entry were checked the same way; this is the only one that fails.
Contributor
There was a problem hiding this comment.
🟢 Approval recommended
The reviewed changelog addition matches the described fix and has no identified blocking issues.
No open issues 🎉
What changed in this PR
Documents the MIOpen gfx1200 backward-weights correctness fix following the earlier system-database cleanup.
Changes:
- Adds a changelog entry describing the affected BF16 problem, incorrect CK results, and fallback to
GemmWrwUniversal.
| File | Description |
|---|---|
| projects/miopen/CHANGELOG.md | Documents the database workaround and replacement solver. |
Review effort: Balanced
✅ 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. |
Brings in the MIOpen 3.6.3 for ROCm 10.3.0 changelog section from #14241. No conflicts. develop has not changed the gfx1200 database files since this branch's merge base, so this PR's versions are kept as they are.
develop now targets MIOpen 3.6.3 for ROCm 10.3.0 (#14241), so the entry moves out of the 3.6.2 section, as requested in review.
…ults Resolve the MIOpen CHANGELOG conflict with #14169 by keeping both entries.
Contributor
TheRock Submodule Bump ActivityNewest first
|
yassinsolim
added a commit
that referenced
this pull request
Nov 2, 2026
… for #13373 (#14038) JIRA ID: ALMIOPEN-2090 ## Stacked on the wrong-results fix This branch sits on #14037, `users/ysoliman/gfx1200-wrw-drop-wrong-results`, and carries that PR's two commits until it merges. Review only the last two commits here: the database commit and the changelog commit. The two changes touch disjoint records, checked record by record on the written files. ## What changed 475 `ConvHipImplicitGemmGroupWrwXdlops` entries go back into the gfx1200 system database, `gfx120020.HIP.fdb.txt.bz2` and `gfx120020.db.txt.bz2`, with perf configs re-tuned on gfx1200 hardware and re-measured against a TheRock CI build. - find-db: 475 records gain one entry each. Record count is unchanged at 90,688. - perf-db: 470 records, of which 193 are extended and 277 are recreated. #13697 had deleted those 277 outright because the CK entry was the only thing they held. Record count goes from 77,298, the count the parent branch leaves, to 77,575. - Every restored perf config carries an explicit `split_k`: 391 at 1, 46 at 2, 24 at 4, 12 at 8 and 2 at 16. None uses auto split-K. - Every other record is untouched. The build script appends text and never re-serializes an existing entry, and it asserts per record that the new content is the old content plus exactly the one appended entry. - Combined with the parent commits, the database this branch produces has 90,688 find-db records and 77,575 perf-db records: 475 records gain a CK entry, one record loses the CK entry that returned wrong weight gradients, and nothing else moves. ## Why #13697 removed the CK backward-weights entry from 838 find-db keys and 829 perf-db records in this database. The configs behind them are all `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted` instances that CK rejects after #13373, which fixed a wave32 support check that read `GridwiseGemm64::BlockwiseGemmPipe::PrefetchStages` where it should read `GridwiseGemm32`. Before that fix those kernels could return without writing the weight gradient, which is the all-zero `dw` in #13305. Removing the entries was correct and had to ship, but it left those problems on whatever solver was next in the record. Measured on a TheRock build, that fallback is 2.64x slower in geomean over the 475 restored keys and up to 50x slower on the worst shape. 835 of the 838 keys went through the re-tune; the candidate list was built before the last three were spotted, so they were never tuned. 475 came back with a config that is both correct and at least 5 percent faster than the fallback. The other 363 keys, the 360 the re-tune did not bring back plus those 3, are left exactly as #13697 left them. ## Why every split_k is pinned The tuner can return `split_k = -1`, which asks CK to derive `k_batch` at run time from an occupancy query. That query is answered by the device and the compiled kernel, so two builds of the same source can disagree, and `TestPerfCfgParams` then rejects an entry that the tuning build accepted. This is not hypothetical. The first candidate database for this change carried 209 auto split-K entries. It passed `StaticFDBSync/10` on the from-source build it was tuned on and failed on a TheRock build with Invalid perf config found fdb-key:1024-6-6-3x3-1024-14-14-2-0x0-2x2-1x1-0-NCHW-FP16-W_g32, Solver: ConvHipImplicitGemmGroupWrwXdlops: DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted<64, 32, 128, 32, Default, 8, 1, 4, 8, 8, 8, 8, 1, 1, 2>+-1 which is the same class of nightly failure #13697 had to fix. So every one of those 209 keys was re-measured at `split_k` 1, 2, 4, 8, 16, 32, 64 and 128 on the TheRock build, and the fastest value that verifies was taken. Pinning costs nothing measurable: over the 208 keys where both forms run, the auto time divided by the pinned time has geomean 1.031 and median 1.001, and 147 of 208 are within 2 percent. 187 of the 209 chose `split_k = 1`. ## How the entries were produced Three builds are used. - TheRock CI run 38988799797 (linux), which pins rocm-libraries 8a29d2e. That commit is 18 commits past the #13373 merge. This is the build the tuning numbers below were measured on, because it is the build shape nightly CI runs. - A from-source build of develop 6bb613a for gfx1200 inside `rocm/dev-ubuntu-22.04@sha256:80de159c`, CK with `MIOPEN_REQ_LIBS_ONLY`, then MIOpen, `MIOpenCKGroupedConv_gfx1200`, `test_db_sync` and `MIOpenDriver`. - TheRock main CI run 39445395419 (linux), which pins rocm-libraries cb26189. This is the build every check in "How it was checked" was re-run on. The gfx120020 database files are the same git blobs at 8585227, at 6bb613a and at develop 03f8851 (`aff3af95` for the find-db, `2c9fa01a` for the perf-db), and both TheRock builds ship that pair unchanged. That pair is the baseline every check below starts from. Per key, every run is a fresh process with the gfx120020 find-db and perf-db hidden unless stated: 1. Fallback. Shipped post-#13697 database, `MIOPEN_FIND_MODE=2`, empty user database. On a find-db hit modes 2, 3 and 5 all take the recorded entry, and every one of these keys is a hit, so this is what a user gets today. 2. Bench. The candidate perf config written into a fresh user perf-db, `MIOPEN_FIND_MODE=1` restricted to `ConvHipImplicitGemmGroupWrwXdlops`. MIOpen benchmarks that one solution and writes the user find-db record, which carries the workspace size and the algorithm. 3. Replay. A fresh process reading both user database files back, so only the recorded winner runs. A Find benchmarks every candidate into the real output buffer, so a Find pass can report a verification the winner alone would not reproduce. Only the replay proves the recorded winner is the one that works. Every driver invocation used `-V 1 -N 1 -t 1 -i 10`. `-N 1` fills the weight-gradient buffer with NaNs first, which is what makes a kernel that writes nothing fail instead of inheriting a previous result. Any run that logged `Invalid config loaded from Perf Db`, or that loaded a config other than the one asked for, was discarded. The find-db and perf-db record strings come out of MIOpen's own `.ufdb.txt` and `.udb.txt`, never hand-built. ## Inclusion rule A key is added back only if all of these hold, otherwise the record is left exactly as #13697 left it: 1. the replay returned `ConvHipImplicitGemmGroupWrwXdlops`; 2. the replay verifies against the reference; 3. the find-db key MIOpen wrote equals the key being targeted; 4. the perf-db key MIOpen wrote equals the one derived from the find-db key; 5. the replay is at least 5 percent faster than the measured fallback. Where the gain landed between 0 and 10 percent, both sides were re-measured as the median of 3 runs. Results on the TheRock build: 475 accepted, 8 rejected, all 8 on the 5 percent bar. Speedup over the fallback for the accepted set: geomean 2.645, median 2.214, p10 1.240, p90 8.687, min 1.054, max 50.23. ## Recorded times A record's times were measured on the machine that originally tuned it, so an absolute millisecond figure measured here would not be on the same scale. The fallback is measured both ways, here and in the record, so its ratio carries this node's measurement onto the record's scale: t_insert = t_fallback_recorded * (t_replay_here / t_fallback_here) Because the fallback is by construction the fastest entry left in each record, and the replay beats it, the inserted time is the new minimum. MIOpen sorts find-db entries by time through `SolutionTimeComparator` before choosing, so CK ranks first. This is asserted per record at build time and confirmed on the written file for all 475. ## Perf-db key collisions 9 pairs among the 835 re-tuned keys share a perf-db key, 18 keys mapping to 826 distinct perf-db keys, because the perf-db key carries the output spatial dims and a 13x13 and a 14x14 input with 3x3, pad 0, stride 2 both produce 6x6. A perf-db record holds one config per solver. In this set the colliding keys agreed on the config, so none had to be dropped. 5 of those pairs are in the accepted set, which is why 475 restored find-db entries land in 470 perf-db records. ## What could break - The 475 configs are validated on one gfx1200 part, which is the part `gfx120020` exists for. They are not validated on any other gfx1200 SKU. - Timings come from one machine. The ordering inside a record is only as good as that measurement, so a record whose top two entries are within a few percent could pick differently elsewhere. The 5 percent bar exists to keep those out. - The entries are new instances, not the ones that were removed. `DeviceGroupedConvBwdWeight_Wmma_CShuffleV3`, `DeviceGroupedConvBwdWeight_Xdl_CShuffle` and `DeviceGroupedConvBwdWeight_Explicit_Xdl` now appear where the removed ones were all `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted`. Any CK change to those families moves these records. - Workspace sizes change for the affected problems, since CK entries carry a workspace and several of the displaced GEMM and Winograd entries carry zero. One related defect, fixed by the parent commits rather than by this one: the exhaustive search produced a wrong result on 4 bf16 keys, all `384-18-18-3x3` shapes, all landing on `DeviceGroupedConvBwdWeight_Wmma_CShuffleV3<96, 96, 96, 48, ...>` with relative error 0.24 to 0.29. Those 4 are excluded by rule 2. The post-#13697 gfx120020 perf-db also used that exact instance in two bf16 entries: `96-18-18-3x3-96-35-35-256` (pad 1, stride 2, split_k 8) and `192-23-1-3x1-192-25-1-1` (stride 1, split_k 1). Both were run. For `96-18-18-3x3-96-35-35-256` the solver that actually ran by default was `ConvHipImplicitGemmGroupWrwXdlops` with that config, in `MIOPEN_FIND_MODE=2` and with the mode unset, three runs each, and it failed verification every time with relative error about 0.306. For `192-23-1-3x1-192-25-1-1` the solver that actually ran is `GemmWrwUniversal`, which verifies; forcing the CK entry on that key verifies as well, so that config is correct for that shape. All 765 records in that database where `ConvHipImplicitGemmGroupWrwXdlops` is the fastest entry were checked the same way on both builds, and `96-18-18-3x3-96-35-35-256` was the only failure. #14037 removes it, and this branch is stacked on that change, so the database here no longer serves it. The CK defect is #14039. ## How it was checked Hardware: a gfx1200 machine with two GPUs, which is the part `gfx120020` serves. Not the single-GPU gfx1200 node used earlier, which has been drained since 2026-10-15 with a prolog error. Every check in this section was re-run on TheRock main CI run 39445395419, against the database this branch produces, which is the parent commits plus this one: the decompressed form of the two `.bz2` blobs this branch carries, find-db sha256 `5b7b43c4675921b9ad7f65fd824c1dbf4c2b1db71e5b7e88fe4079b74af795e7`, perf-db sha256 `431b64f7666b87a842cb768d02ba79a99a017ac24e4a65b5fde74a7b1b36fe9b`. `miopen_gtest` sha256 `f377deb71a777a1da5ecba0bd3404d8a77aa19bb8f32dbce28a8e9170feb9cd6`, `MIOpenDriver` sha256 `56c748f2a962390d33cc95af8a6b3a833e027eddc0cfbaa11a6d582d974e4df9`. - `StaticFDBSync/10` passes on the combined database: zero invalid perf configs, zero not-applicable solvers, zero missing perf-db entries, zero invalid solutions. It passes the same way on the database #13697 shipped and on the parent branch's database, both run in the same session, so the pass is not vacuous. - The other eleven `StaticFDBSync` parameters skip on this hardware, the same on all three databases. - All 476 changed keys, a fresh process each, with `-V 1 -N 1`, both in `MIOPEN_FIND_MODE=2` and with the mode unset. The 475 restored keys each select `ConvHipImplicitGemmGroupWrwXdlops` with exactly the config in the database, verify, and log no `Invalid config loaded from Perf Db` warning. `96-18-18-3x3-96-35-35-256-1x1-2x2-1x1-0-NCHW-BF16-W` selects `GemmWrwUniversal` and verifies. 476 of 476 in each mode. - Positive control in the same session: that key on the unmodified post-#13697 database still selects `ConvHipImplicitGemmGroupWrwXdlops` and still fails verification, relative error 0.3056 to 0.3060, three runs in each mode, so the pick check does catch a wrong result. - A database-loading smoke subset (`CPU_DbPaths_NONE`, `CPU_PerfDb_NONE`, `CPU_SysDbSelectionCu_NONE`, `CPU_Cache_NONE`, `CPU_DBSync_NONE.KDBTargetID`), with the pair copied into the build's own database directory: 39 passed, 1 skipped (`KDBTargetID`, the artifact carries no kernel database), identical on this database, on the parent branch's and on develop's. - The two databases differ from develop's only where they should, checked record by record on the decompressed files: the parent commits change one find-db record and remove one perf-db record, this branch appends one entry to 475 find-db records and to 193 perf-db records and recreates 277 perf-db records, and the two sets of keys do not overlap. Every appended find-db entry is the fastest in its record. Earlier, on TheRock CI run 38988799797 and the from-source build of develop 6bb613a, the same `StaticFDBSync/10` and 476-key pick checks passed in both modes, alongside these controls: - Negative control 1: the earlier candidate database, identical except that one entry keeps its auto split-K, fails `StaticFDBSync/10` on the TheRock build with the `Invalid perf config` line quoted above, while passing on the from-source build. - Negative control 2: a copy of this database with one pinned entry changed to a `split_k` CK rejects for that problem fails `StaticFDBSync/10` on both builds. - All 483 fallbacks verified correctly, so no fallback correctness problem was reproduced there. Closes nothing on its own. Follows #13697, depends on #13373, and restores the performance that #13305 and #13373 forced us to give up.
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-2090
What changed
One Composable Kernel backward-weights entry is removed from the gfx1200 system database,
gfx120020.HIP.fdb.txt.bz2andgfx120020.db.txt.bz2.96-18-18-3x3-96-35-35-256-1x1-2x2-1x1-0-NCHW-BF16-Wloses itsConvHipImplicitGemmGroupWrwXdlopsentry. The record keepsGemmWrwUniversal:7.53519,891648,miopenConvolutionBwdWeightsAlgoGEMM, which now serves it.2x96x18x18x1x3x3x1x96x256x1x1x0x2x2x0x1x1x0x0x1xNCHWxBF16xWis removed. Thatrecord held nothing but the CK config.
This change stands alone and merges first. The gfx1200 backward-weights restore PR,
users/ysoliman/gfx1200-wrw-retune-after-13373, is stacked on top of it.Why
The removed entry made MIOpen run a kernel that computes the wrong weight gradient.
With the shipped database on a gfx1200 part,
MIOpenDriverselects
ConvHipImplicitGemmGroupWrwXdlopswith perf configand fails verification against the GPU reference with a relative error of about 0.306. That is
three runs in
MIOPEN_FIND_MODE=2and three with the mode unset, on two independently builtstacks, with the weight-gradient buffer pre-filled with NaNs (
-N 1) so a kernel that writesnothing cannot inherit an earlier result.
The CK instance is at fault, not the database format or the perf config syntax. Forcing the same
instance through a user perf-db on the same problem reproduces the failure, and forcing a
different instance on the same problem verifies. The CK defect is #14039; this change stops the
wrong kernel from being served while that is fixed.
Scope of the check
Every record in the shipped post-#13697 gfx120020 find-db where
ConvHipImplicitGemmGroupWrwXdlopsis the fastest entry, and therefore the default pick, was runthe same way: 765 records, 475 FP16 and 290 BF16, one fresh process each,
MIOPEN_FIND_MODE=2,empty user database,
-V 1 -N 1. All 765 selected the CK solver. 764 verify. The one in thischange does not.
How it was checked
Hardware: a gfx1200 machine.
Re-verified on TheRock main CI run 39445395419 (linux), which pins rocm-libraries
cb261897980ef5b4f482a7765e904e160d1aa59f. That build ships develop's gfx120020 pair byte forbyte (
aff3af95find-db,2c9fa01aperf-db), so the baseline arm below is the database a usergets today.
miopen_gtestsha256f377deb71a777a1da5ecba0bd3404d8a77aa19bb8f32dbce28a8e9170feb9cd6,MIOpenDriversha25656c748f2a962390d33cc95af8a6b3a833e027eddc0cfbaa11a6d582d974e4df9.StaticFDBSync/10passes on this database: zero invalid perf configs, zero not-applicablesolvers, zero missing perf-db entries, zero invalid solutions. Develop's database passes the
same way in the same session.
StaticFDBSyncparameters skip on this hardware, the same on this database ason develop's.
-V 1 -N 1, three runswith
MIOPEN_FIND_MODE=2and three with the mode unset:GemmWrwUniversalevery time,verifying with a relative error of 0, at 6.51 to 7.45 ms. The same key on develop's database in
the same session still selects
ConvHipImplicitGemmGroupWrwXdlopsand still fails, six runs forsix failures, relative error 0.3056 to 0.3060.
CPU_DbPaths_NONE,CPU_PerfDb_NONE,CPU_SysDbSelectionCu_NONE,CPU_Cache_NONE,CPU_DBSync_NONE.KDBTargetID), run with the pair copied into the build's owndatabase directory so nothing is redirected: 39 passed, 1 skipped (
KDBTargetID, the artifactcarries no kernel database). Identical on develop's database and on this one.
The 765-record sweep above ran earlier, on TheRock CI run 38988799797 (rocm-libraries pin
8a29d2e8) and on a from-source build of develop6bb613ae74for gfx1200, CK withMIOPEN_REQ_LIBS_ONLYthen MIOpen, insiderocm/dev-ubuntu-22.04@sha256:80de159c. Both buildsgave the same single failure.
What could break
7.53519 ms for
GemmWrwUniversal, but the CK number was measured on a kernel that returns thewrong answer, so it was never a real option.
first is not served by CK today, so it was not run, with one exception that was run and
verified.