Skip to content

fix(miopen): drop invalid CK perf configs from the gfx1200 system database - #13420

Merged
yassinsolim merged 1 commit into
developfrom
users/ysoliman/fix-gfx1200-ck-wrw-bytespan
Oct 13, 2026
Merged

yassinsolim merged 1 commit into
developfrom
users/ysoliman/fix-gfx1200-ck-wrw-bytespan

Conversation

@yassinsolim

@yassinsolim yassinsolim commented Oct 9, 2026 •

Copy link
Copy Markdown
Member

JIRA ID: ROCM-29026

Motivation

Smoke/CPU_DBSync_NONE.StaticFDBSync/10 fails on gfx1200 (#13394). The gfx1200 system
database was added by #10050 and has not been touched since. 20 of its find-db entries name a
CK instance that CK now refuses to run, so TestPerfCfgParams rejects the stored perf config,
FindSolutionImpl logs Invalid config loaded from Perf Db, and the shape falls back to an
untuned solver.

#12088 did this cleanup for gfx1201. gfx1200 never got one.

Unlike gfx1201, gfx1200 has a single cause. All 20 are ConvHipImplicitGemmGroupWrwXdlops
FP16 grouped backward-weights entries whose largest tensor holds between 2^30 and 2^31
elements, which puts them in the gap between MIOpen's large-tensor gate and CK's own overflow
guard:

  • MIOpen's RequiresLargeTensorCKInstance (src/ck_impl/implicitgemm_ck_util.hpp, as restored
    by revert(miopen): partial revert of CK byte span gate for large instances #12396) counts elements and fires at 2^31. None of the 20 reaches it, so MIOpen keeps
    offering the regular int32-indexing instances and the tuner recorded one of them.
  • CK's tensor_exceeds_2gb (tensor_size_check.hpp, added by fix(ck): convolution edge cases large tensors #11754 on 2026-09-30) fires at
    2^31 elements or 2^31 bytes. For a 2-byte type the byte arm trips at 2^30 elements, so
    all 20 do reach it. It sets stride_overflow on the argument and the non-large-tensor
    instance bails out of IsSupportedArgument.

The database predates #11754, which added this check. Before #11754, CK accepted these configs without testing for the overflow, so the tuner recorded them unflagged. That does not mean the earlier runs computed correct results for these shapes; #11754 exists precisely because unguarded int32 indexing can silently compute wrong results, which is the same risk #11585 and #12396 are about for the other directions.

The source-side gap is tracked separately in #13419; this PR only removes the rows that
can no longer resolve.

Technical Details

20 find-db entries and 20 perf-db rows removed. Both numbers are produced by
MIOPEN_DBSYNC_CLEAN=1, which is MIOpen's own predicate (solv.IsApplicable then
solv.TestPerfCfgParams), not a hand-built list.

group find-db entries perf-db rows solver
CK byte-span overflow 20 20 ConvHipImplicitGemmGroupWrwXdlops

Nothing outside that solver is touched. By CK instance family:

CK instance entries split_k values
DeviceGroupedConvBwdWeightTwoStage_Wmma_CShuffleV3 12 64, 128
DeviceGroupedConvBwdWeight_Wmma_CShuffleV3 8 16, 32, 64, 128

File deltas: the find database keeps all 90688 keys and changes 20 lines, each losing only the
unusable solver. The perf database goes from 77756 to 77736 lines, because for all 20 the CK
entry was the key's only perf-db row.

Why each entry fails, measured rather than inferred

An instrumented build was used to record, for every IsCKArgsSupported call the test makes,
whether MIOpen's gate fired, whether the stored instance was found, and which CK line rejected
it. Across all 40 failing checks (20 entries, each checked twice: once by TestPerfCfgParams
and once by FindSolution) the answer is identical:

CKSIZE tensor_exceeds_2gb elems=1233125376 bytes=2466250752 dtsz=2
Stride overflow! In .../device_grouped_conv_bwd_weight_two_stage_wmma_cshuffle_v3.hpp:1177, in function: IsSupportedArgument
DBG13394 key=96-224-224-7x7-3-224-224-256-3x3-1x1-1x1-0-NCHW-FP16-W|rlt=0|found=1|large_ok=1|supported=0|split_k=64|
         in_es=1233125376|wei_es=14112|out_es=38535168|in_b=2466250752|wei_b=28224|out_b=77070336|
         kid=DeviceGroupedConvBwdWeightTwoStage_Wmma_CShuffleV3<...>+64
  • rlt=0: RequiresLargeTensorCKInstance returned false on all 40.
  • found=1: the recorded instance still exists in the registry on all 40. No entry fails
    because an instance was renamed or dropped.
  • large_ok=1: the large-tensor filter rejected nothing.
  • supported=0: CK's IsSupportedArgument refused, in all 40 at the arg.stride_overflow
    branch, in device_grouped_conv_bwd_weight_two_stage_wmma_cshuffle_v3.hpp:1180 (24 checks)
    or device_grouped_conv_bwd_weight_wmma_cshuffle_v3.hpp:1054 (16 checks).
  • All 40 have max element span below 2^31 and max byte span above 2^31.

