Repository navigation
fix(miopen): restore gfx1201 backward-data performance lost to the split-K cleanup - #13415
Conversation
…lit-K cleanup JIRA ID: ROCM-29026 #12088 removed 638 invalid CK entries from the gfx1201 system database. 169 of them were the fastest recorded winner for their key. 117 of those are the split-K class, where the surviving entry is usually GemmBwdRest and the recorded cost was a median 6.15x, worst 69.5x. This re-tunes those 117 shapes on gfx1201 hardware and adds the resulting ConvHipImplicitGemmGroupBwdXdlops entries back. Every perf config came from MIOpen's own restricted search, every one carries kbatch 1, and every one was verified numerically against the GPU reference. StaticFDBSync/11 passes on the updated database and MIOPEN_DBSYNC_CLEAN=1 removes nothing from it.
There was a problem hiding this comment.
Copilot wasn't able to review any files in this pull request. Check if the Files changed in this pull request are included in default exclusions.
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
…rent gate no longer rejects (#13421) Restore 228 of the 229 large-tensor CK entries #12088 removed from the gfx1201 system database (165 backward-data, 63 forward find-db entries and 225 perf-db rows). Every restored line is byte-identical to its pre-#12088 original; no lines are removed or newly written. #12396 returned RequiresLargeTensorCKInstance to its backward-weight-only, element-count form once #11754 gave CK its own per-instance overflow checks, so these forward and backward-data entries resolve again. On gfx1201 the pre-#12088 database fails DBSync on 409 split-K entries only; the split-K half is re-tuned separately in #13415. One entry stays out because it was slower. Validated on an AMD Radeon AI PRO R9700 (gfx1201, 64 CU): StaticFDBSync and GPU-reference checks pass for the base, this database and its union with #13415, and a paired MIOPEN_FIND_MODE=2 A/B measured 1.43x aggregate kernel time across the 51 shapes these entries win (single run). JIRA ID: ROCM-29026
#13421 restored large-tensor entries in the same gfx120140 find and perf databases this branch re-tunes. Rebuild both as the key-level union: the two changes touch disjoint keys (228 and 117 find keys, 225 and 108 perf keys), every other line matches develop, and lines stay sorted. The decompressed result matches the union validated natively on gfx1201 (find sha256 cb4d8f30, perf sha256 061f9397).
|
@ROCm/miopen-reviewers could someone review this? It is a database-only change (gfx120140 find-db and perf-db) that restores the gfx1201 backward-data winners #12088 removed. pre-commit, Multi-Arch and Component CI are green. Math CI Summary has been pending since Oct 13 23:00 UTC. The status comes from the rccl status gate ( |
TheRock Submodule Bump ActivityNewest first
|
JIRA ID: ROCM-29026
Motivation
#12088removed 638 invalid CK entries from the gfx1201 system database. 169 of them were thefastest recorded winner for their key, so the removal cost measured performance. That PR said so
and asked for a follow-up re-tune. This is that follow-up for the split-K group.
117 of the 169 are split-K entries. Their surviving winner is
GemmBwdRest95 times, and 19 keyslost their only entry and fall back to a runtime find. The recorded cost was a median 6.15x, worst
69.5x on
1024-1-1-7x7-256-1-1-800-0x0-1x1-1x1-0-NCHW-BF16-B, which went from 0.4843 ms to33.6636 ms.
The other 52 are the large-tensor group and are not touched here. #13421 handles those.
Technical Details
Each of the 117 shapes was re-tuned on gfx1201 with
MIOPEN_FIND_ENFORCE=SEARCHrestricted toConvHipImplicitGemmGroupBwdXdlops. The find-db and perf-db rows added by this PR are the rowsMIOpen itself serialized during that search, not hand-built strings.
ConvHipImplicitGemmGroupBwdXdlops, zero entries outside itEvery tuned config carries kbatch 1. The search did not reselect a split-K configuration for any
of these shapes, so this does not reintroduce the class
#12088removed, and it stays clear ofthe split-K defect in #12942.
117 find-db keys map to only 108 perf-db keys, because the perf-db key drops the output spatial
dimensions. Nine perf-db keys are shared by two find-db keys, and five of those received different
tuned configs. Those five were resolved by measuring every (find-db key, candidate config) pair and
keeping the config with the lower total time, then re-running the find for the ten affected keys
with the chosen config so the recorded time belongs to the config that ships.
Performance consequence
Paired A/B on one node, one build, back to back, default find mode and a fresh user database, which
is what a user gets. Two full replicates.
The solver MIOpen selects changes from
GemmBwdReston 95 shapes toConvHipImplicitGemmGroupBwdXdlopson 110 of 117.Largest recoveries:
128-7-7-5x5-32-7-7-512-2x2-1x1-1x1-0-NCHW-BF16-B340-1-1-3x3-256-1-1-256-1x1-1x1-1x1-0-NCHW-BF16-B340-1-1-3x3-256-1-1-200-1x1-1x1-1x1-0-NCHW-BF16-B1024-1-1-7x7-256-1-1-800-0x0-1x1-1x1-0-NCHW-BF16-BNot every shape improves. Replicate 1 showed 9 shapes below 0.95x and replicate 2 showed 8, but
only 2 are the same shape in both, so the rest is run to run variation on kernels that take under
0.1 ms. The two that persist are
512-7-7-3x3-512-7-7-15-1x1-1x1-1x1-0-NHWC-NHWC-NHWC-BF16-Band340-3-3-3x3-256-3-3-4-1x1-1x1-1x1-0-NCHW-BF16-B, which together cost 0.0243 ms against 210.1 mssaved across the set.
Test Plan
Build CK and MIOpen from the same checkout for gfx1201 and run the DBSync test on matching hardware.
Test Result
Run on a gfx1201 machine reporting 64 compute units, so MIOpen selects the
gfx120140database under test.Before, on the shipped database:
After, on the database in this PR:
The clean run leaves the files byte-identical, so MIOpen's own predicate accepts every entry this
PR adds.
Harness sensitivity was established with a negative control on the same build. Installing the
pre-
#12088database, whose decompressed sha256 arecf621424...for the find-db and49a90431...for the perf-db exactly as#12088recorded, gives:so the harness does detect bad rows, and the zero above is a real zero.
Content checks on the committed files:
f53cd455026861eb, pdbc3a5faab0d738cb5,first 16 hex of sha256 of the decompressed files)
9bb2fac83c2d678a, pdb046c3e66f88dca0aConvHipImplicitGemmGroupBwdXdlopsentry, and zero lines are removed71955 tunable keys
Submission Checklist
References
StaticFDBSync/11fails on 56-CU gfx1201: no system database is shipped for that SKU #13416 is a separate gfx1201 DBSync failure on the 56 CU SKU, unrelated to perf config validityStaticFDBSyncfails for gfx1200: invalid CK perf configs in the system database #13394 is the equivalent gfx1200 report