Repository navigation
revert(miopen): partial revert of CK byte span gate for large instances - #12396
Merged
JonathanLichtnerAMD merged 2 commits intoOct 5, 2026
Merged
JonathanLichtnerAMD merged 2 commits into
JonathanLichtnerAMD merged 2 commits into
Conversation
JIRA ID: ALMIOPEN-2645 Reverts the forward/backward-data portion of PR #11585's RequiresLargeTensorCKInstance() gate back to backward-weights-only, element-space based. That gate was a workaround for byte-span overflow in MI455 (gfx950) benchmarking; CK PR #11754 fixes the underlying overflow in CK itself, so the MIOpen-side forward/backward-data gate is no longer needed and was excluding many efficient large-tensor instances from perf. DB changes from PR #11585 are intentionally kept (handled by a later DB tuning update). This revert is a draft pending CK PR #11754 (#11754) landing.
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
JonathanLichtnerAMD
marked this pull request as ready for review
October 1, 2026 16:03
jdcampbe
approved these changes
Oct 2, 2026
JonathanLichtnerAMD
deleted the
users/jlichtne/almiopen-2645-ck-byte-span-revert
branch
October 5, 2026 21:36
Contributor
TheRock Submodule Bump ActivityNewest first
|
This was referenced Oct 9, 2026
yassinsolim
added a commit
that referenced
this pull request
Oct 13, 2026
…rent gate no longer rejects (#13421) Restore 228 of the 229 large-tensor CK entries #12088 removed from the gfx1201 system database (165 backward-data, 63 forward find-db entries and 225 perf-db rows). Every restored line is byte-identical to its pre-#12088 original; no lines are removed or newly written. #12396 returned RequiresLargeTensorCKInstance to its backward-weight-only, element-count form once #11754 gave CK its own per-instance overflow checks, so these forward and backward-data entries resolve again. On gfx1201 the pre-#12088 database fails DBSync on 409 split-K entries only; the split-K half is re-tuned separately in #13415. One entry stays out because it was slower. Validated on an AMD Radeon AI PRO R9700 (gfx1201, 64 CU): StaticFDBSync and GPU-reference checks pass for the base, this database and its union with #13415, and a paired MIOPEN_FIND_MODE=2 A/B measured 1.43x aggregate kernel time across the 51 shapes these entries win (single run). JIRA ID: ROCM-29026
mattwill-amd
pushed a commit
that referenced
this pull request
Oct 14, 2026
JIRA ID: ALMIOPEN-2913 ## Motivation Only 5 of the 30 commits that landed against `projects/miopen` since the 3.6.2 version bump (`2d20a8f775e`, #12211) were reflected in `CHANGELOG.md`. The remaining 25 were reviewed and the 15 externally-visible ones documented, so the ROCm 10.2.0 section accurately reflects what shipped. 14 entries added (two of the fixes share one entry), plus a duplicate-heading fix. ### Added * [Conv] NHWC (channel-last) support in the GEMM convolution solvers, with new im2col/col2im kernels (#8693) * [Conv] Cross-architecture AI heuristics for immediate mode (`MIOPEN_FIND_MODE=2`), extending data-driven selection to Navi (#9680) * [Conv] Missing 3x3 stride-1 pad-1 tile-28 CK depthwise backward-data instance — ALMIOPEN-2423 (#11271) ### Changed * [Conv] `ConvHipConv` held out of the AI perf-config picker until its model is retrained — ALMIOPEN-2494 (#13337) * [Conv] CK byte-span gate lifted for forward and backward-data, now that CK fixes the 32-bit offset overflow itself — ALMIOPEN-2645 (#12396) ### Optimized * [Conv] `ConvDirectNaiveConv`: per-direction kernel files, float accumulators, spatial tiling, bfloat16 path — ALMIOPEN-2485/2486 (#7529) * Kernel cache reuses cached SQLite connections instead of reopening per load/store — ALMIOPEN-2598 (#11933) ### Resolved Issues * [Conv] Incorrect backward-weights results from the gfx950 direct kernel at `waves_q == 1` (missing cross-lane ordering on LDS staging reuse) — ROCM-31508 (#12406) * [Conv] Up to 8x slower hipconv `direct`/`direct_l1`/`direct_wgrad` under ROCm 10.1's full AMDGCN LTO default; fixed with `-fno-offload-lto` — ROCM-31334 (#12328) * [Conv] gfx942 multi-node training throughput regression from backward GEMM on small-batch 1x1 NHWC shapes — ROCM-31155 (#13232) * [Conv] Invalid CK perf-config entries in the gfx1201 and gfx1100 system databases — ROCM-29026 (#12088, #12117) * [BatchNorm] JIT compilation failures on non-GFX9 targets; Welford DPP assembly restricted to GFX9 — ALMIOPEN-2838 (#12889) * Out-of-bounds write in `miopenLoadSolution` when a serialized solution carries more than three work dimensions — ROCM-31082 (#12242) * `uint64_t` typedef conflict between `miopen_cstdint.hpp` and system `<stdint.h>` on LP64 Linux (#12623) ## Technical Details Single-file change to `projects/miopen/CHANGELOG.md`. The 10.2.0 section had **two** `### Resolved Issues` headings. The two entries under the trailing one (`ConvHipConv` gfx1250 crash/LDS race, and the perf-db reload config loss) were merged into the first, and the trailing heading removed — this accounts for the 2 removed bullets in the diff. Section order is now `Resolved Issues → Added → Changed → Optimized → Removed`, matching the 3.6.1 section. `Optimized` is new to this release section. Already documented, so untouched: #12874 (reduce int32 rejection), #12000 / #13016 / #13213 (hipconv v0.3.1 / v0.3.2 / v0.4.0), #11761 (gfx803 removal). ## Notes Deliberately excluded as not externally visible (10 commits): * **hipDNN convolution forwarding (#12276)** — `MIOPEN_ENABLE_HIPDNN_WRAPPER` defaults to `OFF` (`projects/miopen/CMakeLists.txt:162`), so no shipped build compiles the forwarding code and the runtime `MIOPEN_HIPDNN_FORWARDING` flag has nothing to enable. * **Test-only**: #13289 (gfx1250 depthwise test applicability, ALMIOPEN-2841), #11610 (hipDNN parity harness), #12660 (Jenkins status post). * **Internal kernel-math cleanup, generated ISA unchanged**: #13312 (drop direct `__ocml` calls), #13190 (remove broken single-argument `__half` overloads). * **Add/revert pairs, net no change**: #12466 + #13066 (gfx1101 Navi system DBs), #12496 + #12800 (gfx1250 in the dev image arch lists). ## Risk Assessment Documentation-only change to a single Markdown file; no code, build, or runtime impact. **Risk Level: 🟢 Low** ## Test Plan None required — Markdown-only. Verified that the entry topics are not already documented elsewhere in the file, that no content was lost when merging the duplicate heading, and that every new entry maps to a commit in `2d20a8f775e..origin/develop -- projects/miopen`. ## Test Result 16 bullets added, 2 relocated — 14 net-new entries covering 15 commits. Duplicate heading removed; section ordering matches 3.6.1. Co-authored-by: Claude <noreply@anthropic.com>
yassinsolim
added a commit
that referenced
this pull request
Oct 17, 2026
…#13373 (#13697) JIRA ID: ROCM-29026, ALMIOPEN-2090 Closes #13394. Two fixes to the gfx1200 system database (`gfx120020.HIP.fdb.txt.bz2`, `gfx120020.db.txt.bz2`), one commit each, plus a changelog commit: 1. Re-tune 8 and remove 3 of the 11 CK backward-weights rows that TheRock-built binaries reject (the part of #13394 that #13420 left). 2. Drop the 835 CK backward-weights entries that #13373 (merged Oct 14) invalidates. Before #13373, the configs behind them returned wrong weight gradients on gfx1200. They ship together because the database from part 1 alone fails `StaticFDBSync/10` with 835 invalid configs on any build that has #13373, which develop now does. ## Part 1: the 11 rows TheRock builds reject All 11 are `ConvHipImplicitGemmGroupWrwXdlops` entries whose stored instance has `split_k = -1` ("auto"). At runtime `MakeArgPtr` calls `GetMaxOccupancy()` (`hipOccupancyMaxActiveBlocksPerMultiprocessor`) on the live device, multiplies by the CU count, divides by the kernel's output grid size and picks `k_batch` from that (`get_best_occupancy_k_batch_value`, `split_k_utils.hpp:37`). `IsSupportedArgument` then accepts or refuses the derived `k_batch` depending on the instance's tile shape. Evidence for occupancy as the mechanism, and its limits: - #13420's analysis (prior art, not re-measured here) found TheRock CI artifacts reject 31 entries where a from-source build with the ROCm 7.14.1 toolchain rejects 20. That is a build-to-build pass/fail difference, not a measured occupancy difference. - An `LD_PRELOAD` interposer that forces a fixed value for every occupancy query, run against TheRock's own binaries, shows the failures follow the forced value: 0 of the 236 auto-split-K WRW rows fail at occupancy 1, and 1, 3, 3, 5, 6, 7, 9, 17, 28, 39 and 38 fail at 2, 3, 4, 5, 6, 8, 10, 12, 16, 24 and 32 (not strictly monotonic). Unforced, TheRock's build fails 11. - A run meant to compare the occupancy each build actually measures failed to start its container (`Unable to find group render`) and was not retried. "The two builds measure different occupancy" is the best-supported explanation, not a direct measurement. None of the 11 trips the large-tensor gate that #13419 is about. Each key was re-tuned on the TheRock build (`MIOPEN_FIND_MODE=1 MIOPEN_FIND_ENFORCE=4`, those 11 rows stripped from the system DB first) and each winner was re-run in a fresh process with `MIOpenDriver -V 1`: | key | re-tuned instance (+split_k) | verify | action | |---|---|---|---| | `1024-14-14-1x1-256-14-14-1024-...-BF16-W` | `Explicit_Xdl` | OK | kept | | `128-56-56-1x1-192-56-56-42-...-FP16-W` | `Explicit_Xdl` | OK | kept | | `256-28-28-1x1-512-28-28-1024-...-BF16-W` | `Explicit_Xdl` | OK | kept | | `352-7-7-3x3-352-14-14-10-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` | OK (forced) | kept | | `352-7-7-3x3-352-14-14-11-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` | OK (forced) | kept | | `352-7-7-3x3-352-14-14-22-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` | OK | kept | | `352-7-7-3x3-352-7-7-16-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` | OK | kept | | `512-1-1-1x1-512-1-1-1-...-BF16-W` | `Explicit_Xdl` | OK (forced) | kept | | `60-13-13-3x3-60-28-28-4-...-BF16-W` | `Xdl_CShuffleV3_WmmaPorted<64,64,32,32,...>+8` | **FAILED 0.88 > 0.13** | removed | | `64-28-28-7x1-64-28-28-1-...-BF16-W` | `Xdl_CShuffleV3_WmmaPorted<64,32,64,32,...>+-1` | **FAILED 0.84 > 0.13** | removed | | `64-28-28-7x1-64-28-28-1-...-FP16-W` | `Xdl_CShuffleV3_WmmaPorted<64,32,32,32,...>+8` | **FAILED 0.84 > 0.08** | removed | "Forced" rows: immediate-mode find never selects that key's CK entry on its own (a cheaper fallback wins), so the exact CK solver was forced with `MIOPEN_DEBUG_FIND_ONLY_SOLVER` (Slurm 68090048); all three verify (1.7e-5, 1.5e-5, 0). The three removed rows are the WMMA-ported V3 family that part 2 is about; #13373 rejects them too. Default find mode 2, TheRock build, median of 3 runs, RX 9060 XT, develop's database vs part 1: | key | develop (ms) | part 1 (ms) | ratio | |---|---|---|---| | 1024-...-BF16-W | 6.495 | 5.446 | 0.84x | | 128-...-FP16-W | 2.095 | 0.918 | 0.44x | | 256-...-BF16-W | 17.484 | 12.898 | 0.74x | | 352-...-10-...W_g352 | 0.087 | 0.099 | 1.13x | | 352-...-11-...W_g352 | 0.090 | 0.075 | 0.84x | | 352-...-22-...W_g352 | 0.147 | 0.115 | 0.78x | | 352-...-16-...W_g352 | 0.098 | 0.084 | 0.86x | | 512-...-BF16-W | 0.036 | 0.036 | 1.0x | | 60-...-BF16-W (removed) | 0.081 | 0.166 | 2.05x | | 64-...-BF16-W (removed) | 0.052 | 0.072 | 1.40x | | 64-...-FP16-W (removed) | 0.089 | 0.038 | 0.43x | Removing all 11 rows instead of re-tuning was tried and rejected: two keys then fall back to `GemmWrw1x1_stride1`, which fails verification on both (0.59 and 0.60 against 0.13, BF16 at N=1024). ## Part 2: the entries #13373 invalidates #13373 fixes `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3::IsSupportedArgument`, which compared `num_k_loop` against the wave64 `PrefetchStages` on RDNA. Wave32 pipeline-V2 configs with too few K loops were accepted, the kernel returned without writing its output, and it finished far faster than the real work, so tuning ranked it first. Those configs are in this database. Scope: 1845 find-db records list `ConvHipImplicitGemmGroupWrwXdlops`; 844 resolve to a `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted` perf config, and in 686 of those the CK entry is the fastest in the record, so find mode 2 picks it. Before #13373: all 686 were run on an RX 9060 XT with the TheRock main CI build of 24bdd07 (run 38489560586), one fresh MIOpenDriver process each, `-F 4 -t 1 -i 10 -V 1`. Forcing the solver and letting find mode 2 choose give the same verdict on every key: 680 wrong, 6 correct. Of the 680 failures, 674 have finite relative errors from 0.56 to 2.5e3 (median 0.91), 4 are NaN and 2 are above 1e30. With #13373: `StaticFDBSync/10` on the part 1 database, built from 24bdd07 plus only the #13373 diff, reports 835 invalid perf configs. That set contains all 680 measured-wrong keys and none of the 6 measured-correct ones; the other 155 sit in records where CK was not the fastest, so they were never measured. The committed files are the `MIOPEN_DBSYNC_CLEAN=1` output of that run: 835 solver entries leave the find-db, 826 perf-db records change, and the 435 of them that held no other solver are dropped. No find-db record is left empty. gfx120140 (gfx1201) has no CK backward-weights records and is untouched. ## Validation of this PR's database (parts 1 and 2 together) `Smoke/CPU_DBSync_NONE.StaticFDBSync/10` on `navi4x-hw-10` (Slurm 68091724): | check | build | database | result | |---|---|---|---| | negative control | 24bdd07 + #13373 | part 1 only | FAILED, 835 invalid perf configs | | cleanup | 24bdd07 + #13373, `MIOPEN_DBSYNC_CLEAN=1` | part 1 only | OK, 835 removals | | this PR | 24bdd07 + #13373 | this PR | OK, 0 failures | | this PR | 24bdd07, no #13373 | this PR | OK, 0 failures | | this PR | TheRock CI build 38489560586 | this PR | OK, 0 failures | The last row also covers part 1: develop's database fails 11 on that build. The committed bz2 files decompress to exactly the cleanup output (find-db sha256 `294a5eb5...`, perf-db `b159206a...`). Correctness after the change, #13373 build, default find mode 2, one fresh process per problem, `-V 1`, all 686 problems: 686 verify OK (Slurm 68091782). 680 move off CK (`GemmWrwUniversal` 486, `GemmWrw1x1_stride1` 63, `ConvWinoRageRxS` 41, Winograd variants 89, `TransposedConvBinWinoRxS` 1); the other 6 keep CK. All 63 `GemmWrw1x1_stride1` picks verify. Leaving the stale entries in place is not free once #13373 lands. The same run against the database without part 2 gave no wrong results, but 90 of the 529 problems it reached before its 6-hour limit hit the driver's 1200 s timeout. Against this PR's database all 686 problems finish, the slowest in 493 ms. ## What could break The 680 problems that move off CK are slower than the times recorded for the old entries: geomean 4.74x, median 3.69x, worst 4321x (0.11 ms to 493 ms). That is an upper bound, because the old kernel returned without computing the output, so its recorded time is the cost of a no-op. Re-tuning recovers most of it: on 12 stratified keys re-tuned with #13373, a valid CK config wins every time, all 12 verify, and they run a geomean 6.5x faster than the fallback (1.1x to 64.4x), at 4 to 16 s of search per key. Re-tuning all 835 entries is a follow-up; at that rate it is roughly 1 to 2 hours on one gfx1200 (estimate). ## Not in this PR - The re-tune of the 835 entries. - #13419 (element count vs byte span in `RequiresLargeTensorCKInstance`). - gfx1151 (`/9`) and gfx1102 (`/8`): their reported failures came from the pre-#12396 large-tensor gate and pass on current builds (GPU-free check); no database change. Related: #13305, #12942, #13373. ## Test plan ``` MIOPEN_SYSTEM_DB_PATH=<staged-dir> MIOPEN_USER_DB_PATH=<scratch> \ ./bin/miopen_gtest --gtest_filter='Smoke/CPU_DBSync_NONE.StaticFDBSync/10' ``` The branch was rebased onto develop after these runs; the database files are byte-identical to the validated ones, and develop has not changed them.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
revert(miopen): partial revert of CK byte span gate for large instances
JIRA ID: ALMIOPEN-2645
Motivation
PR #11585 widened
RequiresLargeTensorCKInstance()to gate forward and backward-data (not just backward-weights) on tensor byte span, to work around a 32-bit byte-offset overflow in CK. CK PR #11754 (#11754) fixes that overflow directly in CK, so the MIOpen-side forward/backward-data gate is no longer needed and was unnecessarily excluding large-tensor problems from faster int32 CK instances.Technical Details
Reverts
RequiresLargeTensorCKInstance()inprojects/miopen/src/ck_impl/implicitgemm_ck_util.hppto its pre-#11585 state: the large-tensor check is scoped back to backward-weights only, and uses element-space (GetElementSpace()) rather than byte-span (GetNumBytes()). Test expectations inprojects/miopen/test/gtest/unit_implicitgemm_ck_util.cppare reverted to match (forward/backward-data no longer expected to require large-tensor instances at the >INT_MAX-byte-span shape), and the byte-span-specificByteSpanOverflowGatetest and itsMakeByteSpanConvProblemhelper are removed. The.bz2FDB/PDB database files that PR #11585 also touched are intentionally left untouched — this PR is source-only, with DB changes deferred to a later database-tuning update.Risk Assessment
Source-only change to a solver-selection gate in a single header, reverting an already-reviewed prior state; blast radius is confined to CK instance selection for forward/backward-data conv problems with large tensors, and is gated on a companion CK fix (#11754) landing first — this PR is a draft and not intended to merge before that.
Risk Level: 🟡 Medium
Impacted Components
projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp(RequiresLargeTensorCKInstance)projects/miopen/test/gtest/unit_implicitgemm_ck_util.cppPotential Side Effects
Until CK PR #11754 lands, large-tensor forward/backward-data problems that trip the byte-span overflow (byte span > INT_MAX with element count still < INT_MAX) could silently pick a narrowing int32 CK instance again, reintroducing the correctness bug #11585 fixed for those two directions.
Test Plan
Built and ran
test_unit_implicitgemm_ck_utilin a docker dev container (gfx942).Test Result
All 30 tests pass, including the two reverted large-tensor gate tests (
WrwElementCountOverflowGate,Wrw3DElementCountOverflowGate) back to their pre-#11585 (element-space, backward-weights-only) expectations. A broader multi-shape perf-eval sweep has been run as part of the MI455 performance benchmarking.