So nothing here is a split-K problem, an LDS-size or KPerBlock % KPerXdl rejection from
#13079, or a PrefetchStages problem (#13373). The whole set is the byte/element threshold
gap. Element spans run 1.10e9 to 2.11e9 and byte spans 2.21e9 to 4.12e9; every entry is FP16.

On the issue's count of 31

#13394 reports 31 failures. This PR removes 20. The other 11 are not removed here.

The 20 fail in every build tested: the from-source build above and TheRock's CI artifacts. The 11 (DeviceGroupedConvBwdWeight_Explicit_Xdl x4, DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle x4, DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted x3) fail only in TheRock's build. On the same gfx1200 GPU with byte-identical database files, TheRock artifacts from run 37986471185 (rocm-libraries 3ecaa6512fb, no per-target CK changes) and from run 37986476054 (per-target cutover) both fail the same 31 with a byte-identical gfx1200 CK plugin, while the source build fails 20. So the 11 come from the TheRock-built binaries. The cause is not isolated yet; #13394 has the details.

An earlier version of this description attributed the 11 to the per-target CK branch. That was wrong.

Consequence: after this PR, TheRock's gfx1200 StaticFDBSync/10 still fails on those 11 until they are resolved.

Performance consequence

16 of the 20 removed entries were the fastest recorded winner for their key. No key loses its
last entry, so there is no fallback to a runtime find.

For those 16 the next best entry is a median 1.34x slower and at worst 1.84x. The worst case is
64-222-222-3x3-3-224-224-512-0x0-1x1-1x1-0-NCHW-FP16-W, which goes from
ConvHipImplicitGemmGroupWrwXdlops at 50.6900 ms to GemmWrwUniversal at 93.5055 ms. The new
fastest solver is GemmWrwUniversal in all 16 cases.

None of this is recoverable by re-tuning as the code stands: CK refuses the regular instances
for these shapes, and MIOpen's gate does not route them to a large-tensor instance either, so a
fresh tuning run on gfx1200 would record the same unusable configs again. Closing that is what
#13419 is for.

Test Plan

Build CK and MIOpen from one checkout for gfx1200 and run the DBSync test on gfx1200 hardware.

test_db_sync --gtest_filter='Smoke/CPU_DBSync_NONE.StaticFDBSync/10'

Test Result

node       Slurm job 68055153, AMD Radeon RX 9060 XT, gfx1200
           rocminfo Compute Unit: 32, hipDeviceProp multiProcessorCount 16, warpSize 32
           test prints "Handle CU count: 16", selects the gfx120020 database
container  rocm/dev-ubuntu-22.04@sha256:80de159c03b84ef2551d6230a06a59de5cdad1cde288c136bc2042a93aeeda75
source     ROCm/rocm-libraries e56d6561886b79500d814a3cc039ee477bf45f01, MIOpen and CK built
           from the same checkout, CK installed over /opt/rocm so MIOpen binds the CK it was
           built against
cmake      MIOPEN_TEST_DBSYNC=ON MIOPEN_TEST_DISCRETE=ON MIOPEN_BUILD_CK=OFF
           MIOPEN_USE_COMPOSABLEKERNEL=ON MIOPEN_USE_HIPCONV=ON BUILD_DEV=OFF

Before, against the exact blobs of the commit under test:

FDB: "/work/db/active/gfx120020.HIP.fdb.txt", PDB: "/work/db/active/gfx120020.db.txt"
[  FAILED  ] Smoke/CPU_DBSync_NONE.StaticFDBSync/10, where GetParam() = ("gfx1200", 0) (2773 ms)
20 "Invalid perf config found", 0 "Solver is not applicable", 0 "PDB entry does not exist",
0 "Invalid solution"

Oracle pass, MIOPEN_DBSYNC_CLEAN=1:

[       OK ] Smoke/CPU_DBSync_NONE.StaticFDBSync/10 (2774 ms)
20 "Removal complete for fdb-key:"

After, with the cleaned files bz2-compressed and decompressed again, exactly as committed:

roundtrip OK gfx120020.HIP.fdb.txt 4d7c3116a6e2d4beaf8c5315b0d0f5919da7e407b3279d44cc73c4ae68961bf6
roundtrip OK gfx120020.db.txt      08e62d0fd96438d52b350a4bc288a4ff6cf9c8bb5cf8ab826b7ab28904a0b644
[       OK ] Smoke/CPU_DBSync_NONE.StaticFDBSync/10 (2356 ms)
0 residual failures

Content checks on the committed files:

  • pre-change tree content byte identical to develop: decompressed sha256 ab5e514b728496d9
    (find-db) and 2c94d6a19418a233 (perf-db), the same blobs as at 3ecaa6512fb and at the
    current develop head
  • removal list generated by MIOPEN_DBSYNC_CLEAN=1, so it is MIOpen's own predicate
  • removals confined to ConvHipImplicitGemmGroupWrwXdlops, zero outside it
  • the post-change run reports no PDB entry does not exist for tunable failure, so no
    surviving tunable find-db entry was left without its perf-db row
  • bz2 round trip verified against the cleaned files

Submission Checklist

…abase

Smoke/CPU_DBSync_NONE.StaticFDBSync/10 fails on gfx1200. 20 find-db entries
name a ConvHipImplicitGemmGroupWrwXdlops CK instance that CK refuses to run,
so TestPerfCfgParams rejects the stored perf config and the shape falls back
to an untuned solver.

