fix(conv): gate ASM-GTC fwd/bwd/wrw NHWC solvers off int32-overflow tensors by yassinsolim · Pull Request #9568 · ROCm/rocm-libraries

yassinsolim · 2026-07-17T19:57:33Z

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 feat(conv): large-tensor grouped CK xdlops conv + split_k workspace backport #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 fix(conv): gate ASM-GTC fwd NHWC solver off int32-overflow tensors #9768, and the bwd/wrw half as fix(conv): gate ASM-GTC bwd/wrw NHWC solvers off int32-overflow tensors #10020. This PR is the
develop counterpart of both. The consolidated test file here is byte-identical
to the one merged in fix(conv): gate ASM-GTC bwd/wrw NHWC solvers off int32-overflow tensors #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

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

therock-pr-bot · 2026-07-17T20:04:43Z

✅ 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 · 2026-07-17T20:04:45Z

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

ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC (the ASM-GTC forward solver from the unmaintained MISA project) indexes global tensor memory with 32-bit element indices and does not implement large-tensor support. Its IsApplicable() relied on AllTensorsDimsFitIntoInt(), which only checks that each individual length/stride fits in int32 -- it never bounds the flattened element count. A tensor whose total element count exceeds INT_MAX therefore overflows the kernel's 32-bit indexing and is silently computed incorrectly (observed forward verification error ~0.37 on gfx942 and gfx950). Gate the solver not-applicable when any tensor's element count exceeds INT_MAX, so a large-tensor-capable solver (e.g. grouped CK xdlops) is selected instead. This is a no-op for every int32-safe shape, so existing workloads are unaffected. Adds a CPU device-applicability unit test that shares geometry across a batch-size boundary so the element count crosses INT_MAX, asserting the solver stays applicable just below the limit and is gated off just above. Related to ROCM-27526.

The preceding commit gates ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC not-applicable when any tensor's element count exceeds INT_MAX. The shipped find-dbs still list that solver for such shapes -- it had been recorded there (in some entries as the fastest kernel) even though its 32-bit element indexing computes those shapes incorrectly. The db_sync FDB consistency test asserts every solver listed in a find-db entry is still IsApplicable(), so the now-gated entries fail Dbsync on gfx908/gfx90a/gfx942/gfx950. Prune the solver from exactly the find-db entries whose input or output element count exceeds INT_MAX, mirroring the gate and the MIOPEN_DBSYNC_CLEAN removal path (cf. #3725, "Clean db files to pass db_sync"). Entries at or below INT_MAX are left untouched; lines where it was the only listed solver are dropped. gfx90878.HIP.fdb.txt : 43 gfx90a68.HIP.fdb.txt : 102 gfx90a6e.HIP.fdb.txt : 102 gfx942130.HIP.fdb.txt : 268 gfx942e4.HIP.fdb.txt : 206 gfx950100.HIP.fdb.txt : 309 Related to ROCM-27526.

ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC shares the MISA 32-bit element indexing of the forward solver already gated in this PR. AllTensorsDimsFitIntoInt() bounds each length and stride but not the flattened element count, so a packed NHWC tensor with a large batch passes the check while the total overflows and the kernel silently returns wrong results. Apply the same guard.

ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC shares the MISA 32-bit element indexing of the forward solver already gated in this PR. AllTensorsDimsFitIntoInt() bounds each length and stride but not the flattened element count, so a packed NHWC tensor with a large batch passes the check while the total overflows and the kernel silently returns wrong results. Apply the same guard.

Mirrors the forward test: same geometry, N=140 vs N=141 brackets the INT_MAX element-count boundary, so the int32-safe case must stay applicable and the over-threshold case must be applicable on no device.

Matches PR #10020, already merged to release/rocm-rel-7.0.2.1, which carries the same bwd/wrw gate with a single test file. Taking that file verbatim keeps the two branches identical and defines the four shared helpers once instead of three times.

…ded by the consolidated file

The gates added here compared each tensor's element count against INT_MAX, but these assembly kernels address global memory with 32-bit byte offsets. For fp16 that lets shapes up to twice the safe size through, and fp32 up to four times, where they return silently wrong results. Compare GetNumBytes() instead, matching the byte-based form already used for the equivalent CK solver guard in conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp. The unit test's boundary cases move with the threshold: for the shared 1024x162x92 fp16 geometry the effective element ceiling halves, so the in-range and gated cases become N=70 (2,136,637,440 bytes) and N=71 (2,167,160,832 bytes) rather than N=140 and N=141.

…erflow gate The previous commit changed the ASM-GTC fwd/bwd/wrw NHWC solvers' overflow guard from an element-count comparison to a byte-count comparison, because these MISA-derived assembly kernels address global memory with 32-bit byte offsets, not 32-bit element offsets. develop ships no gate at all on these solvers, so its find-db files contain tuned rows that name one of these solvers as the winner for shapes whose input, output, or weight tensor now exceeds INT_MAX bytes. Those rows were already silently computing wrong results before this PR; the new gate only makes StaticFDBSync notice. Removed the exact (key, solver) pairs identified by a static byte-accurate parser that mirrors ParseProblemKey()/IsApplicable() in db_sync.cpp and the three solver files, cross-validated key-for-key (zero difference either direction) against IsApplicable() run live on gfx942/MI300X-304CU hardware in CI. Re-parsing every cleaned file with the same tool confirms zero stale rows remain. A key keeps every entry that is still valid; a line is dropped outright only when removing its stale entries would leave it empty. Per-file counts (stale pairs removed / find-db lines edited in place / lines dropped entirely): gfx90878 1210 / 1209 / 1 gfx90a68 1227 / 1202 / 25 gfx90a6e 1244 / 1228 / 16 gfx942130 3206 / 3200 / 6 gfx942e4 2654 / 2654 / 0 gfx950100 3366 / 3333 / 33 Perf-db cleanup for the same shapes is a separate commit.

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.

…9568

cderb

One issue isn't tied to any changed line, so it's up here; the rest are inline.

Scope first, since it shapes everything below: the GetWeights() clause of this guard is closing a real hole. igemm_split_batch_size() only ever looks at the activation tensors — it never reads y, x, or the filter — and weights have no batch dimension to split on, so nothing in the tree bounds them. That clause should survive in some form. What I'm questioning is the other two clauses, the constant, and a bug upstream of all of it.

An arithmetic bug in the split helper, upstream of this gate

src/include/miopen/conv/asm_implicit_gemm.hpp:269 (not in this diff):

size_t m = std::ceil(n / max_n); // m >= n * (image_size / max_tensor_size)

n is int, max_n is size_t, so the division is integer and std::ceil is a no-op on an already-truncated value. The loop is while(n % max_n != 0) — divisibility only; it never re-checks image_size against max_tensor_size.

I compiled the exact function body and ran it on this PR's geometry (fp16, C=K=1024, 162x92):

n	splits	per-slice bytes
140	1	4,273,274,880	ok (last safe)
141	1	4,303,798,272	over 4 GiB
200	1	6,104,678,400	over
280	2	4,273,274,880	ok
282	2	4,303,798,272	over

238 of n in [1,512] yield an over-4 GiB slice. For n=141: 141/140 = 1, ceil(1) = 1, max_n = 141, exit, one split. The launch guard (fwd:976 / bwd:1052 / wrw:945) measures threads, not bytes, so it doesn't catch these. Introduced in 2023 by 4c55d07fe12 (#2042); no test covers the function.

std::ceil(static_cast<double>(n) / static_cast<double>(max_n)), plus dropping the max_n = n % max_n branch, fixes the activation path for all four MISA solvers — including conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp, which calls the same helper and gets no gate here. It's also consistent with the N=240 reproducer cited in the description: at 6.8 GiB with two splits, that shape lands in this path rather than in the band the guard removes. The INT_MAX bound masks the bug for NHWC as a side effect, which may be why the fix reads as working empirically.

Combining that with the inline points, the end state I'd suggest is: fix the division; keep a weights guard, at 0xffffffff on the per-group weight size; drop the GetIn()/GetOut() clauses, which the split mechanism handles once the division is right.

This rests on static analysis plus host-side execution — I did not run a GPU. Forcing the solver at N=71 with the guard removed (predict: correct) and at N=141/200 (predict: wrong) would settle both this and the inline over-gating point.

On the six .bz2 files: I verified the scrub is clean and exactly consistent with the guard — applying the predicate to each old file reproduces each new file byte-for-byte, nothing added or retimed, no other solver touched. Not raising it separately, since it's downstream of the guard and likely reverts if the bound changes. Two things if it stays: it's larger than the description says (12,907 entries across all six arches, 4,668 of them rank-1, 81 keys emptied — vs. "606 find keys in gfx942130… no find lines emptied… forward ASM-GTC entries are untouched"), and it'd want a ### Changed CHANGELOG entry per CHANGELOG:17-18.

🤖 Review assisted by Claude Code. Findings were independently verified against the source; the split-helper table above came from compiling and running that function locally.

…/wrw Addresses two review points on #9568: - averinevg: the three IsApplicable() additions were byte-identical copy-paste. Extracted into igemm_tensor_size_exceeds_asm_gtc_addressing() in asm_implicit_gemm.hpp, which the three solvers already include. The bound now lives in one place. - cderb: the comment described a 32-bit element-count bound while the code checked GetNumBytes(). The check was the correct one; the prose was not. Comment and CHANGELOG now both say bytes. No behavioural change: same three tensors, same > INT_MAX byte threshold.

# Conflicts: # projects/miopen/src/kernels/gfx942e4.HIP.fdb.txt.bz2 # projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2

# Conflicts: # projects/miopen/src/kernels/gfx942130.HIP.fdb.txt.bz2 # projects/miopen/src/kernels/gfx942e4.HIP.fdb.txt.bz2 # projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2

cderb noted on #9568 that the find-database prune is a user-visible change to shipped databases with no CHANGELOG entry, and that it belongs under Changed rather than being implied by the Resolved Issues entry about the solver gate. Uses his audited figures: 12,907 entries across the six databases, 4,668 of them rank-1, 81 keys emptied. The PR description previously said 606 keys in gfx942130 with no lines emptied, which was wrong.

cderb

Could relax fwd/bwd bounds. Otherwise LGTM

assistant-librarian · 2026-09-09T18:46:47Z

TheRock Submodule Bump Activity

Newest first

2026-09-09 — Included in TheRock via ROCm/TheRock#8089.

…#11809) JIRA ID: ALMIOPEN-2367 ## Motivation Follow-up to #9568, which gated the three ASM-GTC NHWC solvers off tensors whose byte size exceeds INT_MAX. Review of that PR established the bound is roughly 2x too strict for forward and backward-data: those kernels address a full 4 GiB window per dispatch slice and correctly advance the base pointer per batch-split, so activations are already handled by the split mechanism rather than by the gate. The cost of the over-strict bound was not limited to solver selection. #9568 also pruned the system find-databases using that same bound as its predicate, so the affected shapes lost their tuned entries and resolve to a noticeably slower solver in immediate mode. Relaxing the guard without restoring those entries would leave exactly the shapes this change re-enables pointing at the slower path, so both parts ship together. ## Technical Details `igemm_tensor_size_exceeds_asm_gtc_addressing()` now branches by direction rather than applying one bound to all three. Forward and backward-data: the `GetIn()` and `GetOut()` clauses are dropped. The weight base is advanced per group with 64-bit arithmetic, so the per-group weight size is the only quantity that must fit the window, bounded at `0xffffffff`. Backward-weights: unchanged, all three tensors still bounded at INT_MAX. This asymmetry is deliberate. 292 of its 900 kernels ignore the batch-split dimension and compute only the first slice however many were requested, so the bound is what holds the solver to a single split. Relaxing it would turn on a silent wrong-result path. That is filed as #11805 and recorded in a comment at the gate so it is not tidied away later. Database restore: 8,978 forward and backward-data entries reinstated across the six system find-databases, 164 keys reinstated whole. Entries naming the backward-weights solver stay pruned. | arch | restored | left pruned (wrw) | |---|---|---| | gfx90878 | 795 | 415 | | gfx90a68 | 949 | 278 | | gfx90a6e | 958 | 286 | | gfx942130 | 2,224 | 982 | | gfx942e4 | 1,693 | 961 | | gfx950100 | 2,359 | 1,007 | `ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC` is intentionally out of scope. It has no large-tensor bound today, it is a Dlops rather than Xdlops kernel family, and the defect has not been reproduced against it. Adding a gate there is new behaviour rather than a relaxation and should carry its own evidence. ## Test Plan `conv_asm_gtc_nhwc_large_tensor.cpp` is extended rather than replaced: - `SmokeInt32Safe` (N=70, under the old bound) unchanged for all three directions, as a regression guard. - `SmokeActivationsOverInt32` (N=71, activations over INT_MAX, weights 18 MiB) now expects forward and backward-data to be applicable, which is the behaviour this PR changes. - `SmokeWeightsOverBound` is new: 16384x16384x3x3 fp16 weights are 4.50 GiB per group, over the window, while activations stay at 2 MiB so the weight term is isolated. Forward and backward-data must be gated off. - Backward-weights keeps `SmokeOverInt32` expecting gated at N=71, proving the direction split holds. ## Test Result Database restore verified programmatically against the pre-prune and post-prune trees: - every restored entry was present pre-prune and names only the forward or backward-data ASM-GTC solver - non-target solver entries altered: 0 - entries removed from develop's state: 0 - entries left pruned that are not backward-weights: 0 - largest per-group weight tensor among restored entries: 85.6 MiB against the 4 GiB bound, so the new bound excluded none of them - all six files pass `bzip2 -t` and a decompress/recompress round trip The restored count for gfx942130 (2,224) matches the forward and backward-data split of the DBSync stale-row list that drove the original prune (1,294 forward, 930 backward-data) exactly. ## Submission Checklist - [x] Reviewed the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests

… directions (#11585)" This reverts commit c862d1e. Source files only. The six MIOpen perf-db blobs touched by #11585 (gfx942130/gfx942e4/gfx950100 .db.txt.bz2 and .HIP.fdb.txt.bz2) are deliberately left at their develop state, so the later ASM-GTC large-tensor updates from #9568 and #11809 are preserved.

#9568 removed ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC from 1007 gfx950100 find-db records whose activations exceeded INT_MAX. Restore the 342 FP32 entries whose tensors fit the new 4 GiB bound and whose tuned perf config passes the batch-padding check (5 of them whole keys). Drop the entry from 3 FP16/BF16 records under INT_MAX whose tuned config fails that check. No other entry changes.

…t offsets; restore gfx950 FP32 rows up to 4 GiB (#13640) JIRA ID: ROCM-31718 ## What changed `ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC` (solver 110): - New check `WrwPaddedBatchFitsBufferOffsets()`. For a perf config it computes the last batch row the kernel can read (padding rows of the last batch pack plus the tail of the last gemm-k split) and requires that row to end at or below 4 GiB in both activation tensors. Workgroup `bz` reads exactly the gemm-k range `[bz * gemmk_per_wg, (bz + 1) * gemmk_per_wg)`, so there is no read margin. It is applied in `IsValidPerformanceConfig()` and `IsValid(ctx, problem)` (perf-db configs, LightGBM picks, GenericSearch), in `HeuristicInit()` (falls back to the least-padding valid config that does not wrap), and in `IsApplicable()` (rejects the problem if no such config exists). - `ComputeWrwGemmKSplit()` is factored out of `GetSolution()` so the check and the kernel launch use the same split arithmetic. - On gfx950, FP32 NCHW/NHWC activations may use the full `0xffffffff`-byte window (was INT_MAX). FP16/BF16, mixed layouts and other devices keep INT_MAX. The backward-weights weight tensor keeps INT_MAX in every case, including gfx950 FP32: no 2-4 GiB backward-weights weight tensor has been run against a GPU reference and no shipped find-db row has one. `IsApplicable()` now requires exactly one batch split, so the 292 split-blind kernels from #11805 are never asked to split. - gfx950100 find-db: 342 FP32 ASM-GTC WRW entries removed by #9568 are restored (5 as whole keys). Each one's tensors fit the new bound and its tuned perf-db config passes the new check. 2 BF16 entries under INT_MAX whose tuned config fails the check are dropped. No other entry changes. - `conv_asm_gtc_nhwc_large_tensor.cpp`: FP32 in-window and over-window cases in both NHWC and NCHW, an FP32 case whose weights sit between INT_MAX and 4 GiB and must stay declined, three batch-padding wrap cases (one FP16 under INT_MAX, one FP16 at 2 GiB, one FP32 NCHW inside the 4 GiB window), a case that covers `HeuristicInit()`'s fallback, and a no-padding control. Closes #11805: solver 110 no longer splits the batch, so the split-blind kernels are unreachable. ## Why ROCM-31718: one-GPU FP32 batch-1024 ResNet-50/101 training on MI355X regressed about 7.7x between the 0909 and 0910 nightlies. #9568 gated solver 110 off backward-weights tensors above INT_MAX and pruned its gfx950 find-db rows, so conv1 wgrad at batch 1024 fell from about 5.9 ms (solver 110) to about 2.8 s (GemmWrwUniversal). Simply raising the bound to 4 GiB was not safe. An earlier candidate gave a wrong FP16 result (relative RMS 1.00 against the GPU reference) for N=1, C=K=256, 2048x2048, 3x3. Root cause: the FP16/BF16 kernels with filter padding (`ex1`) load batch rows in packs of 4 (one config packs 8) and rely on the buffer range check to zero rows past N. The row offsets are computed with 32-bit `s_mul_i32`/`v_mul_lo_u32`, so a padding row whose offset reaches 4 GiB wraps back into the tensor and adds real data. The FP32 kernels pack one row, which is why FP32 passed at the same sizes. The same hazard exists under INT_MAX on develop today: N=1, C=K=64, 4096x2731, 3x3 FP16 (1.43 GB per tensor) gives relative RMS 0.92 with solver 110 on unmodified develop, and C=K=248, 2048x2048 (1.94 GB) gives 0.071, which MIOpenDriver's 0.082 FP16 tolerance accepts. The new check rejects both. ## What could break - Problems where solver 110's usual heuristic config would wrap now get another valid config or lose solver 110. On the shipped databases this affects only the 2 gfx950100 BF16 entries listed in the DB commit. Every find-db entry that names solver 110 in gfx90878, gfx90a68, gfx90a6e, gfx942130, gfx942e4 and gfx950100 was re-checked against the new predicate: 215,171 of them have a tuned perf-db config and all pass; the 25,599 gfx90a6e entries have no perf db and their heuristic config passes. - The check still rejects some shapes whose kernels happen to produce correct output, because any row that ends above 4 GiB counts even when the thread offsets inside it do not reach that far. - gfx950 FP32 2-4 GiB WRW problems can select solver 110 again. ## How it was checked MI355X (gfx950, 256 CU), rocm-libraries develop 52ddd47, CK built from the same commit, in a ROCm container. - Root cause, forced solver 110 with the GPU reference: N=1, C=K=256, 2048x2048, 3x3, FP16 NCHW relative RMS 1.002, FP16 NHWC 1.013, BF16 1.000. Controls with the same kernel family: C=K=128 (no wrap) 1.9e-4, FP32 C=K=128 2.8e-7, N=4 no padding rows 6.4e-3. - Unmodified develop through the normal Find path (`MIOPEN_FIND_MODE=2`) selects solver 110 for four sub-INT_MAX FP16/BF16 shapes and returns wrong weights: relative RMS 0.92, 0.072, 0.064, 0.039. Only the first fails MIOpenDriver's own check. With this PR the same Find picks ConvHipImplicitGemmGroupWrwXdlops: 2.4e-4, 3.2e-3, 1.9e-2, 2.4e-4. - All 342 restored FP32 entries, solver 110 forced with the tuned config: relative RMS at most 2.5e-6 against the GPU reference. Through Find, all 342 now select solver 110 (before: 283 GemmWrwUniversal, 31 GemmWrw1x1_stride1, 19 Winograd, 9 ConvHipImplicitGemmGroupWrwXdlops). Single-iteration kernel time is lower for all 342; median 22x, range 1.07x to 967x. - StaticFDBSync gfx950 (index 5): pass on base and on this PR (231,075 keys). - `test_conv_asm_gtc_nhwc_large_tensor`: 18/18 pass, nothing skipped. At the previous 17-test revision the same test binary against the base library failed 4 (`SmokeFp32InWindow`, `SmokeFp32NchwInWindow`, `SmokeBatchPaddingWrapsUnderInt32`, `SmokeFp32NchwPaddingFallback`), as expected; the base control was not re-run after `SmokeFp32WeightsOverInt32` was added. By the base source, `SmokeFp32NchwBatchPaddingWraps` and `SmokeFp32WeightsOverInt32` are declined there through the INT_MAX bound, and this PR declines them through the padded-batch check and the weight bound. - ResNet-50/101 backward-weights, batch 1024, NCHW, `MIOPEN_FIND_MODE=2`, median of 3 single-GPU MIOpenDriver runs, clocks not locked: | layer (FP32) | develop ms | this PR ms | solver develop -> PR | |---|---|---|---| | conv1 7x7 s2 3->64 | 2794.2 | 5.54 | 102 -> 110 | | layer1 1x1 64->256 (and downsample) | 27.09 | 2.51 | 101 -> 110 | | layer1 1x1 256->64 | 22.17 | 2.52 | 101 -> 110 | | layer2 1x1 256->128 | 22.04 | 3.60 | 101 -> 110 | | layer2 downsample 256->512 s2 | 22.57 | 3.72 | 102 -> 110 | The other 18 unique FP32 layers and all 23 FP16 layers keep solver 110 and their time (ratios 0.99 to 1.01). Sum over unique FP32 layers: 2925.0 ms to 54.8 ms. The branch was then rebased onto develop 8db16c8 for this PR. The only conflict was `CHANGELOG.md`; develop has not changed any other file this PR touches since 52ddd47. The root-cause, Find, 342-entry and ResNet numbers above were not re-run after the rebase. ### Follow-up round after averinevg's review (dropping the gemm-k margin) Same MI355X (gfx950, 256 CU, ROCm 10.2.0), built from this branch. Everything below was measured against a library built from the exact committed bytes. - The margin was dropped because the kernels do not read past their own gemm-k range. In all 900 wrw NHWC `.s` files under `src/kernels/dynamic_igemm/igemm_gtc_xdlops_nhwc{,_gfx90a,_gfx940,_gfx950}/wrw_*`: `s_knum = round_up(s_gemmk_per_wg << log2(ta1), gemm_k_per_block)`, there is one `s_kitr` setup and one body decrement, every `buffer_load` against `s_p_in`/`s_p_out` is in the prologue or in the body ahead of the offset advance with equal counts and none after the decrement, and `tensor_a_thread_lengths[1] * tensor_a_cluster_lengths[1] == gemm_k_per_block`. - Forced solver 110 against the GPU reference (`-S ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC -o 1`): the `SmokeFp32NchwPaddingFallback` shape (N=130) is now applicable and gives relative RMS 3.5e-07; N=129 of the same shape 3.5e-07; the new `SmokeFp32NchwBatchPaddingWraps` shape at N=64 2.1e-07 and at N=65 no solution. The restored FP16 find-db row runs with its tuned config `61[9]` at 2.2e-03, against 2.0e-03 for `ConvHipImplicitGemmGroupWrwXdlops` on the same shape and 2.2e-03 for the same forced config on the base library. - Still declined on gfx950, with and without the margin: N=1 C=K=256 2048x2048 3x3 in FP16 NCHW, FP16 NHWC and BF16; N=1 C=K=248 2048x2048 FP16; N=1 C=K=64 4096x2731 FP16; and both BF16 find-db rows that stay dropped. As a control that the comparison detects a wrong result, the base library forced onto the same solver gives relative RMS 0.92 for C=K=64 4096x2731 and 0.072 for C=K=248. - Database audit with the new predicate over every find-db entry that names solver 110 in gfx90878, gfx90a68, gfx90a6e, gfx942130, gfx942e4 and gfx950100 (240,770 entries): 215,171 have a tuned perf-db config and all pass, the 25,599 gfx90a6e entries have no perf db and their heuristic config passes. The same scan with the margin gives the same result. - The gfx950 window applies to the activations only. An earlier revision of this branch compared the weight tensor against the same relaxed bound, so a gfx950 FP32 NHWC shape with weights between INT_MAX and 4 GiB (N=1, C=K=8192, 14x14, 3x3, weights 2,415,919,104 bytes) became applicable on the weight term alone. The weight comparison is back on INT_MAX and `SmokeFp32WeightsOverInt32` pins it. Not tested: 8-GPU DDP and PyTorch end to end, gfx942/gfx90a/gfx908 hardware, FP16/BF16 above INT_MAX beyond the 105 shapes run (they stay gated).

yassinsolim requested review from a team as code owners July 17, 2026 19:57

github-actions Bot added documentation project: miopen labels Jul 17, 2026

therock-pr-bot Bot added the Not ready to Review label Jul 17, 2026

yassinsolim force-pushed the users/ysoliman/gate-asm-gtc-fwd-large-tensor branch from ac74d92 to 420bf8a Compare July 17, 2026 20:13

therock-pr-bot Bot removed the Not ready to Review label Jul 17, 2026

assistant-librarian Bot added the organization: ROCm label Jul 17, 2026

yassinsolim changed the title ~~fix(conv): gate ASM-GTC fwd NHWC solver off int32-overflow tensors~~ fix(conv): gate ASM-GTC fwd/bwd/wrw NHWC solvers off int32-overflow tensors Jul 31, 2026

yassinsolim added 15 commits August 19, 2026 12:24

test(conv): applicability test for the ASM-GTC bwd NHWC INT_MAX gate

48fd8c5

Mirrors the forward test: same geometry, N=140 vs N=141 brackets the INT_MAX element-count boundary, so the int32-safe case must stay applicable and the over-threshold case must be applicable on no device.

test(conv): applicability test for the ASM-GTC wrw NHWC INT_MAX gate

f9b183f

Mirrors the forward test: same geometry, N=140 vs N=141 brackets the INT_MAX element-count boundary, so the int32-safe case must stay applicable and the over-threshold case must be applicable on no device.

docs(changelog): cover bwd and wrw in the ASM-GTC INT_MAX gate entry

33a00f0

test(conv): drop test_conv_asm_gtc_fwd_nhwc_large_tensor.cpp, superse…

24f7ebb

…ded by the consolidated file

test(conv): drop test_conv_asm_gtc_bwd_nhwc_large_tensor.cpp, superse…

c270ba5

…ded by the consolidated file

test(conv): drop test_conv_asm_gtc_wrw_nhwc_large_tensor.cpp, superse…

77b9e33

…ded by the consolidated file

fix(conv): DBSync prune gfx90a68 (auto)

4f198de

fix(conv): DBSync prune gfx90a68 (auto)

8393599

fix(conv): DBSync prune gfx90a6e (auto)

e1ff811

yassinsolim force-pushed the users/ysoliman/gate-asm-gtc-fwd-large-tensor branch 2 times, most recently from f0152e6 to c376df7 Compare August 20, 2026 15:58

yassinsolim force-pushed the users/ysoliman/gate-asm-gtc-fwd-large-tensor branch from 3dd0181 to 03ef09d Compare August 21, 2026 15:21

yassinsolim mentioned this pull request Aug 21, 2026

DO NOT MERGE: CI-only rerun of reverted perf-db cleaning content (flake vs regression check) #11143

Closed

yassinsolim added 3 commits August 21, 2026 16:54

Merge branch 'develop' into users/ysoliman/gate-asm-gtc-fwd-large-tensor

eb5a76b

Merge remote-tracking branch 'refs/remotes/origin/develop-x' into mrg…

ce70fff

…9568

cderb requested changes Aug 31, 2026

View reviewed changes

Comment thread projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp Resolved

Comment thread projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp Resolved

averinevg reviewed Sep 1, 2026

View reviewed changes

Comment thread projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp Resolved

yassinsolim added 4 commits September 2, 2026 10:49

Merge branch 'develop' into gate-asm-gtc-fwd-large-tensor

bc3ffa3

# Conflicts: # projects/miopen/src/kernels/gfx942e4.HIP.fdb.txt.bz2 # projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2

Merge branch 'develop' into gate-asm-gtc-fwd-large-tensor

8d2dd81

# Conflicts: # projects/miopen/src/kernels/gfx942130.HIP.fdb.txt.bz2 # projects/miopen/src/kernels/gfx942e4.HIP.fdb.txt.bz2 # projects/miopen/src/kernels/gfx950100.HIP.fdb.txt.bz2

yassinsolim requested a review from cderb September 4, 2026 03:07

cderb approved these changes Sep 8, 2026

View reviewed changes

yassinsolim enabled auto-merge (squash) September 8, 2026 18:03

yassinsolim disabled auto-merge September 8, 2026 18:09

yassinsolim mentioned this pull request Sep 8, 2026

[MIOpen] ASM-GTC wrw NHWC: 292 of 900 kernels ignore the batch-split dimension, masked today by the int32 bound #11805

Open

yassinsolim merged commit 545bfd3 into develop Sep 8, 2026
129 checks passed

yassinsolim deleted the users/ysoliman/gate-asm-gtc-fwd-large-tensor branch September 8, 2026 19:13

yassinsolim mentioned this pull request Sep 8, 2026

perf(conv): relax the ASM-GTC NHWC large-tensor guard for fwd and bwd #11809

Merged

1 task done

Navigation Menu

Repository navigation

fix(conv): gate ASM-GTC fwd/bwd/wrw NHWC solvers off int32-overflow tensors - #9568

fix(conv): gate ASM-GTC fwd/bwd/wrw NHWC solvers off int32-overflow tensors#9568
yassinsolim merged 23 commits into
developfrom
users/ysoliman/gate-asm-gtc-fwd-large-tensor

yassinsolim commented Jul 17, 2026 •

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

cderb left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

cderb left a comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Sep 9, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

yassinsolim commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

cderb left a comment

Choose a reason for hiding this comment

An arithmetic bug in the split helper, upstream of this gate

Uh oh!

Uh oh!

Uh oh!

Uh oh!

cderb left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Sep 9, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

yassinsolim commented Jul 17, 2026 •

edited

Loading

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

edited

Loading

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

edited

Loading