Repository navigation
fix(MIOpen): gate CK large-tensor instances on byte span, all directions - #11585
Merged
JonathanLichtnerAMD merged 4 commits intoSep 3, 2026
Conversation
RequiresLargeTensorCKInstance() compared element count (not byte span) against INT_MAX, and only applied the check to backward-weights. A grouped conv tensor whose byte span exceeds 2 GiB but whose element count and per-dim strides fit in int32 (e.g. forward/backward-data) could still be routed to an int32 CK instance and silently corrupt results (ALMIOPEN-2570). Use GetNumBytes() instead of GetElementSpace(), and apply the check to all directions; CK already ships int64 Large_Tensor instances for fwd/bwd-data/wrw, so no CK-side changes are needed. Co-Authored-By: Claude <noreply@anthropic.com>
✅ 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. |
…IOPEN-2570-tighten-miopen-large-tensor-def
…nsor byte-span gate ALMIOPEN-2570 changed RequiresLargeTensorCKInstance to gate on byte span (GetNumBytes() > INT_MAX) across all directions, correctly excluding non-large CK instances for grouped-conv problems whose byte span exceeds 2 GiB. The shipped SystemDBs were tuned under the old (backward-weights-only, element-count) behavior and pointed those >2 GiB forward/backward-data/wrw problems at non-large instances, so StaticFDBSync now flags them as invalid perf configs. Regenerate the three affected SystemDBs via the MIOPEN_DBSYNC_CLEAN process (rocjitsu CU-corrected for the GPU-free arches, per the #11283 runner), iterated to a fixpoint (a second pass removes fdb entries whose shared pdb key was collaterally dropped): - gfx942130 (304 CU): removed 2349 invalid grouped-conv perf-config entries - gfx942e4 (228 CU): removed 2098 - gfx950100 (256 CU): removed 2181 Re-verified: a fresh StaticFDBSync run (CLEAN off) passes on all three with 0 invalid perf configs, 0 dangling PDB entries, 0 inapplicable solvers. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…en-large-tensor-def
JonathanLichtnerAMD
enabled auto-merge (squash)
September 3, 2026 03:50
JonathanLichtnerAMD
deleted the
users/jlichtne/ALMIOPEN-2570-tighten-miopen-large-tensor-def
branch
September 3, 2026 05:12
Contributor
TheRock Submodule Bump ActivityNewest first
|
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
Vsevolod1983
pushed a commit
that referenced
this pull request
Sep 28, 2026
… 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.
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Oct 5, 2026
…es (#12396) # 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()` in `projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp` to 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 in `projects/miopen/test/gtest/unit_implicitgemm_ck_util.cpp` are 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-specific `ByteSpanOverflowGate` test and its `MakeByteSpanConvProblem` helper are removed. The `.bz2` FDB/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.cpp` ### Potential 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_util` in 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.
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.
fix(MIOpen): gate CK large-tensor instances on byte span, all directions
JIRA ID: ALMIOPEN-2570
Motivation
RequiresLargeTensorCKInstanceonly forced the int64 CK path for backward-weights problems whose element count exceeded INT_MAX. Non-large-tensor CK instances narrow MakeArgumentPointer's int64 overload back to int32 and compute byte offsets in 32-bit arithmetic, so a tensor can have an in-range element count and per-dim strides yet still overflow if its byte span exceeds INT_MAX — and this applies to forward and backward-data too, not just backward-weights. The old check missed both of these cases, silently returning wrong results for real (non-synthetic) large-tensor shapes.Tightening the gate also invalidated SystemDB entries tuned under the old behavior: the shipped gfx942/gfx950 find/perf DBs point >2 GiB grouped forward/backward-data/wrw problems at now-excluded non-large instances, which
StaticFDBSyncflags as invalid perf configs. Those stale entries have to be cleaned alongside the code change.Technical Details
implicitgemm_ck_util.hpp: replaced the backward-weights-onlyGetElementSpace() > INT_MAXcheck with a direction-agnosticGetNumBytes() > INT_MAXcheck on In/Out/Weights, matching the actual quantity (byte span) that overflows the narrowed int32 arithmetic. Updatedunit_implicitgemm_ck_util.cppaccordingly: existing element-count tests now expectTRUEfor forward/backward-data too, and a newByteSpanOverflowGatetest isolates the byte-span check from the element-count check using an fp32 shape that crosses the 2 GiB byte boundary while staying under INT_MAX elements.SystemDB clean: regenerated the three affected SystemDBs via the
MIOPEN_DBSYNC_CLEANprocess (rocjitsu CU-corrected for the GPU-free arches, following #11283's runner), iterated to a fixpoint — gfx942130/304: 2349 entries, gfx942e4/228: 2098, gfx950100/256: 2181. The removed entries are exactly the >2 GiB grouped-conv perf configs the new gate excludes; a second pass drops fdb entries whose shared pdb key was collaterally removed. Also merged currentdevelop(picks up #11283's separate gfx950100/gfx942e4 toolchain-staleness fixes).Risk Assessment
Low blast radius: an internal gating helper and its unit test, no public API change. The gate is strictly broader/more conservative than before (adds forward/backward-data + closes a byte-span gap), so any behavior change is diverting previously-miscomputed shapes onto the (slower but correct) large-tensor CK path. This PR additionally regenerates three shipped SystemDBs; the removed perf configs only affect >2 GiB shapes, which fall back to runtime find on the correct large-tensor path (perf-only, no correctness impact).
Risk Level: 🟢 Low
Test Plan
Updated and extended
unit_implicitgemm_ck_util.cpp: existing element-count regression tests now assert forward/backward-data also require the large-tensor instance, and a newByteSpanOverflowGatetest checks all three directions on both sides of the byte-span boundary with an otherwise-INT_MAX-safe element count/stride shape.Re-verified the cleaned DBs with a fresh
StaticFDBSyncrun (CLEAN off) for gfx942130/304, gfx942e4/228, and gfx950100/256: 0 invalid perf configs, 0 dangling PDB entries, 0 inapplicable solvers.Test Result
The three cleaned SystemDBs pass
StaticFDBSync; the gate unit tests were updated to match the new logic. CI (the gfx942Test miopenshard + TheRockmiopen-dbsynccomponent) re-running on the DB commit.