All 20 are FP16 grouped backward-weights whose largest tensor holds between
2^30 and 2^31 elements. MIOpen's RequiresLargeTensorCKInstance counts elements
and so leaves them on the regular int32-indexing instances, while CK's own
tensor_exceeds_2gb guard (added by #11754) also tests the byte span and refuses
those instances at 2^31 bytes. Every one of the 20 sits in that gap.

Removal list produced by MIOPEN_DBSYNC_CLEAN=1 on gfx1200 hardware, not by
hand. 20 find-db entries and 20 perf-db rows removed, nothing outside
ConvHipImplicitGemmGroupWrwXdlops, and the test passes afterwards.
@yassinsolim
yassinsolim requested a review from a team as a code owner October 9, 2026 21:26
Copilot AI balanced review requested due to automatic review settings October 9, 2026 21:26

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

@therock-pr-bot

therock-pr-bot Bot commented Oct 9, 2026 •

Copy link
Copy Markdown

⏳ Waiting on Required Checks

⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass PR does not contain code files — Unit Test auto-passed
🔎 pre-commit ⏳ Pending ⏳ Still running…
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

⏳ Required checks are still running or have not started yet.
This PR is not ready for review until they complete.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Oct 9, 2026 •

Copy link
Copy Markdown

⏳ Waiting for required checks to start or finish:

  • ⏳ pre-commit

@SreecharanGundaboluAMD SreecharanGundaboluAMD left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@yassinsolim
yassinsolim merged commit 9cffe57 into develop Oct 13, 2026
214 of 225 checks passed
@yassinsolim
yassinsolim deleted the users/ysoliman/fix-gfx1200-ck-wrw-bytespan branch October 13, 2026 22:49
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

yassinsolim added a commit that referenced this pull request Oct 17, 2026
…#13373 (#13697)

JIRA ID: ROCM-29026, ALMIOPEN-2090

Closes #13394.

Two fixes to the gfx1200 system database (`gfx120020.HIP.fdb.txt.bz2`,
`gfx120020.db.txt.bz2`), one commit each, plus a changelog commit:

1. Re-tune 8 and remove 3 of the 11 CK backward-weights rows that
TheRock-built binaries reject (the part of #13394 that #13420 left).
2. Drop the 835 CK backward-weights entries that #13373 (merged Oct 14)
invalidates. Before #13373, the configs behind them returned wrong
weight gradients on gfx1200.

They ship together because the database from part 1 alone fails
`StaticFDBSync/10` with 835 invalid configs on any build that has
#13373, which develop now does.

## Part 1: the 11 rows TheRock builds reject

All 11 are `ConvHipImplicitGemmGroupWrwXdlops` entries whose stored
instance has `split_k = -1` ("auto"). At runtime `MakeArgPtr` calls
`GetMaxOccupancy()` (`hipOccupancyMaxActiveBlocksPerMultiprocessor`) on
the live device, multiplies by the CU count, divides by the kernel's
output grid size and picks `k_batch` from that
(`get_best_occupancy_k_batch_value`, `split_k_utils.hpp:37`).
`IsSupportedArgument` then accepts or refuses the derived `k_batch`
depending on the instance's tile shape.

Evidence for occupancy as the mechanism, and its limits:
- #13420's analysis (prior art, not re-measured here) found TheRock CI
artifacts reject 31 entries where a from-source build with the ROCm
7.14.1 toolchain rejects 20. That is a build-to-build pass/fail
difference, not a measured occupancy difference.
- An `LD_PRELOAD` interposer that forces a fixed value for every
occupancy query, run against TheRock's own binaries, shows the failures
follow the forced value: 0 of the 236 auto-split-K WRW rows fail at
occupancy 1, and 1, 3, 3, 5, 6, 7, 9, 17, 28, 39 and 38 fail at 2, 3, 4,
5, 6, 8, 10, 12, 16, 24 and 32 (not strictly monotonic). Unforced,
TheRock's build fails 11.
- A run meant to compare the occupancy each build actually measures
failed to start its container (`Unable to find group render`) and was
not retried. "The two builds measure different occupancy" is the
best-supported explanation, not a direct measurement.

None of the 11 trips the large-tensor gate that #13419 is about.

Each key was re-tuned on the TheRock build (`MIOPEN_FIND_MODE=1
MIOPEN_FIND_ENFORCE=4`, those 11 rows stripped from the system DB first)
and each winner was re-run in a fresh process with `MIOpenDriver -V 1`:

| key | re-tuned instance (+split_k) | verify | action |
|---|---|---|---|
| `1024-14-14-1x1-256-14-14-1024-...-BF16-W` | `Explicit_Xdl` | OK |
kept |
| `128-56-56-1x1-192-56-56-42-...-FP16-W` | `Explicit_Xdl` | OK | kept |
| `256-28-28-1x1-512-28-28-1024-...-BF16-W` | `Explicit_Xdl` | OK | kept
|
| `352-7-7-3x3-352-14-14-10-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1`
| OK (forced) | kept |
| `352-7-7-3x3-352-14-14-11-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1`
| OK (forced) | kept |
| `352-7-7-3x3-352-14-14-22-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1`
| OK | kept |
| `352-7-7-3x3-352-7-7-16-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` |
OK | kept |
| `512-1-1-1x1-512-1-1-1-...-BF16-W` | `Explicit_Xdl` | OK (forced) |
kept |
| `60-13-13-3x3-60-28-28-4-...-BF16-W` |
`Xdl_CShuffleV3_WmmaPorted<64,64,32,32,...>+8` | **FAILED 0.88 > 0.13**
| removed |
| `64-28-28-7x1-64-28-28-1-...-BF16-W` |
`Xdl_CShuffleV3_WmmaPorted<64,32,64,32,...>+-1` | **FAILED 0.84 > 0.13**
| removed |
| `64-28-28-7x1-64-28-28-1-...-FP16-W` |
`Xdl_CShuffleV3_WmmaPorted<64,32,32,32,...>+8` | **FAILED 0.84 > 0.08**
| removed |

"Forced" rows: immediate-mode find never selects that key's CK entry on
its own (a cheaper fallback wins), so the exact CK solver was forced
with `MIOPEN_DEBUG_FIND_ONLY_SOLVER` (Slurm 68090048); all three verify
(1.7e-5, 1.5e-5, 0). The three removed rows are the WMMA-ported V3
family that part 2 is about; #13373 rejects them too.

Default find mode 2, TheRock build, median of 3 runs, RX 9060 XT,
develop's database vs part 1:

| key | develop (ms) | part 1 (ms) | ratio |
|---|---|---|---|
| 1024-...-BF16-W | 6.495 | 5.446 | 0.84x |
| 128-...-FP16-W | 2.095 | 0.918 | 0.44x |
| 256-...-BF16-W | 17.484 | 12.898 | 0.74x |
| 352-...-10-...W_g352 | 0.087 | 0.099 | 1.13x |
| 352-...-11-...W_g352 | 0.090 | 0.075 | 0.84x |
| 352-...-22-...W_g352 | 0.147 | 0.115 | 0.78x |
| 352-...-16-...W_g352 | 0.098 | 0.084 | 0.86x |
| 512-...-BF16-W | 0.036 | 0.036 | 1.0x |
| 60-...-BF16-W (removed) | 0.081 | 0.166 | 2.05x |
| 64-...-BF16-W (removed) | 0.052 | 0.072 | 1.40x |
| 64-...-FP16-W (removed) | 0.089 | 0.038 | 0.43x |

Removing all 11 rows instead of re-tuning was tried and rejected: two
keys then fall back to `GemmWrw1x1_stride1`, which fails verification on
both (0.59 and 0.60 against 0.13, BF16 at N=1024).

## Part 2: the entries #13373 invalidates

#13373 fixes
`DeviceGroupedConvBwdWeight_Xdl_CShuffleV3::IsSupportedArgument`, which
compared `num_k_loop` against the wave64 `PrefetchStages` on RDNA.
Wave32 pipeline-V2 configs with too few K loops were accepted, the
kernel returned without writing its output, and it finished far faster
than the real work, so tuning ranked it first. Those configs are in this
database.

Scope: 1845 find-db records list `ConvHipImplicitGemmGroupWrwXdlops`;
844 resolve to a `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted`
perf config, and in 686 of those the CK entry is the fastest in the
record, so find mode 2 picks it.

Before #13373: all 686 were run on an RX 9060 XT with the TheRock main
CI build of 24bdd07 (run 38489560586), one fresh MIOpenDriver
process each, `-F 4 -t 1 -i 10 -V 1`. Forcing the solver and letting
find mode 2 choose give the same verdict on every key: 680 wrong, 6
correct. Of the 680 failures, 674 have finite relative errors from 0.56
to 2.5e3 (median 0.91), 4 are NaN and 2 are above 1e30.

With #13373: `StaticFDBSync/10` on the part 1 database, built from
24bdd07 plus only the #13373 diff, reports 835 invalid perf configs.
That set contains all 680 measured-wrong keys and none of the 6
measured-correct ones; the other 155 sit in records where CK was not the
fastest, so they were never measured. The committed files are the
`MIOPEN_DBSYNC_CLEAN=1` output of that run: 835 solver entries leave the
find-db, 826 perf-db records change, and the 435 of them that held no
other solver are dropped. No find-db record is left empty. gfx120140
(gfx1201) has no CK backward-weights records and is untouched.

## Validation of this PR's database (parts 1 and 2 together)

`Smoke/CPU_DBSync_NONE.StaticFDBSync/10` on `navi4x-hw-10` (Slurm
68091724):

| check | build | database | result |
|---|---|---|---|
| negative control | 24bdd07 + #13373 | part 1 only | FAILED, 835
invalid perf configs |
| cleanup | 24bdd07 + #13373, `MIOPEN_DBSYNC_CLEAN=1` | part 1 only
| OK, 835 removals |
| this PR | 24bdd07 + #13373 | this PR | OK, 0 failures |
| this PR | 24bdd07, no #13373 | this PR | OK, 0 failures |
| this PR | TheRock CI build 38489560586 | this PR | OK, 0 failures |

The last row also covers part 1: develop's database fails 11 on that
build. The committed bz2 files decompress to exactly the cleanup output
(find-db sha256 `294a5eb5...`, perf-db `b159206a...`).

Correctness after the change, #13373 build, default find mode 2, one
fresh process per problem, `-V 1`, all 686 problems: 686 verify OK
(Slurm 68091782). 680 move off CK (`GemmWrwUniversal` 486,
`GemmWrw1x1_stride1` 63, `ConvWinoRageRxS` 41, Winograd variants 89,
`TransposedConvBinWinoRxS` 1); the other 6 keep CK. All 63
`GemmWrw1x1_stride1` picks verify.

Leaving the stale entries in place is not free once #13373 lands. The
same run against the database without part 2 gave no wrong results, but
90 of the 529 problems it reached before its 6-hour limit hit the
driver's 1200 s timeout. Against this PR's database all 686 problems
finish, the slowest in 493 ms.

## What could break

The 680 problems that move off CK are slower than the times recorded for
the old entries: geomean 4.74x, median 3.69x, worst 4321x (0.11 ms to
493 ms). That is an upper bound, because the old kernel returned without
computing the output, so its recorded time is the cost of a no-op.
Re-tuning recovers most of it: on 12 stratified keys re-tuned with
#13373, a valid CK config wins every time, all 12 verify, and they run a
geomean 6.5x faster than the fallback (1.1x to 64.4x), at 4 to 16 s of
search per key. Re-tuning all 835 entries is a follow-up; at that rate
it is roughly 1 to 2 hours on one gfx1200 (estimate).

## Not in this PR

- The re-tune of the 835 entries.
- #13419 (element count vs byte span in
`RequiresLargeTensorCKInstance`).
- gfx1151 (`/9`) and gfx1102 (`/8`): their reported failures came from
the pre-#12396 large-tensor gate and pass on current builds (GPU-free
check); no database change.

Related: #13305, #12942, #13373.

## Test plan

```
MIOPEN_SYSTEM_DB_PATH=<staged-dir> MIOPEN_USER_DB_PATH=<scratch> \
  ./bin/miopen_gtest --gtest_filter='Smoke/CPU_DBSync_NONE.StaticFDBSync/10'
```

The branch was rebased onto develop after these runs; the database files
are byte-identical to the validated ones, and develop has not changed
them.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants