Skip to content

fix(conv): require large-tensor CK instance for grouped wrw on >INT_MAX element count - #10043

Merged
HereThereBeDragons merged 7 commits into
developfrom
users/jlichtne/fix-wrw-large-tensor-elemcount
Aug 18, 2026
Merged

HereThereBeDragons merged 7 commits into
developfrom
users/jlichtne/fix-wrw-large-tensor-elemcount

Conversation

@JonathanLichtnerAMD

@JonathanLichtnerAMD JonathanLichtnerAMD commented Jul 28, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID: ALMIOPEN-2367
JIRA ID: ALMIOPEN-2386

Motivation

The grouped backward-weights CK xdlops solver (ConvHipImplicitGemmGroupWrwXdlops,
and its 3D counterpart ConvHipImplicitGemm3DGroupWrwXdlops) returns silently wrong
results
for NHWC/NDHWC problems whose tensor element extent exceeds INT_MAX while
every individual length and stride still fits int32.

Repro (fp16, NHWC, N=240, C=K=1024, 162x92, 3x3, full tune):

Backward Convolution Weights FAILED: 0.2138 > 0.082

After this fix the same command selects a large-tensor instance and verifies:

Backward Convolution Weights Verifies OK (0.0025 < 0.082)

Technical Details

RequiresLargeTensorCKInstance() gated only on AllTensorsDimsFitIntoInt(), which
validates that each individual length/stride fits int32 — it does not bound
the maximum flat element offset the kernel indexes. A tensor whose addressable
extent exceeds INT_MAX (but whose dims/strides fit) therefore stayed eligible for a
fast int32-indexing CK instance, whose flat element offset overflows once
GenericSearch selects it by timing.

  • RequiresLargeTensorCKInstance() now also returns true when
    GetIn()/GetOut()/GetWeights().GetElementSpace() > INT_MAX. GetElementSpace() (the
    stride-aware max addressable offset, as GetNumBytes() uses) is the correct bound,
    not GetElementSize() (the packed element count): for a non-packed tensor the
    element space is larger and is the true bound on the int32 offset that overflows.
  • Scoped to backward-weights (IsDirectionBackwardWrW()): forward and
    backward-data verify correctly on the int32 instances and keep them — the
    large-tensor int64 instances are measurably slower (e.g. 168 ms -> 323 ms on the
    repro), so gating fwd/bwd would be an unnecessary perf regression.
  • The gate is rank-agnostic, so it covers both 2D
    (ConvHipImplicitGemmGroupWrwXdlops) and 3D
    (ConvHipImplicitGemm3DGroupWrwXdlops) grouped backward-weights.

This is the same class of overflow the ASM-GTC NHWC gate (#10020) addressed; the CK
grouped-solver gate had the same gap.

System DB cleanup

The gate correctly invalidates existing tuning entries that recorded a non-large
(int32)
CK instance for >INT_MAX-element wrw problems — precisely the
silently-wrong combination this fix rejects. On a 304-CU MI300X, StaticFDBSync
flagged 28 such entries in the gfx942130 system DB (find + perf), so this PR
removes them.

  • Only ConvHipImplicitGemmGroupWrwXdlops is stripped from each affected key; other
    solvers (GemmWrwUniversal, ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC) are
    preserved — every affected problem retains a valid, correct find/perf entry, and
    none falls back to a naive solver.
  • Cost of removal is negligible: for 26 of 28 a correct solver was already faster
    than the removed (buggy) entry; for the other 2, the next-best correct solver
    (ConvAsm…WrwXdlopsNHWC) is within ~15% and ~0.4%. Those ConvHip times were
    the wrong-answer int32 kernel anyway; the correct large-tensor instance would be
    slower, so no correct performance is lost.

TODO (follow-up, not in this PR): gfx950 — and other gfx942 CU variants —
very likely carry the same stale large-tensor wrw entries and still need the
equivalent MIOPEN_DBSYNC_CLEAN pass, but that must be run on the corresponding
hardware. gfx950's StaticFDBSync param was skipped in this PR's CI (no gfx950
runner); the non-304-CU gfx942 variants have no StaticFDBSync param at all, so
they are not validated here.

Test Plan / Result

  • CPU unit test WrwElementCountOverflowGate in unit_implicitgemm_ck_util.cpp:
    asserts the gate fires for wrw and not for fwd/bwd on a synthetic
    >INT_MAX-element shape (dims/strides fit int32), and not for a small wrw shape.
    Fast (no GPU/tune). The existing Large2D test does not catch this because its
    tuning is capped at 5 iterations and never selects the offending instance.
  • CPU unit test Wrw3DElementCountOverflowGate in the same file: the NDHWC 3D
    sibling of the above, confirming the rank-agnostic gate fires for 3D grouped wrw
    and stays off for fwd/bwd and a small 3D wrw shape. Also CPU-only/fast.
  • GPU end-to-end numerical test ElementCountOverflow in
    unit_conv_solver_ConvHipImplicitGemmGroupWrwXdlops_ElementCountOverflow.cpp:
    runs ConvHipImplicitGemmGroupWrwXdlops on x=(9,256,1024,1024),
    w=(64,256,3,3), NHWC fp16 and verifies against a GPU naive reference. This
    synthetic shape has element count 2.42B (> INT_MAX) while its largest
    per-dimension stride (C*H*W ~= 268M) stays well under INT_MAX, so it exercises
    the new element-count gate specifically — distinct from the existing per-stride
    gate covered by the _LargeStride sibling (which uses N=1 so the stride itself
    overflows). Verified on MI300X (gfx942): fails pre-fix (RMS ~= 0.20) and
    passes post-fix; the pre-fix CK candidate pool contains int32 instances,
    which the fix removes (_Large_Tensor-only post-fix). Restricted to
    gfx90A/gfx94X/gfx950, gated by a device-memory estimate, and placed in the
    full-only tier (~60s runtime), mirroring the _LargeStride tests.
  • System DB: StaticFDBSync (gfx942, 304-CU MI300X) reproduced the 28 failures
    pre-cleanup and passes after removing the stale entries (verified locally end
    to end: fail -> clean -> pass).
  • Manually validated across confidential large-tensor
    shapes (portrait/landscape, N=240): wrw fixed (e.g. 0.2138 -> 0.00876), fwd/bwd
    unaffected (verify OK without the gate).

@therock-pr-bot

therock-pr-bot Bot commented Jul 28, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 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 Jul 28, 2026 •

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@therock-pr-bot

Copy link
Copy Markdown

Pre-commit check failed

⛔ pre-commit failed

Please run locally:

  • python -m pip install pre-commit
  • pre-commit install
  • pre-commit run --all-files --show-diff-on-failure

This repo uses .pre-commit-config.yaml.

@JonathanLichtnerAMD
JonathanLichtnerAMD force-pushed the users/jlichtne/fix-wrw-large-tensor-elemcount branch 2 times, most recently from f435f00 to 2582806 Compare July 31, 2026 01:59
…AX element count

The grouped backward-weights CK solver (ConvHipImplicitGemmGroupWrwXdlops)
returned silently wrong results (RMS error ~0.21 vs a ~0.08 threshold) for NHWC
problems whose flattened tensor element count exceeds INT_MAX while every
individual length and stride still fits int32 (e.g. N=240, C=K=1024, 162x92).

RequiresLargeTensorCKInstance only gated on AllTensorsDimsFitIntoInt(), which
validates per-dimension/stride int32 fit but not the flattened element count.
Such problems stayed eligible for a fast int32-indexing CK instance whose flat
element offset overflows once the tuner selected it. Add an element-count >
INT_MAX check, scoped to the backward-weights direction: forward and
backward-data verify correctly on the int32 instances and keep them (the
large-tensor instances are measurably slower).

Adds two regression tests:
- A CPU unit test asserting the gate fires for wrw and not for fwd/bwd on a
  synthetic >INT_MAX-element shape (dims/strides fit int32).
- A GPU end-to-end numerical test (ElementCountOverflow) that runs the solver
  on x=(9,256,1024,1024) w=(64,256,3,3) NHWC FP16 and verifies against a GPU
  naive reference. This shape's element count (2.42B) exceeds INT_MAX while its
  largest per-dimension stride (C*H*W ~= 268M) stays well under it, so it
  exercises the new element-count gate specifically rather than the existing
  per-stride gate covered by the LargeStride sibling. Verified to fail pre-fix
  (RMS ~0.20) and pass post-fix. Restricted to gfx90A/gfx94X/gfx950, gated by a
  device-memory estimate, and placed in the full-only tier (~60s runtime).

Also removes 28 stale ConvHipImplicitGemmGroupWrwXdlops entries from the
gfx942130 system DB (find + perf). These recorded a non-large (int32) CK
instance for >INT_MAX-element wrw problems -- exactly the silently-wrong
combination this fix rejects -- so StaticFDBSync (on 304-CU MI300X) flagged
them as invalid once the gate was applied. Only the ConvHipImplicitGemmGroupWrwXdlops
solver was stripped from each affected key; other solvers (GemmWrwUniversal,
ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC) are preserved, so every affected
problem retains a valid, correct find/perf entry (no fallback to naive). For 26
of the 28, a correct solver was already faster than the removed entry; for the
other 2, ConvAsm is within ~15% and ~0.4%.

NOTE: gfx950 (and other gfx942 CU variants) very likely carry the same stale
large-tensor wrw entries and still need the equivalent cleanup, but that must be
run on the corresponding hardware and is NOT done in this PR. gfx950's
StaticFDBSync param was skipped in CI here (no gfx950 runner); gfx942 CU variants
other than 304 have no StaticFDBSync param at all.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
@JonathanLichtnerAMD
JonathanLichtnerAMD force-pushed the users/jlichtne/fix-wrw-large-tensor-elemcount branch from 2582806 to 28e3066 Compare July 31, 2026 02:24
@JonathanLichtnerAMD
JonathanLichtnerAMD marked this pull request as ready for review July 31, 2026 16:54
@JonathanLichtnerAMD
JonathanLichtnerAMD requested a review from a team as a code owner July 31, 2026 16:54
JonathanLichtnerAMD added a commit that referenced this pull request Aug 15, 2026
…ld-essential, arch-matched rocjitsu config)

Validated on gfx950 (run 31865606916): fetch + dist contents + GPU-free StaticFDBSync
all work with the cdna4 config; surfaced real invalid-perf-config gfx950 grouped-wrw
entries (the #2266 / PR #10043-predicted stale class). Key correction: CK IsApplicable
does a HIP device call, so the rocjitsu config must match the arch (cdna4=gfx950).

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.

Do other DBs not need to be updated?
Surprised that only gfx942 needs updates, and not gfx950 + others.

Might be that the DB_SYNC tests only run on gfx942 right now?

@JonathanLichtnerAMD JonathanLichtnerAMD Aug 17, 2026 •

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Yes, it is also broken for gfx950. I was going to fix that as a follow-on change

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Yes, I think we only run this for gfx942, but I will double-check the nightlies

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

This would probably run for gfx950 as part of the nighlies, except that gfx950 is intentionally skipped due to a runner capacity issues (tracked in ROCm/TheRock#3288)

Comment on lines +198 to +200
return problem.GetIn().GetElementSize() > max_int32 ||
problem.GetOut().GetElementSize() > max_int32 ||
problem.GetWeights().GetElementSize() > max_int32;

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.

Should this be GetElementSpace? I think space acutally checks the strides so covers you as well for non-packed cases.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Yes, thank you for catching this. It's fixed now

@@ -179,7 +182,24 @@ typename ConvPtrsType::iterator FindConvPtrByID(ConvPtrsType& conv_ptrs,
template <typename ProblemDescriptionType>
inline bool RequiresLargeTensorCKInstance(const ProblemDescriptionType& problem)

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.

I think this includes 2D & 3D, ck_grouped_conv_3d_wrw_impl.cpp will hit this as well.
I imagine this is intended, but the tests don't cover it, and it's not in the description of the change.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Yes, intended. Comments and some tests were added.

@BrianHarrisonAMD BrianHarrisonAMD 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.

Some minor comments to resolve before merge.
Might also want to update the changelog to include this fix.

JonathanLichtnerAMD and others added 3 commits August 17, 2026 22:17
RequiresLargeTensorCKInstance's element-count check used GetElementSize()
(the packed element count), which ignores strides. The int32 overflow it
guards against is on the maximum flat element *offset* a CK instance
indexes, which for a non-packed tensor is bounded by GetElementSpace()
(as GetNumBytes() already uses), not the packed count. Switch the three
GetIn/GetOut/GetWeights checks to GetElementSpace() so strided problems
whose element count fits int32 but whose addressable extent does not are
also routed to a large-tensor (int64) instance.

Also document that the gate is rank-agnostic and covers both 2D
(ConvHipImplicitGemmGroupWrwXdlops) and 3D
(ConvHipImplicitGemm3DGroupWrwXdlops) grouped backward-weights.

Co-Authored-By: Claude <noreply@anthropic.com>
Add Wrw3DElementCountOverflowGate, a CPU-only sibling of the 2D gate
test, exercising RequiresLargeTensorCKInstance on an NDHWC 3D grouped
backward-weights problem whose element extent exceeds INT_MAX while all
lengths/strides still fit int32. Confirms the gate fires for 3D wrw and
stays off for forward, backward-data, and small wrw.

Co-Authored-By: Claude <noreply@anthropic.com>
Add a ROCm 10.1.0 section documenting the grouped backward-weights
large-tensor element-count gate fix.

Co-Authored-By: Claude <noreply@anthropic.com>
@JonathanLichtnerAMD
JonathanLichtnerAMD requested a review from a team as a code owner August 17, 2026 22:17
JonathanLichtnerAMD added a commit that referenced this pull request Aug 17, 2026
…#2266)

StaticFDBSync[gfx950] has failed on 2 stale ConvHipImplicitGemmGroupWrwXdlops
perf-config entries in the gfx950 SystemDB (fdb-keys
128-7-7-1x1-1888-7-7-128-... and 240-14-14-3x3-240-14-14-16-...) every time
it's been run, on both the old container CK and fresh source-built CK --
genuinely stale entries, not a CK-version artifact. Same class of issue PR
#10043 fixed for gfx942, but no gfx950 hardware was available to run the
usual real-hardware MIOPEN_DBSYNC_CLEAN pass.

Fixed via the rocjitsu GPU-free dbsync CI (ALMIOPEN-1779): fetched a
source-CK TheRock gfx950 artifact, ran StaticFDBSync[gfx950] under rocjitsu
with MIOPEN_DEVICE_ARCH=gfx950 MIOPEN_DBSYNC_CLEAN=1 (rewrites the find/perf
DB files in place, stripping only entries that fail applicability/perf-config
validation), then verified a fresh process re-reading the cleaned files now
passes with 0 failures. Diffed the cleaned files against the committed
versions: exactly the 2 known-bad fdb-keys changed, with only the
ConvHipImplicitGemmGroupWrwXdlops entry removed from each -- every other
solver's entry for those keys is untouched.

Re-enables gfx950 dbsync in TheRock CI (test/therock/test_miopen.py),
removing the TODO(rocm-libraries#2266) exclusion -- dbsync now runs and
passes for gfx950 like every other family.

ALMIOPEN-1779

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
@HereThereBeDragons
HereThereBeDragons merged commit 1a14df1 into develop Aug 18, 2026
99 of 107 checks passed
@HereThereBeDragons
HereThereBeDragons deleted the users/jlichtne/fix-wrw-large-tensor-elemcount branch August 18, 2026 16:14
JonathanLichtnerAMD added a commit that referenced this pull request Aug 19, 2026
…#2266)

StaticFDBSync[gfx950] flagged 2 stale ConvHipImplicitGemmGroupWrwXdlops entries in the gfx950
SystemDB (fdb-keys 128-7-7-1x1-1888-7-7-128-0x0-1x1-1x1-0-NCHW-FP16-W and
240-14-14-3x3-240-14-14-16-1x1-1x1-1x1-0-NCHW-FP16-W_g240) whose perf config is no longer valid
for the current solver -- the same class of issue PR #10043 fixed for gfx942. Removed those 2
entries via a MIOPEN_DBSYNC_CLEAN run (verified by diff that ONLY the bad solver entry was
stripped from each of the 2 affected lines; all other solvers on those lines preserved), then
confirmed a fresh StaticFDBSync[gfx950] run passes clean. Fix produced GPU-free via rocjitsu (no
gfx950 hardware required).

With the DB clean, re-enable gfx950 dbsync on TheRock: drop the gfx950-dcgpu *DBSync* exclusion
and the TODO(rocm-libraries#2266) guard in test_miopen.py.

ALMIOPEN-1779, rocm-libraries#2266

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 20, 2026
The rebase onto develop pulled in #10043 ("require large-tensor CK instance for grouped wrw on
>INT_MAX element count"), which changed RequiresLargeTensorCKInstance() so large-tensor
ConvHipImplicitGemmGroupWrwXdlops problems now require the int64 CK instance. That invalidated 130
previously int32-tuned perf configs in the gfx950 SystemDB (StaticFDBSync[gfx950] flagged them all
in run 32301189658). #10043 cleaned gfx942's DB but not gfx950's; the earlier 2-entry clean here
predated #10043 so it only caught 2.

Re-cleaned via the rocjitsu MIOPEN_DBSYNC_CLEAN process against this PR's post-#10043 gfx950
source-CK artifact (clean run 32327440705): removed the 130 invalid entries in place, then a fresh
StaticFDBSync[gfx950] run PASSED. Diff-verified against the committed DBs: exactly 130 lines changed
in each of gfx950100.HIP.fdb.txt / gfx950100.db.txt, each removing ONLY the
ConvHipImplicitGemmGroupWrwXdlops entry (0 keys added/removed wholesale, no other solver's entry
altered). Recompressed bzip2 -9, round-trip verified byte-identical.

ALMIOPEN-1779, rocm-libraries#2266

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 20, 2026
…acle

Root cause (validated, prototype run 32404021201): CK's grouped-wrw auto-split-k ("+-1") derives
k_batch from max_capacity = GetMaxOccupancy() * DeviceProperties.num_cu_, where
num_cu_ = hipGetDeviceProperties().multiProcessorCount -- read straight from the device, bypassing
MIOpen's Handle CU-override. rocjitsu's stock cdna configs declare a minimal (~4 CU) device for
fast functional emulation, so under rocjitsu the auto k_batch is wrong and StaticFDBSync
false-flags valid small grouped-wrw perf configs (the 3 gfx942 / 14 gfx950 "+-1" entries that are
valid on real hardware -- #10043's real-MI300X run left the gfx942 ones untouched).

Fix (config-only, no CK/rocjitsu-source change): before the run, rewrite the cdna config's
device-block CU fields to the arch's real CU count (gfx942->304, gfx950->256) via
therock_dbsync_cu_correct_config.py, leaving the topology untouched (no emulation-cost change).
The prototype A/B'd this on real gfx942 dbsync: stock config -> 3 invalid; CU-corrected -> 0
invalid, StaticFDBSync PASSED. Genuine issues (e.g. the gfx950 >INT_MAX #10043 entries) are still
flagged -- only the occupancy false-positives clear. Wired into both the automatic gate
(therock-ci.yml) and the standalone diagnostic.

Also restores the original gfx950 SystemDB: the earlier clean here removed 130 entries, but 14 of
those were these occupancy false-positives (wrongly stripped). Reverting to develop's DB; a
follow-up commit re-cleans only the ~116 genuine >INT_MAX large-tensor entries under the
CU-corrected config.

ALMIOPEN-1779

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 21, 2026
Temporarily vendor the real-gfx950-hardware SystemDB clean from PR #11130
(users/jlichtne/ALMIOPEN-2471-clean-gfx950-dbs) so this PR's rocjitsu gfx950 dbsync gate can run
against the correct DB. #11130 removed 130 genuinely-invalid grouped-wrw entries (116 >INT_MAX
#10043 large-tensor + 14 small); to be dropped once #11130 lands on develop and this branch rebases.

Ground-truth comparison (real HW vs my rocjitsu CU-corrected re-clean): real HW removed 130, my
rocjitsu removed 132 -- the 2 extra (128-7-7-1x1-1888-7-7-128 FP16, 240-14-14-3x3-...-g240) are
residual rocjitsu occupancy false-positives that real HW keeps (valid). CU-correction cleared the
gfx942 false-positives and 14/16 of gfx950's; these 2 persist.

ALMIOPEN-1779

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 23, 2026
…#2266)

StaticFDBSync[gfx950] flagged 2 stale ConvHipImplicitGemmGroupWrwXdlops entries in the gfx950
SystemDB (fdb-keys 128-7-7-1x1-1888-7-7-128-0x0-1x1-1x1-0-NCHW-FP16-W and
240-14-14-3x3-240-14-14-16-1x1-1x1-1x1-0-NCHW-FP16-W_g240) whose perf config is no longer valid
for the current solver -- the same class of issue PR #10043 fixed for gfx942. Removed those 2
entries via a MIOPEN_DBSYNC_CLEAN run (verified by diff that ONLY the bad solver entry was
stripped from each of the 2 affected lines; all other solvers on those lines preserved), then
confirmed a fresh StaticFDBSync[gfx950] run passes clean. Fix produced GPU-free via rocjitsu (no
gfx950 hardware required).

With the DB clean, re-enable gfx950 dbsync on TheRock: drop the gfx950-dcgpu *DBSync* exclusion
and the TODO(rocm-libraries#2266) guard in test_miopen.py.

ALMIOPEN-1779, rocm-libraries#2266

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 23, 2026
The rebase onto develop pulled in #10043 ("require large-tensor CK instance for grouped wrw on
>INT_MAX element count"), which changed RequiresLargeTensorCKInstance() so large-tensor
ConvHipImplicitGemmGroupWrwXdlops problems now require the int64 CK instance. That invalidated 130
previously int32-tuned perf configs in the gfx950 SystemDB (StaticFDBSync[gfx950] flagged them all
in run 32301189658). #10043 cleaned gfx942's DB but not gfx950's; the earlier 2-entry clean here
predated #10043 so it only caught 2.

Re-cleaned via the rocjitsu MIOPEN_DBSYNC_CLEAN process against this PR's post-#10043 gfx950
source-CK artifact (clean run 32327440705): removed the 130 invalid entries in place, then a fresh
StaticFDBSync[gfx950] run PASSED. Diff-verified against the committed DBs: exactly 130 lines changed
in each of gfx950100.HIP.fdb.txt / gfx950100.db.txt, each removing ONLY the
ConvHipImplicitGemmGroupWrwXdlops entry (0 keys added/removed wholesale, no other solver's entry
altered). Recompressed bzip2 -9, round-trip verified byte-identical.

ALMIOPEN-1779, rocm-libraries#2266

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 23, 2026
…acle

Root cause (validated, prototype run 32404021201): CK's grouped-wrw auto-split-k ("+-1") derives
k_batch from max_capacity = GetMaxOccupancy() * DeviceProperties.num_cu_, where
num_cu_ = hipGetDeviceProperties().multiProcessorCount -- read straight from the device, bypassing
MIOpen's Handle CU-override. rocjitsu's stock cdna configs declare a minimal (~4 CU) device for
fast functional emulation, so under rocjitsu the auto k_batch is wrong and StaticFDBSync
false-flags valid small grouped-wrw perf configs (the 3 gfx942 / 14 gfx950 "+-1" entries that are
valid on real hardware -- #10043's real-MI300X run left the gfx942 ones untouched).

Fix (config-only, no CK/rocjitsu-source change): before the run, rewrite the cdna config's
device-block CU fields to the arch's real CU count (gfx942->304, gfx950->256) via
therock_dbsync_cu_correct_config.py, leaving the topology untouched (no emulation-cost change).
The prototype A/B'd this on real gfx942 dbsync: stock config -> 3 invalid; CU-corrected -> 0
invalid, StaticFDBSync PASSED. Genuine issues (e.g. the gfx950 >INT_MAX #10043 entries) are still
flagged -- only the occupancy false-positives clear. Wired into both the automatic gate
(therock-ci.yml) and the standalone diagnostic.

Also restores the original gfx950 SystemDB: the earlier clean here removed 130 entries, but 14 of
those were these occupancy false-positives (wrongly stripped). Reverting to develop's DB; a
follow-up commit re-cleans only the ~116 genuine >INT_MAX large-tensor entries under the
CU-corrected config.

ALMIOPEN-1779

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 23, 2026
Temporarily vendor the real-gfx950-hardware SystemDB clean from PR #11130
(users/jlichtne/ALMIOPEN-2471-clean-gfx950-dbs) so this PR's rocjitsu gfx950 dbsync gate can run
against the correct DB. #11130 removed 130 genuinely-invalid grouped-wrw entries (116 >INT_MAX
#10043 large-tensor + 14 small); to be dropped once #11130 lands on develop and this branch rebases.

Ground-truth comparison (real HW vs my rocjitsu CU-corrected re-clean): real HW removed 130, my
rocjitsu removed 132 -- the 2 extra (128-7-7-1x1-1888-7-7-128 FP16, 240-14-14-3x3-...-g240) are
residual rocjitsu occupancy false-positives that real HW keeps (valid). CU-correction cleared the
gfx942 false-positives and 14/16 of gfx950's; these 2 persist.

ALMIOPEN-1779

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 23, 2026
The header claimed gfx942's 3 originally-flagged entries were "genuinely stale... not a
CK-version-skew artifact" -- the opposite is true. They were a rocjitsu config bug (stock cdna3
declaring ~4 CU vs real 304, feeding a wrong CK auto-split-k occupancy calc), fixed by
CU-correction and confirmed valid on real gfx942 hardware by PR #10043. gfx942 is now confirmed
green end-to-end (rebase past #10043 + CU-correction, real CI run).

Replaces the stale investigation-diary narrative (old stall/thread-count history, no longer
relevant) with the current, accurate status for both arches, including gfx950's 2 still-unresolved
residual entries (a CK-version-drift hypothesis did not survive retest -- same entries, same
flagged CK instances, still fail with the identified fix commit present in the tree).

ALMIOPEN-1779

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 23, 2026
…entries, re-enable dbsync (#2266)

StaticFDBSync[gfx950] flags 132 ConvHipImplicitGemmGroupWrwXdlops / ConvHipImplicitGemm3DGroupWrwXdlops
entries whose committed perf config is invalid against current CK: 116 are the large-tensor
(>INT_MAX element-space) class PR #10043 fixed for gfx942 (auto-split-k now requires the int64 CK
instance); the other 16 are smaller auto-split-k ("+-1") entries where CK's occupancy-derived
split-k value no longer matches what was tuned.

All 132 were confirmed genuinely invalid -- not a rocjitsu artifact -- via real gfx950 hardware
running this same pipeline's actual TheRock-packaged artifact: TheRock's shipping toolchain
allocates fewer VGPRs for the borderline kernels than an earlier toolchain used for a prior
real-HW clean, which raises occupancy, raises the auto-split-k value, and genuinely invalidates
the config on TheRock's real compiler. (Two of the 16 looked like possible rocjitsu false
positives for a while; ruled in as genuine only after matching the toolchain, not just the CK
source commit or "ran on real hardware" -- see ALMIOPEN-1779 history.)

Cleaned via the rocjitsu MIOPEN_DBSYNC_CLEAN process (CU-corrected config, see the CI mechanism
commit) against this pipeline's own TheRock-packaged gfx950 artifact: removed the 132 entries in
place, then a fresh StaticFDBSync[gfx950] run PASSED. Diff-verified against the previously
committed DBs: exactly 132 lines changed in each of gfx950100.HIP.fdb.txt / gfx950100.db.txt, each
removing ONLY the ConvHip(3D)?GroupWrwXdlops entry (0 keys added/removed wholesale, no other
solver's entry altered). Recompressed bzip2 -9, round-trip verified byte-identical.

With the DB clean, re-enable gfx950 dbsync on TheRock: drop the gfx950-dcgpu *DBSync* exclusion
and the TODO(rocm-libraries#2266) guard in test_miopen.py.

ALMIOPEN-1779, rocm-libraries#2266

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 24, 2026
…entries

StaticFDBSync[gfx950] flags 132 ConvHipImplicitGemmGroupWrwXdlops / ConvHipImplicitGemm3DGroupWrwXdlops
entries whose committed perf config is invalid against current CK: 116 are the large-tensor
(>INT_MAX element-space) class PR #10043 fixed for gfx942 (auto-split-k now requires the int64 CK
instance); the other 16 are smaller auto-split-k ("+-1") entries where CK's occupancy-derived
split-k value no longer matches what was tuned.

All 132 were confirmed genuinely invalid -- not a rocjitsu artifact -- via real gfx950 hardware
running this same pipeline's actual TheRock-packaged artifact: TheRock's shipping toolchain
allocates fewer VGPRs for the borderline kernels than an earlier toolchain used for a prior
real-HW clean, which raises occupancy, raises the auto-split-k value, and genuinely invalidates
the config on TheRock's real compiler. (Two of the 16 looked like possible rocjitsu false
positives for a while; ruled in as genuine only after matching the toolchain, not just the CK
source commit or "ran on real hardware" -- see ALMIOPEN-1779 history.)

Cleaned via the rocjitsu MIOPEN_DBSYNC_CLEAN process (CU-corrected config, see the CI mechanism
commit) against this pipeline's own TheRock-packaged gfx950 artifact: removed the 132 entries in
place, then a fresh StaticFDBSync[gfx950] run PASSED. Diff-verified against the previously
committed DBs: exactly 132 lines changed in each of gfx950100.HIP.fdb.txt / gfx950100.db.txt, each
removing ONLY the ConvHip(3D)?GroupWrwXdlops entry (0 keys added/removed wholesale, no other
solver's entry altered). Recompressed bzip2 -9, round-trip verified byte-identical.

Does NOT touch test/therock/test_miopen.py's gfx950-dcgpu *DBSync* exclusion: that script is not
actually invoked by TheRock's real-hardware CI for MIOpen (TheRock's fetch_test_configurations.py
always uses the generic test_runner.py for the "miopen" project, never test_miopen.py), so that
file has no effect on TheRock's real test path. The actual gate for gfx950 dbsync on real hardware
is projects/miopen/test/gtest/test_categories.yaml's own exclude_gpu_gfx950 entry (still present,
unrelated to this PR's GPU-free rocjitsu gate) -- left untouched here as a separate decision.

ALMIOPEN-1779, rocm-libraries#2266

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 24, 2026
…entries

StaticFDBSync[gfx950] flags 132 ConvHipImplicitGemmGroupWrwXdlops / ConvHipImplicitGemm3DGroupWrwXdlops
entries whose committed perf config is invalid against current CK: 116 are the large-tensor
(>INT_MAX element-space) class PR #10043 fixed for gfx942 (auto-split-k now requires the int64 CK
instance); the other 16 are smaller auto-split-k ("+-1") entries where CK's occupancy-derived
split-k value no longer matches what was tuned.

All 132 were confirmed genuinely invalid -- not a rocjitsu artifact -- via real gfx950 hardware
running this same pipeline's actual TheRock-packaged artifact: TheRock's shipping toolchain
allocates fewer VGPRs for the borderline kernels than an earlier toolchain used for a prior
real-HW clean, which raises occupancy, raises the auto-split-k value, and genuinely invalidates
the config on TheRock's real compiler. (Two of the 16 looked like possible rocjitsu false
positives for a while; ruled in as genuine only after matching the toolchain, not just the CK
source commit or "ran on real hardware" -- see ALMIOPEN-1779 history.)

Cleaned via the rocjitsu MIOPEN_DBSYNC_CLEAN process (CU-corrected config, see the CI mechanism
commit) against this pipeline's own TheRock-packaged gfx950 artifact: removed the 132 entries in
place, then a fresh StaticFDBSync[gfx950] run PASSED. Diff-verified against the previously
committed DBs: exactly 132 lines changed in each of gfx950100.HIP.fdb.txt / gfx950100.db.txt, each
removing ONLY the ConvHip(3D)?GroupWrwXdlops entry (0 keys added/removed wholesale, no other
solver's entry altered). Recompressed bzip2 -9, round-trip verified byte-identical.

Does NOT touch test/therock/test_miopen.py's gfx950-dcgpu *DBSync* exclusion: that script is not
actually invoked by TheRock's real-hardware CI for MIOpen (TheRock's fetch_test_configurations.py
always uses the generic test_runner.py for the "miopen" project, never test_miopen.py), so that
file has no effect on TheRock's real test path. The actual gate for gfx950 dbsync on real hardware
is projects/miopen/test/gtest/test_categories.yaml's own exclude_gpu_gfx950 entry (still present,
unrelated to this PR's GPU-free rocjitsu gate) -- left untouched here as a separate decision.

ALMIOPEN-1779, rocm-libraries#2266

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
JonathanLichtnerAMD added a commit that referenced this pull request Aug 25, 2026
…entries

StaticFDBSync[gfx950] flags 132 ConvHipImplicitGemmGroupWrwXdlops / ConvHipImplicitGemm3DGroupWrwXdlops
entries whose committed perf config is invalid against current CK: 116 are the large-tensor
(>INT_MAX element-space) class PR #10043 fixed for gfx942 (auto-split-k now requires the int64 CK
instance); the other 16 are smaller auto-split-k ("+-1") entries where CK's occupancy-derived
split-k value no longer matches what was tuned.

All 132 were confirmed genuinely invalid -- not a rocjitsu artifact -- via real gfx950 hardware
running this same pipeline's actual TheRock-packaged artifact: TheRock's shipping toolchain
allocates fewer VGPRs for the borderline kernels than an earlier toolchain used for a prior
real-HW clean, which raises occupancy, raises the auto-split-k value, and genuinely invalidates
the config on TheRock's real compiler. (Two of the 16 looked like possible rocjitsu false
positives for a while; ruled in as genuine only after matching the toolchain, not just the CK
source commit or "ran on real hardware" -- see ALMIOPEN-1779 history.)

Cleaned via the rocjitsu MIOPEN_DBSYNC_CLEAN process (CU-corrected config, see the CI mechanism
commit) against this pipeline's own TheRock-packaged gfx950 artifact: removed the 132 entries in
place, then a fresh StaticFDBSync[gfx950] run PASSED. Diff-verified against the previously
committed DBs: exactly 132 lines changed in each of gfx950100.HIP.fdb.txt / gfx950100.db.txt, each
removing ONLY the ConvHip(3D)?GroupWrwXdlops entry (0 keys added/removed wholesale, no other
solver's entry altered). Recompressed bzip2 -9, round-trip verified byte-identical.

Does NOT touch test/therock/test_miopen.py's gfx950-dcgpu *DBSync* exclusion: that script is not
actually invoked by TheRock's real-hardware CI for MIOpen (TheRock's fetch_test_configurations.py
always uses the generic test_runner.py for the "miopen" project, never test_miopen.py), so that
file has no effect on TheRock's real test path. The actual gate for gfx950 dbsync on real hardware
is projects/miopen/test/gtest/test_categories.yaml's own exclude_gpu_gfx950 entry (still present,
unrelated to this PR's GPU-free rocjitsu gate) -- left untouched here as a separate decision.

ALMIOPEN-1779, rocm-libraries#2266

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
yassinsolim added a commit that referenced this pull request Aug 25, 2026
Develop commit 1a14df1 (#10043, ALMIOPEN-2367) tightened
ConvHipImplicitGemmGroupWrwXdlops so it requires a large-tensor CK
instance once GetElementSpace() exceeds INT_MAX. That gate is
solver-level and rank-agnostic, so it applies to every architecture,
but #10043 only cleaned the gfx942130 system DB and noted in its own
message that other architectures likely carry the same stale rows.

The gfx1151 and gfx1200 databases added by this PR were tuned before
that gate existed, so they still record a non-large (int32) CK instance
for wrw problems above INT_MAX elements. That is precisely the
silently-wrong combination the gate now rejects, and StaticFDBSync
fails on it.

Removed, find DB and paired perf DB:

  gfx115128  1 entry   64-445-445-7x7-3-889-889-256-...-NCHW-FP16-W
  gfx120020  5 entries 64-445-445-7x7-3-889-889-208-...-NCHW-FP16-W
                       64-597-597-7x7-3-1193-1193-{112,114,120,128}-...

Only ConvHipImplicitGemmGroupWrwXdlops is stripped from each affected
key. Every one retains GemmWrwUniversal, so no problem loses its
find/perf entry and none falls back to a naive solver.

Confidence differs between the two architectures, deliberately noted:

  gfx1151 is confirmed by CI. Smoke/CPU_DBSync_NONE.StaticFDBSync/8
  with GetParam() = ("gfx1151", 0) failed on the Windows gfx1151 shard
  of this PR's own build with "Solver is not applicable" and "Invalid
  perf config found" on exactly that key.

  gfx1200 is derived by element-count arithmetic over the find DB, not
  observed. No Windows gfx1200 shard exists in this repo's CI, so those
  five entries are not exercised by any check. They match the gfx1151
  case exactly in shape and solver, but a reviewer should treat them as
  reasoned rather than measured.

Performance note. On gfx1151 the removed entry was also the slower of
the two (2158.87 ms versus 1459.64 ms for GemmWrwUniversal), so removal
costs nothing. On gfx1200 the removed entries were faster on paper
(250-306 ms versus 383-445 ms), so those five shapes get roughly 1.5x
slower. Those timings came from the int32 kernel that returns wrong
results above INT_MAX, which is the reason the gate rejects it, so the
comparison is not like for like.

The upstream MIOPEN_DBSYNC_CLEAN pass could not be used to produce this
change: the dbsync tooling's default container pins CK at 2026-04-07 and
no longer compiles MIOpen's ck_impl wrappers from a develop-based branch
(index_t versus long_index_t array width mismatch introduced by the
widen_runtime_index_to_long change in #10547). Entries were removed by
the same rule the gate applies, with find and perf edits verified 1:1.
JonathanLichtnerAMD added a commit that referenced this pull request Aug 25, 2026
JIRA ID: ALMIOPEN-2471

## Motivation

PRs #10043 ("require large-tensor CK instance for grouped wrw on
>INT_MAX element count") and #10547 ("prevent int32 overflow in tensor
descriptor element space size") tightened the
`RequiresLargeTensorCKInstance()` gating logic in
`projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp` for grouped
backward-weights (WrW) CK convolution instances, fixing an int32
overflow bug for large tensors. That gating change invalidated some
previously-tuned SystemDB entries for the
`ConvHipImplicitGemmGroupWrwXdlops` solver: perf configs tuned before
the gate landed are no longer applicable now that certain large-tensor
problem shapes are routed away from that solver.

#10043 already refreshed gfx942's SystemDB (find-db/perf-db) to drop the
entries invalidated by this gating change, but did not touch gfx950,
because there was no gfx950 CI hardware available at the time to run
MIOpen's `test_db_sync` GTest (`projects/miopen/test/gtest/db_sync.cpp`)
validation. This PR is that overdue follow-up for gfx950, verified by
hand on a real gfx950 (MI355X) node since no CI coverage exists for this
architecture yet.

A second, unrelated batch of 2 entries was cleaned in a
follow-up commit. Those 2 are not stale due to the int32-overflow gating above; they
turned out to be a compiler-toolchain-dependent validity issue.

## Technical Details

No source code changes are included. This PR is purely a data-file fix that regenerates two
SystemDB blobs for gfx950:

- `projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2` (find-db,
bzip2-compressed)
- `projects/miopen/src/kernels/gfx950100.db.txt.bz2` (perf-db,
bzip2-compressed)

The cleaning mechanism used is MIOpen's own `MIOPEN_DBSYNC_CLEAN` env
var support in `test_db_sync`: when set, the `StaticFDBSync` test checks
each find-db entry's `IsApplicable()`/perf-config validity against the
current (already-fixed) solver code, and if a stale/invalid entry is
found, it removes only that specific solver's segment from the affected
find-db/perf-db line in place. This is a precision, applicability-driven
removal — not a blanket "delete this solver everywhere" script edit — so
other solvers sharing the same line (e.g.
`ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC`, `GemmWrwUniversal`) are
left untouched.



The gfx950 clean is now **132 entries total** (130 from the original
int32-overflow-gating stale set, 2 from the
compiler-toolchain-dependent set).

## Risk Assessment

**Risk Level: 🟢 Low**

This change only removes previously-invalid entries on one specific GPU
architecture (gfx950) — 130 stale from the int32-overflow gating fix,
plus 2 genuinely invalid under the production compiler toolchain. No
code paths change, and no other architecture's SystemDB is touched.

### Impacted Components

- `projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2` (find-db for
gfx950)
- `projects/miopen/src/kernels/gfx950100.db.txt.bz2` (perf-db for
gfx950)
- Runtime effect limited to MIOpen's find-db lookup for grouped
backward-weights convolutions on gfx950 (MI355X) hardware.

### Potential Side Effects

For the 132 affected problem-shape entries, MIOpen will no longer offer
`ConvHipImplicitGemmGroupWrwXdlops` as a candidate solver via the
find-db fast path on gfx950. This is the intended effect (all 132
entries are invalid — 130 under the current gating logic, 2 under the
production compiler toolchain). Every other solver previously present on
those same find-db lines (e.g.
`ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC`, `GemmWrwUniversal`)
remains present and untouched, so MIOpen falls back to the next-best
already-listed solver for those shapes rather than losing find-db
coverage entirely. No behavioral change is expected for any problem
shape outside these 132 entries, and no other architecture is affected.

### Mitigation Steps

- Verified via `test_db_sync` (`StaticFDBSync[gfx950]`) both before and
after each clean pass, confirming the exact failure set matches the
exact set removed, and that the mutated DB passes cleanly afterward —
for both the original 130-entry pass and the follow-up 2-entry pass.
- Diffed mutated files against originals line-by-line to confirm only
the targeted solver segment was stripped from each affected line, with
no unrelated lines touched, for both passes.
- Byte-for-byte round-trip verification of the recompressed `.bz2` blobs
against the underlying mutated text before committing, to rule out
compression-introduced corruption, for both passes.
- The follow-up 2-entry clean was run against TheRock's actual
production-toolchain build (not a manually-built equivalent), to make
sure the validity determination matches what ships, not just a local
build.
- Revert plan: trivial — this is a two-file data-only change; reverting
either commit restores the corresponding prior gfx950 SystemDB state.

## Test Plan

Manual verification on a real gfx950 (MI355X) node using MIOpen's
`test_db_sync` GTest (`projects/miopen/test/gtest/db_sync.cpp`), since
no gfx950 CI hardware exists yet to automate this:

1. Ran `StaticFDBSync[gfx950]` pre-clean to establish the baseline
failure set (130 entries).
2. Ran `StaticFDBSync[gfx950]` with `MIOPEN_DBSYNC_CLEAN=1` to perform
the in-place removal.
3. Ran `StaticFDBSync[gfx950]` again post-clean (without the clean flag)
to confirm the test now passes.
4. Manually diffed the mutated find-db/perf-db text files against the
originals to validate the scope of the edits.
5. Verified the recompressed `.bz2` artifacts decompress to
byte-identical content as the mutated text files.
6. **Follow-up:** repeated steps 1–5 for the additional 2 entries, using
TheRock's actual production-toolchain build (downloaded via
`install_rocm_from_artifacts.py`) rather than a local build, run
natively (no rocjitsu) on the same real gfx950 hardware.

No new automated tests were added — this PR only regenerates existing
SystemDB data files; there is no code path to unit test, and the
validation is exactly what `test_db_sync` (an existing GTest) already
performs when pointed at this data.

## Test Result

- Pre-clean `StaticFDBSync[gfx950]`: failed with exactly 130
`ConvHipImplicitGemmGroupWrwXdlops` entries flagged as
no-longer-applicable/invalid perf configs.
- `MIOPEN_DBSYNC_CLEAN=1` run: removed exactly those 130 entries, in
place, from the find-db/perf-db lines.
- Post-clean `StaticFDBSync[gfx950]` (no clean flag): passed cleanly at
that point.
- Manual diff of mutated vs. original files: confirmed all 130 changes
are in-place line edits (no lines fully added/removed/emptied); every
other solver on the affected lines was preserved untouched.
- **Follow-up (2 additional entries):** re-ran `StaticFDBSync[gfx950]`
against TheRock's production-toolchain artifact — confirmed exactly
these 2 entries (and only these 2) were now flagged invalid;
`MIOPEN_DBSYNC_CLEAN=1` removed them in place; a fresh non-clean run
afterward passed with 0 invalid entries; diff confirmed only the
targeted solver segment was stripped from each of the 2 lines, with all
other solvers preserved; recompressed and round-trip verified.
- Known limitation: this validation was performed manually because no
gfx950 CI hardware is currently available to run `test_db_sync`
automatically; follow-up work to add gfx950 CI coverage for this test
would let future gating/toolchain changes be validated (and SystemDBs
refreshed) without manual intervention.
yassinsolim added a commit that referenced this pull request Aug 25, 2026
Resolves a binary conflict in gfx950100.HIP.fdb.txt.bz2, where develop and
this branch cleaned the same find DB for different reasons.

  develop 98395e7 (#11130, ALMIOPEN-2471) removed stale
  ConvHipImplicitGemmGroupWrwXdlops rows invalidated by the
  RequiresLargeTensorCKInstance() tightening in #10043 / #10547. It is the
  gfx950 counterpart of the gfx942 cleanup #10043 already did.

  This branch (03ef09d) removed ConvAsmImplicitGemmGTCDynamic{Fwd,Bwd,Wrw}
  NHWC rows for problems above the byte-based INT_MAX gate added here.

The two touch different solver families, so neither supersedes the other and
git cannot merge the compressed payload. Resolved semantically rather than by
picking a side: starting from develop's file, this branch's solver removals
were re-applied on top.

Verification against the common ancestor (09e2cae):

  develop removed ConvHipImplicitGemmGroupWrwXdlops on 132 keys
  this branch removed ASM-GTC NHWC solvers on 3,366 keys
  116 keys are touched by both

  merged result: 231,048 keys
    vs develop  differs only by the three ASM-GTC NHWC solver names
    vs this branch differs only by ConvHipImplicitGemmGroupWrwXdlops
    nothing added relative to either side

  33 keys are emptied and dropped, exactly the same 33 this branch already
  dropped on its own. The combination empties no additional key, and no key
  in the result has zero solvers. All 116 doubly-touched keys retain
  GemmWrwUniversal, so no problem loses its find entry or falls back to a
  naive solver.

gfx950100.db.txt.bz2 is taken from develop unchanged: #11130 cleaned the perf
side and this branch does not modify perf DBs.
yassinsolim added a commit that referenced this pull request Sep 8, 2026
…ensors (#9568)

JIRA ID: ALMIOPEN-2367

## Motivation

The ASM-GTC NHWC solvers come from the unmaintained MISA project and
address global
tensor memory with **32-bit element indices**. They have no large-tensor
support. On any
problem whose flattened element count exceeds `INT_MAX` the index
overflows and the solver
**silently returns wrong results** — no crash, no warning; it runs,
reports a time, and
produces incorrect output.

The existing `AllTensorsDimsFitIntoInt()` check does not catch this. It
validates that each
individual length and stride fits in int32; it does **not** bound the
flattened element
count. For a packed NHWC tensor the largest stride is approximately
`elements / N`, so a
shape with a large batch can have every stride comfortably inside int32
while the total
element count overflows. That is precisely the regime this PR gates.

Because find selects a solver by timing alone and does not verify
numerics, a fast-but-wrong
solver wins selection over correct ones. Neither a version bump nor a
find-mode choice
avoids it reliably.

This PR applies the same guard to all three directions:

| solver | file |
|---|---|
| `ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC` |
`conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp` |
| `ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC` |
`conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp` |
| `ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC` |
`conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp` |

All three carried the identical hole: `AllTensorsDimsFitIntoInt()`
present, element-count
bound absent.

## Technical Details

- Added an element-count guard to `IsApplicable()` in each of the three
solvers: returns
not-applicable when `GetIn()`, `GetOut()` or `GetWeights()`
`.GetElementSize()` exceeds
  `INT_MAX`.
- **No-op for every int32-safe shape** — the predicate is false there,
so existing workloads
take the identical code path. Zero functional change outside the
overflow regime.
- Gating is the right mechanism rather than a database change: a find-db
or perf-db refresh
can only steer shapes that are enumerated in it, whereas the guard
covers every
  over-threshold shape, including ones nobody has measured.
- The heuristic still lists the solver; `IsApplicable()` returning false
makes selection fall
  through to the next candidate, which is the intended behaviour.
- Complementary to #9427, which keys CK large-tensor instances on
`!AllTensorsDimsFitIntoInt()`
(the stride regime) and does not touch MISA. This PR covers the
element-count regime.
- Relationship to the release branch: the forward half is merged to
`release/rocm-rel-7.0.2.1` as #9768, and the bwd/wrw half as #10020.
This PR is the
`develop` counterpart of both. The consolidated test file here is
**byte-identical**
to the one merged in #10020 (same blob `6bc292dfada0`), so the two
branches do not
  drift.

Files:
`src/solver/conv/conv_asm_implicit_gemm_gtc_{fwd,bwd,wrw}_nhwc.cpp`,
`test/gtest/conv_asm_gtc_nhwc_large_tensor.cpp`, `CHANGELOG.md`.

## Test Plan

- One CPU device-applicability gtest file covering all three directions,
with the shared
test-case and parameter helpers defined once. Each direction uses the
same batch-size
boundary pair, identical geometry so the element count crosses
`INT_MAX`, asserting the
solver stays applicable just below the limit and is gated off just
above:
  - `C*H*W` per sample = `1024*162*92` = 15,261,696
  - `N=140` → 2,136,637,440 ≤ `INT_MAX` → applicable
  - `N=141` → 2,151,899,136 > `INT_MAX` → gated off
  MockHandle-based, host-side, no allocation.
- End-to-end validation on hardware from from-source builds:
gfx942/MI300X and gfx950/MI355X,
  ROCm 7.0.2.1 and 7.14, MIOpenDriver `-t 1` verify.

## Test Result

Measured on MI300X against the full 42-shape set (14 customer production
geometries x
fwd/bwd/wrw), ROCm 7.0.2 base container, `FIND_MODE=2`, empty user DB,
one
verification pass per shape. Two arms: as shipped, and with all three
MISA ASM-GTC
NHWC solvers withdrawn, which is what this PR's guards do above INT_MAX.

| direction | wrong before | wrong after |
|---|---|---|
| forward | **6 of 14** | **0 of 14** |
| backward-data | 0 of 14 | 0 of 14 |
| backward-weights | **6 of 14** | **6 of 14** |

**Forward: the gate resolves it.** All six over-INT_MAX shapes move from
solver 107
to CK 137 and verify. It is also not a performance cost - across all 14
forward shapes
the gated arm is 1.12x faster, because the shapes it rescues were
running a kernel that
was both slow and wrong. Individual rescued shapes improve by up to
1.4x.

**Backward-data: no defect was observed on this shape set.** The bwd
guard is correct
by the same source argument as fwd and costs nothing, but these shapes
do not exercise
it. Stated explicitly rather than implied.

**Backward-weights: this gate is necessary but not sufficient.**
Withdrawing solver 110
correctly removes a solver that returns wrong answers, but selection
then falls through
to solver 156 `ConvHipImplicitGemmGroupWrwXdlops`, which is *also*
incorrect above
INT_MAX on the same shapes:

```
ungated  shape 31:  110 ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC  FAILED 0.223096
gated    shape 31:  156 ConvHipImplicitGemmGroupWrwXdlops           FAILED 0.213800
gated    shape 33:  156 ConvHipImplicitGemmGroupWrwXdlops           FAILED 0.640220
gated    shape 29:  156 ConvHipImplicitGemmGroupWrwXdlops           Verifies OK 5.5e-04
```

Shape 29 uses the same solver and passes, so 156 is sound below the
threshold and
breaks above it - the same failure signature as 107, in a different
solver. Closing
backward-weights therefore also requires #10043, which keys CK
large-tensor instances
for grouped wrw. The two changes are complementary, not alternatives.

**Supporting controls.** A paired control isolates the mechanism: at N=1
with a stride
that also overflows, the solver already withdraws correctly; at N=240
with small strides
and an overflowing product it is accepted and silently wrong. Same
solver, same host,
only the overflow mode differs. Forcing the solver on a gated shape
returns "no suitable
algorithm was found" on both gfx942 and gfx950.

**Database impact.** Gating bwd/wrw orphaned 606 find keys in
`gfx942130` whose recorded
winner is now inapplicable, which `StaticFDBSync` correctly reported.
Every one of those
keys has an output tensor exceeding INT_MAX, so those rows were
selecting a solver that
computes them incorrectly. Removed by exact-pair prune taken from the
DBSync report
(606/606 keys matched, 0 absent, no find lines emptied, 234 orphaned
perf lines dropped);
forward ASM-GTC entries are untouched at 79,172 before and after. Other
architectures and
the non-304-CU gfx942 variants need the equivalent pass on matching
hardware.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
yassinsolim added a commit that referenced this pull request Sep 15, 2026
…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 added a commit that referenced this pull request Sep 16, 2026
…abase

Smoke/CPU_DBSync_NONE.StaticFDBSync/11 fails on gfx1201 with 638 invalid
perf configs. They fall into two disjoint groups.

409 are ConvHipImplicitGemmGroupBwdXdlops entries carrying a split-K
suffix above 1. None of them exceeds the large-tensor threshold.

229 are entries whose tensor byte span exceeds INT_MAX. MIOpen gates CK
off there because CK forms byte offsets in 32-bit arithmetic and returns
silently wrong results. #10043 introduced that guard for grouped wrw on
element count and #11585 extended it to byte span across all directions.
#11585 cleaned the CDNA system databases it invalidated (gfx942130,
gfx942e4, gfx950100) but not the Navi ones.

The split is exact. Of the 638 failing find-db keys, 409 are under the
threshold and 229 are over it, and every one of the 229 is over. An
earlier revision of this branch removed only the 409, because it was
produced on a build that predates #11585.

The removal list was generated by running the test with
MIOPEN_DBSYNC_CLEAN=1 on a gfx1201 machine against a develop build, so
it comes from MIOpen's own predicate. A clean re-run afterwards reports
zero invalid configs and the test passes.
yassinsolim added a commit that referenced this pull request Sep 21, 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 21, 2026
…abase (#12088)

JIRA ID: ROCM-29026

## Motivation

`Smoke/CPU_DBSync_NONE.StaticFDBSync/11` fails on gfx1201 with 638
invalid perf configs. MIOpen cannot resolve them, logs "Invalid config
loaded from Perf Db", and falls back to an untuned default.

They fall into two disjoint groups.

**409 split-K entries.** `ConvHipImplicitGemmGroupBwdXdlops` perf
configs carrying a split-K suffix above 1, with kbatch values of 2, 4,
8, 16 and 32. None of these exceeds the large-tensor threshold, so the
split-K suffix is the only thing distinguishing them from the 32347
backward-data entries that still resolve.

**229 large-tensor entries.** Their tensor byte span exceeds INT_MAX.
MIOpen gates CK off there because CK forms byte offsets in 32-bit
arithmetic and returns 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 also fixes the gfx1201 instance of #11065. The perf config named in
that report, `DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3<64, 32,
32, 128, 8, 8, Filter1x1Stride1Pad0, 16, 16, 2, 1, 8, 8, 1, 1>+2` on key
`960-7-7-1x1-160-7-7-1-0x0-1x1-1x1-0-NCHW-BF16-B`, is one of the 409.

## Technical Details

The split between the two groups is exact. Of the 638 failing find-db
keys, 409 are under the INT_MAX byte threshold and 229 are over it, with
no overlap in either direction.

| group | keys | over INT_MAX | solver |
|---|---|---|---|
| split-K | 409 | 0 | `ConvHipImplicitGemmGroupBwdXdlops` |
| large-tensor | 229 | 229 | 165 Bwd, 64 Fwd |

Note that kbatch 1 appears in both descriptions without contradiction.
All 409 split-K entries carry kbatch above 1. The 165 backward-data
entries in the large-tensor group carry kbatch 1 and fail for the byte
span reason instead.

An earlier revision of this branch removed only the 409, because it was
produced against `therock-dist-linux-multiarch-tests-10.1.0a20260822`,
which predates `#11585`. Re-running on develop surfaces the other 229.
The 409 are a strict subset of the 638.

Perf database entries removed: 611, of which 547 are
`ConvHipImplicitGemmGroupBwdXdlops` and 64 are
`ConvHipImplicitGemmGroupFwdXdlops`. Zero removals fall outside those
two solvers.

File deltas: the perf database loses 374 lines and the find database
loses 19. Most find-db lines keep their remaining winners and only lose
the unusable solver. The entry count and line count differ because
several find-db keys can share one perf-db key, since the perf-db key
drops the output spatial dimensions.

## Performance consequence

Removing these entries costs measured performance on some shapes, and
that cost is not recoverable by this PR.

169 of the 638 removed find-db entries were the fastest recorded winner
for their key.

For the 229 large-tensor entries this is cheap and unavoidable. 52 were
fastest, and the next best entry is a median 1.39x slower, at worst
3.56x. These entries cannot be kept at any price, because the gate
exists to prevent silently wrong results.

For the 409 split-K entries the cost is higher. 117 were fastest. 19 of
those keys lose their only entry and will fall back to a runtime find.
For the 98 that keep a survivor, the next best is a median 6.15x slower,
and 34 of them are 10x or worse. The worst case is key
`1024-1-1-7x7-256-1-1-800-0x0-1x1-1x1-0-NCHW-BF16-B`, which goes from
`ConvHipImplicitGemmGroupBwdXdlops` at 0.4843 ms to `GemmBwdRest` at
33.6636 ms, a factor of 69.5. Across those 117 keys the new fastest is
`GemmBwdRest` 95 times.

`GemmBwdRest` is not tunable, so there is no perf config that recovers
this. The only way to get the performance back is to re-tune the
affected shapes on gfx1201 so the tuner emits a currently valid
configuration along with its measured time and workspace size. Rewriting
the existing configs to kbatch 1 by hand is not a safe substitute: 389
of the 409 carry a non-zero recorded workspace produced with the split-K
configuration, median 811008 bytes and max 305135616, and split-K
workspace comes from `GetCKSplitkWorkspaceSize(ptr, split_k)`. A kbatch
1 row needs a different workspace figure, and the recorded time would no
longer describe the named configuration.

A follow-up to re-tune these 409 backward-data shapes should be filed.
Until then this PR trades measured performance on 117 keys for correct
behaviour and a passing test.

## Test Plan

Build CK and MIOpen from develop for gfx1201 and run the DBSync test on
matching hardware.

```
test_db_sync --gtest_filter='*StaticFDBSync*'
```

## Test Result

Run on `ctr-navi4x-aj71-ws10`, a gfx1201 machine reporting 64 compute
units, so MIOpen
selects the `gfx120140` database under test.

```
container  rocm/dev-ubuntu-22.04@sha256:a60cffc2d079dbdc7a54948766e21793abdfeaa2ea95af60214f975a1ce51d06
source     ROCm/rocm-libraries develop, MIOpen and CK built from the same checkout
CK         built from source; /opt/rocm/include/ck and ck_tile moved aside first, and
           MIOpen pointed at it with composable_kernel_DIR, so it binds the CK it was
           built against rather than the container's prebuilt copy
cmake      MIOPEN_TEST_DBSYNC=ON MIOPEN_TEST_DISCRETE=ON MIOPEN_USE_HIPCONV=OFF
           BUILD_DEV=OFF SKIP_KDB_TESTING=1
```

The CK binding detail matters. A MIOpen that loads a CK library from a
different revision
than it was built against reports a large number of unrelated invalid
perf configs, which
swamps the signal from these 638.

This change removes exactly 638 find-db entries and 611 perf-db entries,
and the test passes afterwards in that environment.

The reproducible claim is the delta, which an independent
reimplementation of MIOpen's
predicate also predicts at exactly 638.

Content checks on the committed files:

- pre-change tree content byte identical to develop (pdb
`49a90431c797a1b7`, fdb `cf621424ac57c7f0`, first 16 hex of sha256 of
the decompressed files)
- removal list generated by `MIOPEN_DBSYNC_CLEAN=1`, so it comes from
MIOpen's own predicate rather than a hand-built list
- removals confined to the two affected solvers, zero outside them
- no surviving find-db entry for a tunable solver is left without a
perf-db row
- 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
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
evedovelli pushed a commit that referenced this pull request Sep 23, 2026
…abase (#12088)

JIRA ID: ROCM-29026

## Motivation

`Smoke/CPU_DBSync_NONE.StaticFDBSync/11` fails on gfx1201 with 638
invalid perf configs. MIOpen cannot resolve them, logs "Invalid config
loaded from Perf Db", and falls back to an untuned default.

They fall into two disjoint groups.

**409 split-K entries.** `ConvHipImplicitGemmGroupBwdXdlops` perf
configs carrying a split-K suffix above 1, with kbatch values of 2, 4,
8, 16 and 32. None of these exceeds the large-tensor threshold, so the
split-K suffix is the only thing distinguishing them from the 32347
backward-data entries that still resolve.

**229 large-tensor entries.** Their tensor byte span exceeds INT_MAX.
MIOpen gates CK off there because CK forms byte offsets in 32-bit
arithmetic and returns 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 also fixes the gfx1201 instance of #11065. The perf config named in
that report, `DeviceGroupedConvBwdDataMultipleD_Wmma_CShuffleV3<64, 32,
32, 128, 8, 8, Filter1x1Stride1Pad0, 16, 16, 2, 1, 8, 8, 1, 1>+2` on key
`960-7-7-1x1-160-7-7-1-0x0-1x1-1x1-0-NCHW-BF16-B`, is one of the 409.

## Technical Details

The split between the two groups is exact. Of the 638 failing find-db
keys, 409 are under the INT_MAX byte threshold and 229 are over it, with
no overlap in either direction.

| group | keys | over INT_MAX | solver |
|---|---|---|---|
| split-K | 409 | 0 | `ConvHipImplicitGemmGroupBwdXdlops` |
| large-tensor | 229 | 229 | 165 Bwd, 64 Fwd |

Note that kbatch 1 appears in both descriptions without contradiction.
All 409 split-K entries carry kbatch above 1. The 165 backward-data
entries in the large-tensor group carry kbatch 1 and fail for the byte
span reason instead.

An earlier revision of this branch removed only the 409, because it was
produced against `therock-dist-linux-multiarch-tests-10.1.0a20260822`,
which predates `#11585`. Re-running on develop surfaces the other 229.
The 409 are a strict subset of the 638.

Perf database entries removed: 611, of which 547 are
`ConvHipImplicitGemmGroupBwdXdlops` and 64 are
`ConvHipImplicitGemmGroupFwdXdlops`. Zero removals fall outside those
two solvers.

File deltas: the perf database loses 374 lines and the find database
loses 19. Most find-db lines keep their remaining winners and only lose
the unusable solver. The entry count and line count differ because
several find-db keys can share one perf-db key, since the perf-db key
drops the output spatial dimensions.

## Performance consequence

Removing these entries costs measured performance on some shapes, and
that cost is not recoverable by this PR.

169 of the 638 removed find-db entries were the fastest recorded winner
for their key.

For the 229 large-tensor entries this is cheap and unavoidable. 52 were
fastest, and the next best entry is a median 1.39x slower, at worst
3.56x. These entries cannot be kept at any price, because the gate
exists to prevent silently wrong results.

For the 409 split-K entries the cost is higher. 117 were fastest. 19 of
those keys lose their only entry and will fall back to a runtime find.
For the 98 that keep a survivor, the next best is a median 6.15x slower,
and 34 of them are 10x or worse. The worst case is key
`1024-1-1-7x7-256-1-1-800-0x0-1x1-1x1-0-NCHW-BF16-B`, which goes from
`ConvHipImplicitGemmGroupBwdXdlops` at 0.4843 ms to `GemmBwdRest` at
33.6636 ms, a factor of 69.5. Across those 117 keys the new fastest is
`GemmBwdRest` 95 times.

`GemmBwdRest` is not tunable, so there is no perf config that recovers
this. The only way to get the performance back is to re-tune the
affected shapes on gfx1201 so the tuner emits a currently valid
configuration along with its measured time and workspace size. Rewriting
the existing configs to kbatch 1 by hand is not a safe substitute: 389
of the 409 carry a non-zero recorded workspace produced with the split-K
configuration, median 811008 bytes and max 305135616, and split-K
workspace comes from `GetCKSplitkWorkspaceSize(ptr, split_k)`. A kbatch
1 row needs a different workspace figure, and the recorded time would no
longer describe the named configuration.

A follow-up to re-tune these 409 backward-data shapes should be filed.
Until then this PR trades measured performance on 117 keys for correct
behaviour and a passing test.

## Test Plan

Build CK and MIOpen from develop for gfx1201 and run the DBSync test on
matching hardware.

```
test_db_sync --gtest_filter='*StaticFDBSync*'
```

## Test Result

Run on `ctr-navi4x-aj71-ws10`, a gfx1201 machine reporting 64 compute
units, so MIOpen
selects the `gfx120140` database under test.

```
container  rocm/dev-ubuntu-22.04@sha256:a60cffc2d079dbdc7a54948766e21793abdfeaa2ea95af60214f975a1ce51d06
source     ROCm/rocm-libraries develop, MIOpen and CK built from the same checkout
CK         built from source; /opt/rocm/include/ck and ck_tile moved aside first, and
           MIOpen pointed at it with composable_kernel_DIR, so it binds the CK it was
           built against rather than the container's prebuilt copy
cmake      MIOPEN_TEST_DBSYNC=ON MIOPEN_TEST_DISCRETE=ON MIOPEN_USE_HIPCONV=OFF
           BUILD_DEV=OFF SKIP_KDB_TESTING=1
```

The CK binding detail matters. A MIOpen that loads a CK library from a
different revision
than it was built against reports a large number of unrelated invalid
perf configs, which
swamps the signal from these 638.

This change removes exactly 638 find-db entries and 611 perf-db entries,
and the test passes afterwards in that environment.

The reproducible claim is the delta, which an independent
reimplementation of MIOpen's
predicate also predicts at exactly 638.

Content checks on the committed files:

- pre-change tree content byte identical to develop (pdb
`49a90431c797a1b7`, fdb `cf621424ac57c7f0`, first 16 hex of sha256 of
the decompressed files)
- removal list generated by `MIOPEN_DBSYNC_CLEAN=1`, so it comes from
MIOpen's own predicate rather than a hand-built list
- removals confined to the two affected solvers, zero outside them
- no surviving find-db entry for a tunable solver is left without a
perf-db row
- 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
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