Repository navigation
fix(ck): PrefetchStages on host for wave32 instances - #13373
Merged
Merged
Conversation
github-actions
Bot
added
project: composablekernel
ck: convolution
✅ 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. |
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
The pipeline-v2 wave32 regression lacks targeted automated coverage.
1 open finding
What changed in this PR
Corrects wave32 grouped-convolution backward-weight capability checks.
Changes:
- Uses the wave32 GEMM pipeline’s prefetch-stage count on RDNA devices.
| File | Description |
|---|---|
device_grouped_conv_bwd_weight_xdl_cshuffle_v3.hpp |
Corrects wave32 pipeline validation. |
🧠 Review effort: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
1 task done
BrianHarrisonAMD
approved these changes
Oct 10, 2026
BrianHarrisonAMD
left a comment
BrianHarrisonAMD
left a comment
Contributor
There was a problem hiding this comment.
Verified this fixes the issues logged.
JP-Fernando
approved these changes
Oct 13, 2026
JP-Fernando
left a comment
JP-Fernando
left a comment
Contributor
There was a problem hiding this comment.
This PR repairs the PrefetchStages selection for wave32 instances. Tests have been added to check for correctness and regressions. Approved.
bartekxk
approved these changes
Oct 14, 2026
bartekxk
enabled auto-merge (squash)
October 14, 2026 21:24
bartekxk
disabled auto-merge
October 14, 2026 21:25
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Oct 14, 2026
fix(ck): PrefetchStages on host for wave32 instances ## Motivation ISSUE: #13305 ## Technical Details `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3` uses `GridwiseGemm64::BlockwiseGemmPipe::PrefetchStages` on RDNA (wave32) in `IsSupportedArgument`. This is incorrect for pipeline V2 where the PrefetchStages depend on the wave size. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
This was referenced Oct 15, 2026
Merged
1 task
Contributor
TheRock Submodule Bump ActivityNewest first
|
yassinsolim
added a commit
that referenced
this pull request
Oct 17, 2026
…#13373 (#13697) JIRA ID: ROCM-29026, ALMIOPEN-2090 Closes #13394. Two fixes to the gfx1200 system database (`gfx120020.HIP.fdb.txt.bz2`, `gfx120020.db.txt.bz2`), one commit each, plus a changelog commit: 1. Re-tune 8 and remove 3 of the 11 CK backward-weights rows that TheRock-built binaries reject (the part of #13394 that #13420 left). 2. Drop the 835 CK backward-weights entries that #13373 (merged Oct 14) invalidates. Before #13373, the configs behind them returned wrong weight gradients on gfx1200. They ship together because the database from part 1 alone fails `StaticFDBSync/10` with 835 invalid configs on any build that has #13373, which develop now does. ## Part 1: the 11 rows TheRock builds reject All 11 are `ConvHipImplicitGemmGroupWrwXdlops` entries whose stored instance has `split_k = -1` ("auto"). At runtime `MakeArgPtr` calls `GetMaxOccupancy()` (`hipOccupancyMaxActiveBlocksPerMultiprocessor`) on the live device, multiplies by the CU count, divides by the kernel's output grid size and picks `k_batch` from that (`get_best_occupancy_k_batch_value`, `split_k_utils.hpp:37`). `IsSupportedArgument` then accepts or refuses the derived `k_batch` depending on the instance's tile shape. Evidence for occupancy as the mechanism, and its limits: - #13420's analysis (prior art, not re-measured here) found TheRock CI artifacts reject 31 entries where a from-source build with the ROCm 7.14.1 toolchain rejects 20. That is a build-to-build pass/fail difference, not a measured occupancy difference. - An `LD_PRELOAD` interposer that forces a fixed value for every occupancy query, run against TheRock's own binaries, shows the failures follow the forced value: 0 of the 236 auto-split-K WRW rows fail at occupancy 1, and 1, 3, 3, 5, 6, 7, 9, 17, 28, 39 and 38 fail at 2, 3, 4, 5, 6, 8, 10, 12, 16, 24 and 32 (not strictly monotonic). Unforced, TheRock's build fails 11. - A run meant to compare the occupancy each build actually measures failed to start its container (`Unable to find group render`) and was not retried. "The two builds measure different occupancy" is the best-supported explanation, not a direct measurement. None of the 11 trips the large-tensor gate that #13419 is about. Each key was re-tuned on the TheRock build (`MIOPEN_FIND_MODE=1 MIOPEN_FIND_ENFORCE=4`, those 11 rows stripped from the system DB first) and each winner was re-run in a fresh process with `MIOpenDriver -V 1`: | key | re-tuned instance (+split_k) | verify | action | |---|---|---|---| | `1024-14-14-1x1-256-14-14-1024-...-BF16-W` | `Explicit_Xdl` | OK | kept | | `128-56-56-1x1-192-56-56-42-...-FP16-W` | `Explicit_Xdl` | OK | kept | | `256-28-28-1x1-512-28-28-1024-...-BF16-W` | `Explicit_Xdl` | OK | kept | | `352-7-7-3x3-352-14-14-10-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` | OK (forced) | kept | | `352-7-7-3x3-352-14-14-11-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` | OK (forced) | kept | | `352-7-7-3x3-352-14-14-22-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` | OK | kept | | `352-7-7-3x3-352-7-7-16-...-FP16-W_g352` | `TwoStage_Xdl_CShuffle+1` | OK | kept | | `512-1-1-1x1-512-1-1-1-...-BF16-W` | `Explicit_Xdl` | OK (forced) | kept | | `60-13-13-3x3-60-28-28-4-...-BF16-W` | `Xdl_CShuffleV3_WmmaPorted<64,64,32,32,...>+8` | **FAILED 0.88 > 0.13** | removed | | `64-28-28-7x1-64-28-28-1-...-BF16-W` | `Xdl_CShuffleV3_WmmaPorted<64,32,64,32,...>+-1` | **FAILED 0.84 > 0.13** | removed | | `64-28-28-7x1-64-28-28-1-...-FP16-W` | `Xdl_CShuffleV3_WmmaPorted<64,32,32,32,...>+8` | **FAILED 0.84 > 0.08** | removed | "Forced" rows: immediate-mode find never selects that key's CK entry on its own (a cheaper fallback wins), so the exact CK solver was forced with `MIOPEN_DEBUG_FIND_ONLY_SOLVER` (Slurm 68090048); all three verify (1.7e-5, 1.5e-5, 0). The three removed rows are the WMMA-ported V3 family that part 2 is about; #13373 rejects them too. Default find mode 2, TheRock build, median of 3 runs, RX 9060 XT, develop's database vs part 1: | key | develop (ms) | part 1 (ms) | ratio | |---|---|---|---| | 1024-...-BF16-W | 6.495 | 5.446 | 0.84x | | 128-...-FP16-W | 2.095 | 0.918 | 0.44x | | 256-...-BF16-W | 17.484 | 12.898 | 0.74x | | 352-...-10-...W_g352 | 0.087 | 0.099 | 1.13x | | 352-...-11-...W_g352 | 0.090 | 0.075 | 0.84x | | 352-...-22-...W_g352 | 0.147 | 0.115 | 0.78x | | 352-...-16-...W_g352 | 0.098 | 0.084 | 0.86x | | 512-...-BF16-W | 0.036 | 0.036 | 1.0x | | 60-...-BF16-W (removed) | 0.081 | 0.166 | 2.05x | | 64-...-BF16-W (removed) | 0.052 | 0.072 | 1.40x | | 64-...-FP16-W (removed) | 0.089 | 0.038 | 0.43x | Removing all 11 rows instead of re-tuning was tried and rejected: two keys then fall back to `GemmWrw1x1_stride1`, which fails verification on both (0.59 and 0.60 against 0.13, BF16 at N=1024). ## Part 2: the entries #13373 invalidates #13373 fixes `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3::IsSupportedArgument`, which compared `num_k_loop` against the wave64 `PrefetchStages` on RDNA. Wave32 pipeline-V2 configs with too few K loops were accepted, the kernel returned without writing its output, and it finished far faster than the real work, so tuning ranked it first. Those configs are in this database. Scope: 1845 find-db records list `ConvHipImplicitGemmGroupWrwXdlops`; 844 resolve to a `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted` perf config, and in 686 of those the CK entry is the fastest in the record, so find mode 2 picks it. Before #13373: all 686 were run on an RX 9060 XT with the TheRock main CI build of 24bdd07 (run 38489560586), one fresh MIOpenDriver process each, `-F 4 -t 1 -i 10 -V 1`. Forcing the solver and letting find mode 2 choose give the same verdict on every key: 680 wrong, 6 correct. Of the 680 failures, 674 have finite relative errors from 0.56 to 2.5e3 (median 0.91), 4 are NaN and 2 are above 1e30. With #13373: `StaticFDBSync/10` on the part 1 database, built from 24bdd07 plus only the #13373 diff, reports 835 invalid perf configs. That set contains all 680 measured-wrong keys and none of the 6 measured-correct ones; the other 155 sit in records where CK was not the fastest, so they were never measured. The committed files are the `MIOPEN_DBSYNC_CLEAN=1` output of that run: 835 solver entries leave the find-db, 826 perf-db records change, and the 435 of them that held no other solver are dropped. No find-db record is left empty. gfx120140 (gfx1201) has no CK backward-weights records and is untouched. ## Validation of this PR's database (parts 1 and 2 together) `Smoke/CPU_DBSync_NONE.StaticFDBSync/10` on `navi4x-hw-10` (Slurm 68091724): | check | build | database | result | |---|---|---|---| | negative control | 24bdd07 + #13373 | part 1 only | FAILED, 835 invalid perf configs | | cleanup | 24bdd07 + #13373, `MIOPEN_DBSYNC_CLEAN=1` | part 1 only | OK, 835 removals | | this PR | 24bdd07 + #13373 | this PR | OK, 0 failures | | this PR | 24bdd07, no #13373 | this PR | OK, 0 failures | | this PR | TheRock CI build 38489560586 | this PR | OK, 0 failures | The last row also covers part 1: develop's database fails 11 on that build. The committed bz2 files decompress to exactly the cleanup output (find-db sha256 `294a5eb5...`, perf-db `b159206a...`). Correctness after the change, #13373 build, default find mode 2, one fresh process per problem, `-V 1`, all 686 problems: 686 verify OK (Slurm 68091782). 680 move off CK (`GemmWrwUniversal` 486, `GemmWrw1x1_stride1` 63, `ConvWinoRageRxS` 41, Winograd variants 89, `TransposedConvBinWinoRxS` 1); the other 6 keep CK. All 63 `GemmWrw1x1_stride1` picks verify. Leaving the stale entries in place is not free once #13373 lands. The same run against the database without part 2 gave no wrong results, but 90 of the 529 problems it reached before its 6-hour limit hit the driver's 1200 s timeout. Against this PR's database all 686 problems finish, the slowest in 493 ms. ## What could break The 680 problems that move off CK are slower than the times recorded for the old entries: geomean 4.74x, median 3.69x, worst 4321x (0.11 ms to 493 ms). That is an upper bound, because the old kernel returned without computing the output, so its recorded time is the cost of a no-op. Re-tuning recovers most of it: on 12 stratified keys re-tuned with #13373, a valid CK config wins every time, all 12 verify, and they run a geomean 6.5x faster than the fallback (1.1x to 64.4x), at 4 to 16 s of search per key. Re-tuning all 835 entries is a follow-up; at that rate it is roughly 1 to 2 hours on one gfx1200 (estimate). ## Not in this PR - The re-tune of the 835 entries. - #13419 (element count vs byte span in `RequiresLargeTensorCKInstance`). - gfx1151 (`/9`) and gfx1102 (`/8`): their reported failures came from the pre-#12396 large-tensor gate and pass on current builds (GPU-free check); no database change. Related: #13305, #12942, #13373. ## Test plan ``` MIOPEN_SYSTEM_DB_PATH=<staged-dir> MIOPEN_USER_DB_PATH=<scratch> \ ./bin/miopen_gtest --gtest_filter='Smoke/CPU_DBSync_NONE.StaticFDBSync/10' ``` The branch was rebased onto develop after these runs; the database files are byte-identical to the validated ones, and develop has not changed them.
pschang-phy
added a commit
that referenced
this pull request
Oct 19, 2026
…ix the regressions (#12766) JIRA ID : AIHPBLAS-1405 ## Motivation `WaitAwareScheduleRepairPass` is a count-based repair that runs after wait-count insertion and reorders instructions the DAG scheduler has already placed, without modeling WMMA latency, co-issue slots or wait stalls. So, this PR disables it and leaves instruction placement to the scheduler. Disabling it exposes places where the scheduler relied on the pass. The following commits fix those regressions in the scheduler, one commit per fix. ## Technical Details ### 1. Keep the single-stage prefetch lead on short stages **Problem.** For loops without HalfPLR, KernelWriter asks for single-stage prefetch grouping (`PrefetchLeadWmmas = 4`): hold each iteration's prefetches and issue them as a group in front of the stage barrier. StinkyTofu dropped every prefetch lead on stages shorter than 64 WMMAs (`PrefetchLeadMinStageWmmas`), single-stage leads included. MAB loops have 16 WMMAs per stage, so their three `global_prefetch_b8` issued as soon as their addresses were ready, in the middle of the stage among the LDS reads. The WMMAs that need that LDS data then wait longer: in `bbs_tn_mab`, the WMMA that starts window 10 waits 125 cycles per iteration instead of 11 (hardware trace, typical iteration). The repair pass had hidden this by moving the prefetches to the end of the stage. **Change.** The short-stage rule now drops only multi-stage leads (8 or more), since only they need room within the stage; single-stage leads are kept at any stage length: ```cpp // StageWmmaCounter::effectiveLead (ReadyQueue.hpp) return lead >= kSingleStageLeadBelow && shortStage ? 0 : lead; // was: lead > 0 && shortStage ? 0 : lead ``` `kSingleStageLeadBelow` (8) moves into `StageWmmaCounter`, and `CDNA5ReadyQueue` refers to it, so the rule and the scheduler agree on what a single-stage lead is. Comments and the scheduler docs are updated. In `bbs_tn_mab` the three prefetches now issue together in window 10, and the WMMA above waits 17 cycles. ## Test Plan ## Test Result ### Perf YAMLs - (develop) 0eedc5e fix(ck): PrefetchStages on host for wave32 instances (#13373) ``` Best percent versus 0eedc5e test 0eedc5e c6334c1 source neither stinkytofu mab-bf16 19145.30 19062.70 (-0.43%) mab-FP8 18843.30 18830.40 (-0.07%) mab-MXFP8 18402.90 18428.00 (+0.14%) mab-MXFP4 18617.90 18640.50 (+0.12%) mab-NVFP4 18595.50 18598.70 (+0.02%) mab-bf16-m64 16972.80 17061.40 (+0.52%) maf-bf16 3377.25 3380.12 (+0.08%) maf-FP8 10030.20 10002.20 (-0.28%) maf-FP8-henry 10028.80 10040.20 (+0.11%) maf-MXFP8 8968.30 8957.38 (-0.12%) maf-MXFP4 17412.00 17389.20 (-0.13%) maf-NVFP4 16498.90 16530.90 (+0.19%) medium-bf16 2599.83 2605.49 (+0.22%) medium-FP8 6159.78 6158.45 (-0.02%) medium-MXFP8 5015.56 5022.01 (+0.13%) medium-MXFP4 8234.32 8204.69 (-0.36%) Average with SD test 0eedc5e c6334c1 source neither stinkytofu mab-bf16 19091.28 +/- 32.35 19049.17 +/- 9.75 mab-FP8 18800.28 +/- 28.29 18800.15 +/- 31.02 mab-MXFP8 18368.33 +/- 32.46 18381.75 +/- 33.96 mab-MXFP4 18523.60 +/- 56.23 18580.40 +/- 42.08 mab-NVFP4 18565.33 +/- 28.02 18562.92 +/- 32.80 mab-bf16-m64 16937.42 +/- 44.32 17007.18 +/- 51.41 maf-bf16 3376.61 +/- 0.48 3376.16 +/- 2.53 maf-FP8 9975.45 +/- 36.62 9980.06 +/- 13.52 maf-FP8-henry 10021.30 +/- 6.40 10024.25 +/- 13.15 maf-MXFP8 8956.65 +/- 18.43 8951.98 +/- 9.15 maf-MXFP4 17371.40 +/- 44.54 17332.22 +/- 39.76 maf-NVFP4 16488.60 +/- 14.38 16487.22 +/- 49.88 medium-bf16 2595.23 +/- 2.73 2604.18 +/- 0.89 medium-FP8 6117.63 +/- 55.33 6119.66 +/- 55.57 medium-MXFP8 4991.92 +/- 23.77 5007.50 +/- 10.68 medium-MXFP4 8229.05 +/- 3.92 8197.28 +/- 5.01 ``` ### hipblaslt-test ``` [----------] Global test environment tear-down [==========] 68208 tests from 65 test suites ran. (1102772 ms total) [ PASSED ] 68036 tests. [ SKIPPED ] 172 tests, listed below: [ SKIPPED ] StreamK5HybridDebug.OffModeSelectsStaticTile [ SKIPPED ] StreamK5HybridDebug.OnModeSelectsDynamicWorkQueue [ SKIPPED ] StreamK5HybridDebug.AutoModeRunsHeuristic [ SKIPPED ] StreamK5HybridDebug.OffWithSmCountTargetRunsHeuristic [ SKIPPED ] StreamK5HybridDebug.ForceEnvOverridesOffMode [ SKIPPED ] RowUniformity_pre_checkin.Regression_6144x5120x8192 [ SKIPPED ] RowUniformity_pre_checkin.Regression_3072x3072x12288 [ SKIPPED ] RowUniformity_pre_checkin.Control_4096x4096x8192 [ SKIPPED ] RowUniformity_pre_checkin.Control_8192x8192x8192 [ SKIPPED ] RowUniformity_pre_checkin.ClampRepaired_6144x5120x8192 [ SKIPPED ] RowUniformityClauseTwo_pre_checkin.SplitRegime_4096x32x10240_SmCount80 [ SKIPPED ] RowUniformityClauseTwo_pre_checkin.SplitRegime_4096x32x10240_AllCUs [ SKIPPED ] RowUniformityClauseTwo_pre_checkin.SplitRegime_4096x32x10240_SmCount80_Bias [ SKIPPED ] RowUniformityClauseTwo_pre_checkin.SplitRegime_4096x32x10240_AllCUs_Bias [ SKIPPED ] TuningCache_pre_checkin.NamedEntryReplays [ SKIPPED ] TuningCache_pre_checkin.EntryWhoseNameNoLongerMatchesIsRejected [ SKIPPED ] TuningCache_pre_checkin.StaleFirstEntryDoesNotHideAValidSecond [ SKIPPED ] TuningCache_pre_checkin.NamedEntrySurvivesAnotherBuildsVersionLine [ SKIPPED ] TuningCache_pre_checkin.UnnamedEntryFromThisBuildIsUsed [ SKIPPED ] TuningCache_pre_checkin.UnnamedEntryFromAnotherBuildIsIgnored [ SKIPPED ] TuningCache_pre_checkin.UnnamedEntryWithoutAVersionLineIsIgnored [ SKIPPED ] TuningCache_pre_checkin.TruncatedNamedEntryIsIgnored [ SKIPPED ] TuningCache_pre_checkin.ExtApiIgnoresUnnamedEntryFromAnotherBuild [ SKIPPED ] TuningCache_pre_checkin.ExtApiReplaysNamedEntry [ SKIPPED ] TuningCache_pre_checkin.GroupedGemmUsesDefaultSelection [ SKIPPED ] TuningCache_pre_checkin.SingleAlgoRequestServedByTheFileReturnsOne [ SKIPPED ] TuningCache_pre_checkin.FileCreatedAfterTheFirstQueryIsRead [ SKIPPED ] TuningCache_pre_checkin.CacheModeReplaysANamedRow [ SKIPPED ] TuningCache_pre_checkin.CacheModeReadsAFileCreatedAfterTheFirstQuery [ SKIPPED ] TuningCache_pre_checkin.CacheModeWithoutPathIsInert [ SKIPPED ] TuningCache_pre_checkin.CacheModeIgnoresTheOverrideFile [ SKIPPED ] TuningCache_pre_checkin.CachedEntryWhoseNameNoLongerMatchesIsRejected [ SKIPPED ] TuningCache_pre_checkin.StaleEntryIsCountedAsOneInvalidation [ SKIPPED ] TuningCache_pre_checkin.ExplicitAlgoLaunchesAsGivenInCacheMode [ SKIPPED ] TuningCache_pre_checkin.NullAlgoLaunchesTheCachedKernel [ SKIPPED ] TuningCache_pre_checkin.NullAlgoReplaysAnEntryLoadedAfterItsShapeFellBack [ SKIPPED ] TuningCache_pre_checkin.NullAlgoFindsAFileCreatedAfterItsFirstLook [ SKIPPED ] TuningCache_pre_checkin.LookupTallyCountsShapesNotCalls [ SKIPPED ] TuningCache_pre_checkin.StartupLineReportsWhatLoaded [ SKIPPED ] TuningCache_pre_checkin.RepeatedCacheHitsAddNoOutput [ SKIPPED ] TuningCache_pre_checkin.OverrideFilePrintsNoTuningNotices [ SKIPPED ] TuningStore.NumericLocaleDoesNotChangeReadTiming [ SKIPPED ] AlgoWorkspace.smoke_GemmInitializeRejectsUndersizedWorkspace [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_special_fp_64x64_f32_rf32_rf32_rf32_rxf32_r_NN_64_64_64_1_64_64_0_64_64_1_t3 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_special_fp_64x64_f32_rf32_rf32_rf32_rxf32_r_NN_64_64_64_1_64_64_0_64_64_1_t4 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_tf32_inf_ROCM1545_f32_rf32_rf32_rf32_rxf32_r_NN_64_64_64_1_64_64_0_64_64_1 [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_placement_all_solutions_f32_rf32_rf32_rf32_rf32_r_NN_1031_517_263_2_1031_263_n2_1031_1031_1_PLACE_workspace_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_f32_rf32_rf32_rf32_rf32_r_NN_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_f16_rf16_rf16_rf16_rf32_r_NN_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_NN_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_f32_rf32_rf32_rf32_rf32_r_NT_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_f16_rf16_rf16_rf16_rf32_r_NT_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_NT_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_f32_rf32_rf32_rf32_rf32_r_TN_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_f16_rf16_rf16_rf16_rf32_r_TN_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_TN_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_f32_rf32_rf32_rf32_rf32_r_TT_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_f16_rf16_rf16_rf16_rf32_r_TT_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/pre_checkin_matmul_fast_check_streamk_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_TT_1031_517_263_2_1040_1033_n2_1031_1100_1 [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f32_rf32_rf32_rf32_rf32_r_NN_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f16_rf16_rf16_rf16_rf32_r_NN_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_NN_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f32_rf32_rf32_rf32_rf32_r_NN_4099_2053_129_2_4099_129_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f16_rf16_rf16_rf16_rf32_r_NN_4099_2053_129_2_4099_129_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_NN_4099_2053_129_2_4099_129_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f32_rf32_rf32_rf32_rf32_r_NT_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f16_rf16_rf16_rf16_rf32_r_NT_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_NT_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f32_rf32_rf32_rf32_rf32_r_NT_4099_2053_129_2_4099_2053_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f16_rf16_rf16_rf16_rf32_r_NT_4099_2053_129_2_4099_2053_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_NT_4099_2053_129_2_4099_2053_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f32_rf32_rf32_rf32_rf32_r_TN_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f16_rf16_rf16_rf16_rf32_r_TN_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_TN_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f32_rf32_rf32_rf32_rf32_r_TN_4099_2053_129_2_129_129_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f16_rf16_rf16_rf16_rf32_r_TN_4099_2053_129_2_129_129_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_TN_4099_2053_129_2_129_129_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f32_rf32_rf32_rf32_rf32_r_TT_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f16_rf16_rf16_rf16_rf32_r_TT_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_TT_1031_517_263_2_1040_1033_n2_1031_1100_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f32_rf32_rf32_rf32_rf32_r_TT_4099_2053_129_2_129_2053_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_f16_rf16_rf16_rf16_rf32_r_TT_4099_2053_129_2_129_2053_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_every_solution_repeat_bf16_rbf16_rbf16_rbf16_rf32_r_TT_4099_2053_129_2_129_2053_n2_4099_4099_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_lda_f32_rf32_rf32_rf32_rf32_r_NN_64_64_32_2_69273664_32_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_lda_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_32_2_69273664_32_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_lda_f32_rf32_rf32_rf32_rf32_r_NN_64_64_32_2_69273667_32_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_lda_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_32_2_69273667_32_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_lda_f32_rf32_rf32_rf32_rf32_r_NN_64_64_32_2_138547330_32_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_lda_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_32_2_138547330_32_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_lda_f32_rf32_rf32_rf32_rf32_r_NN_64_64_32_2_138547333_32_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_lda_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_32_2_138547333_32_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldb_f32_rf32_rf32_rf32_rf32_r_NT_64_64_32_2_64_69273664_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldb_bf16_rbf16_rbf16_rbf16_rf32_r_NT_64_64_32_2_64_69273664_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldb_f32_rf32_rf32_rf32_rf32_r_NT_64_64_32_2_64_69273667_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldb_bf16_rbf16_rbf16_rbf16_rf32_r_NT_64_64_32_2_64_69273667_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldb_f32_rf32_rf32_rf32_rf32_r_NT_64_64_32_2_64_138547330_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldb_bf16_rbf16_rbf16_rbf16_rf32_r_NT_64_64_32_2_64_138547330_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldb_f32_rf32_rf32_rf32_rf32_r_NT_64_64_32_2_64_138547333_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldb_bf16_rbf16_rbf16_rbf16_rf32_r_NT_64_64_32_2_64_138547333_n2_64_64_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_32_2_64_64_n2_34087041_34087041_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_NT_64_64_32_2_64_64_n2_34087041_34087041_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_TN_64_64_32_2_64_64_n2_34087041_34087041_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_TT_64_64_32_2_64_64_n2_34087041_34087041_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_32_2_64_64_n2_34087043_34087043_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_NT_64_64_32_2_64_64_n2_34087043_34087043_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_TN_64_64_32_2_64_64_n2_34087043_34087043_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_TT_64_64_32_2_64_64_n2_34087043_34087043_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_32_2_64_64_n2_68174083_68174083_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_NT_64_64_32_2_64_64_n2_68174083_68174083_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_TN_64_64_32_2_64_64_n2_68174083_68174083_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_TT_64_64_32_2_64_64_n2_68174083_68174083_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_32_2_64_64_n2_68174085_68174085_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_NT_64_64_32_2_64_64_n2_68174085_68174085_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_TN_64_64_32_2_64_64_n2_68174085_68174085_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_ldd_bf16_rbf16_rbf16_rbf16_rf32_r_TT_64_64_32_2_64_64_n2_68174085_68174085_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_batch_stride_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_batch_stride_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2_t2 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_batch_stride_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2_t3 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_batch_stride_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2_t4 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_batch_stride_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2_t5 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_batch_stride_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2_t6 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_fast_check_threshold_batch_stride_2_32_ROCM32140_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_fast_check_threshold_batch_stride_2_32_ROCM32140_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2_t2 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_fast_check_threshold_batch_stride_2_32_ROCM32140_bf16_rbf16_rbf16_rbf16_rf32_r_NN_64_64_64_2_64_64_n2_64_64_2_t3 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_f32_f32_rf32_rf32_rf32_rf32_r_NN_2097151_256_48_2_2097151_48_0_2097151_2097151_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_f32_f32_rf32_rf32_rf32_rf32_r_NN_2097153_256_48_2_2097153_48_0_2097153_2097153_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_f32_f32_rf32_rf32_rf32_rf32_r_NN_4194303_256_48_2_4194303_48_0_4194303_4194303_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_f32_f32_rf32_rf32_rf32_rf32_r_NN_4194305_256_48_2_4194305_48_0_4194305_4194305_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_bf16_bf16_rbf16_rbf16_rbf16_rf32_r_NN_4194303_256_48_2_4194303_48_0_4194303_4194303_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_bf16_bf16_rbf16_rbf16_rbf16_rf32_r_NN_4194305_256_48_2_4194305_48_0_4194305_4194305_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_bf16_bf16_rbf16_rbf16_rbf16_rf32_r_NN_8388607_256_48_2_8388607_48_0_8388607_8388607_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_bf16_bf16_rbf16_rbf16_rbf16_rf32_r_NN_8388609_256_48_2_8388609_48_0_8388609_8388609_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_int8_i8_ri8_ri8_ri8_ri32_r_TN_8388607_256_48_1_48_48_0_8388607_8388607_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_int8_i8_ri8_ri8_ri8_ri32_r_TN_8388609_256_48_1_48_48_0_8388609_8388609_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_int8_i8_ri8_ri8_ri8_ri32_r_TN_16777215_256_48_1_48_48_0_16777215_16777215_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_int8_i8_ri8_ri8_ri8_ri32_r_TN_16777217_256_48_1_48_48_0_16777217_16777217_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_all_solutions_bf16_rbf16_rbf16_rbf16_rf32_r_NN_4194303_256_48_2_4194303_48_0_4194303_4194303_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_d_bytes_all_solutions_bf16_rbf16_rbf16_rbf16_rf32_r_NN_4194305_256_48_2_4194305_48_0_4194305_4194305_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_fast_check_threshold_cd_4gib_ROCM32050_f32_rf32_rf32_rf32_rf32_r_NN_32768_32767_4096_2_32768_4096_n2_32768_32768_1 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_fast_check_threshold_cd_4gib_ROCM32050_f32_rf32_rf32_rf32_rf32_r_NN_32768_32768_4096_2_32768_4096_n2_32768_32768_1 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_fast_check_threshold_cd_4gib_ROCM32050_f32_rf32_rf32_rf32_rf32_r_NN_32768_32769_4096_2_32768_4096_n2_32768_32768_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_tiles_bf16_rbf16_rbf16_rbf16_rf32_r_NN_65536_65536_16_2_65536_16_0_65536_65536_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_tiles_bf16_rbf16_rbf16_rbf16_rf32_r_NN_65536_65552_16_2_65536_16_0_65536_65536_1 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_tiles_streamk_bf16_rbf16_rbf16_rbf16_rf32_r_NN_65536_65536_16_2_65536_16_0_65536_65536_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_tiles_streamk_bf16_rbf16_rbf16_rbf16_rf32_r_NN_65536_65552_16_2_65536_16_0_65536_65536_1_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_batch_count_bf16_rbf16_rbf16_rbf16_rf32_r_NN_16_16_16_2_16_16_n2_16_16_65535 [ SKIPPED ] _/matmul_test.matmul/stress_matmul_fast_check_threshold_batch_count_bf16_rbf16_rbf16_rbf16_rf32_r_NN_16_16_16_2_16_16_n2_16_16_65537 [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_batches_f32_rf32_rf32_rf32_rf32_r_NT_1031_517_2049_2_1031_517_n2_1031_1031_3_IE_sparse_k_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_batches_bf16_rbf16_rbf16_rbf16_rf32_r_NT_1031_517_2049_2_1031_517_n2_1031_1031_3_IE_sparse_k_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_batches_f32_rf32_rf32_rf32_rf32_r_NT_4099_2053_1025_2_4099_2053_n2_4099_4099_3_IE_sparse_k_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/nightly_matmul_fast_check_streamk_batches_bf16_rbf16_rbf16_rbf16_rf32_r_NT_4099_2053_1025_2_4099_2053_n2_4099_4099_3_IE_sparse_k_APIFindAllAlgo [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_f32_rf32_rf32_rf32_rf32_r_NT_4096_8192_4096_2_4096_8192_n2_4096_4096_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_bf16_rbf16_rbf16_rbf16_rf32_r_NT_4096_8192_4096_2_4096_8192_n2_4096_4096_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_bf16_rbf16_rf32_rf32_rf32_r_NT_4096_8192_4096_2_4096_8192_n2_4096_4096_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_f32_rf32_rf32_rf32_rf32_r_NT_8192_16384_8192_2_8192_16384_n2_8192_8192_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_bf16_rbf16_rbf16_rbf16_rf32_r_NT_8192_16384_8192_2_8192_16384_n2_8192_8192_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_bf16_rbf16_rf32_rf32_rf32_r_NT_8192_16384_8192_2_8192_16384_n2_8192_8192_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_f32_rf32_rf32_rf32_rf32_r_NT_12288_20480_12288_2_12288_20480_n2_12288_12288_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_bf16_rbf16_rbf16_rbf16_rf32_r_NT_12288_20480_12288_2_12288_20480_n2_12288_12288_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_bf16_rbf16_rf32_rf32_rf32_r_NT_12288_20480_12288_2_12288_20480_n2_12288_12288_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_fp8_f8_rf8_rf32_rf32_rf32_r_TN_2560_16384_2560_2_2560_2560_n2_2560_2560_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_fp8_f8_rbf8_rf32_rf32_rf32_r_TN_2560_16384_2560_2_2560_2560_n2_2560_2560_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_fp8_bf8_rf8_rf32_rf32_rf32_r_TN_2560_16384_2560_2_2560_2560_n2_2560_2560_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_fp8_bf8_rbf8_rf32_rf32_rf32_r_TN_2560_16384_2560_2_2560_2560_n2_2560_2560_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_fp8_f8_rf8_rf32_rf32_rf32_r_TN_5120_12288_5120_2_5120_5120_n2_5120_5120_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_fp8_f8_rbf8_rf32_rf32_rf32_r_TN_5120_12288_5120_2_5120_5120_n2_5120_5120_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_fp8_bf8_rf8_rf32_rf32_rf32_r_TN_5120_12288_5120_2_5120_5120_n2_5120_5120_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/sdc_hunt_matmul_fast_check_sdc_hunt_fp8_bf8_rbf8_rf32_rf32_rf32_r_TN_5120_12288_5120_2_5120_5120_n2_5120_5120_1_IE_sparse_k [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_batch_offset_transpose_bf16_rbf16_rbf16_rbf16_rf32_r_TN_256_256_128_1_256_256_0p5_256_256_3 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_batch_offset_all_solutions_bf16_rbf16_rbf16_rbf16_rf32_r_TN_256_256_128_1_256_256_1_256_256_2 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_batch_offset_negative_bf16_rbf16_rbf16_rbf16_rf32_r_TN_256_256_128_1_256_256_1_256_256_2 [ SKIPPED ] _/matmul_test.matmul/known_bug_matmul_batch_offset_mixed_bf16_rbf16_rbf16_rbf16_rf32_r_TN_256_256_128_1_256_256_1_256_256_2 hipBLASLt version: 100500 hipBLASLt git version: 5a269f7 command line: ./clients/hipblaslt-test -j ``` ### tox ``` =============================== warnings summary =============================== tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.custom - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.characterization - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.LibraryIO - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.data - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.TensileCreateLibraryRun - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.gfx90a - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.gfx908 - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.gfx1201 - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.gfx942 - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.gfx1100 - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.bigfiles - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.gfx940 - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.gfx90 - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.skip-stinkytofu - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.core - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) tensilelite/Tests/common/config_helpers.py:93 tensilelite/Tests/common/config_helpers.py:93 /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests/common/config_helpers.py:93: PytestUnknownMarkWarning: Unknown pytest.mark.benchmark_runs - is this a typo? You can register custom marks to avoid this warning - for details, see https://docs.pytest.org/en/stable/how-to/mark.html return getattr(pytest.mark, name) -- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html - generated xml file: /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/python_tests.xml - --------------------------- snapshot report summary ---------------------------- 12 snapshots passed. = 776 passed, 1474 skipped, 3 xfailed, 1 xpassed, 32 warnings in 2372.28s (0:39:32) = py3: exit 0 (2452.23 seconds) /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite> sh -c 'ARCH=$(invoke get-gpu-arch); export TENSILELITE_CLIENT_ARGS='"'"' --gpu-targets '"'"'$ARCH; invoke build-client --build-dir /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/build_tmp --rocm-path "$ROCM_PATH" $TENSILELITE_CLIENT_ARGS || exit $?; pytest -v --basetemp=/home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/.tensile-tox/py3/tmp --junit-xml=/home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/python_tests.xml --junit-prefix=py3 --color=yes -n $TENSILE_NUM_PYTEST_WORKERS --prebuilt-client=/home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/build_tmp/tensilelite/client/tensilelite-client --global-parameters CheckASMCodeSize=True /home/pochang/rocm-libraries.2/projects/hipblaslt/tensilelite/tensilelite/Tests Tensile/Tests -k '"'"'gfx1250 and not sk_mxf4_force_dp_only_cluster_multicast_ext and not sk_mxf8_force_dp_only_pap_cluster_multicast_ext and not sk_mxf4_force_dp_only_pap_cluster_multicast_ext and not sk_mxf8_force_dp_only_cluster_multicast_ext'"'"'; rc=$?; test $rc -eq 5 && exit 0; exit $rc' pid=1732823 py3: OK (2508.95=setup[9.29]+cmd[0.30,39.56,7.38,0.19,2452.23] seconds) congratulations :) (2508.98 seconds) ``` ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
yassinsolim
added a commit
that referenced
this pull request
Nov 2, 2026
… for #13373 (#14038) JIRA ID: ALMIOPEN-2090 ## Stacked on the wrong-results fix This branch sits on #14037, `users/ysoliman/gfx1200-wrw-drop-wrong-results`, and carries that PR's two commits until it merges. Review only the last two commits here: the database commit and the changelog commit. The two changes touch disjoint records, checked record by record on the written files. ## What changed 475 `ConvHipImplicitGemmGroupWrwXdlops` entries go back into the gfx1200 system database, `gfx120020.HIP.fdb.txt.bz2` and `gfx120020.db.txt.bz2`, with perf configs re-tuned on gfx1200 hardware and re-measured against a TheRock CI build. - find-db: 475 records gain one entry each. Record count is unchanged at 90,688. - perf-db: 470 records, of which 193 are extended and 277 are recreated. #13697 had deleted those 277 outright because the CK entry was the only thing they held. Record count goes from 77,298, the count the parent branch leaves, to 77,575. - Every restored perf config carries an explicit `split_k`: 391 at 1, 46 at 2, 24 at 4, 12 at 8 and 2 at 16. None uses auto split-K. - Every other record is untouched. The build script appends text and never re-serializes an existing entry, and it asserts per record that the new content is the old content plus exactly the one appended entry. - Combined with the parent commits, the database this branch produces has 90,688 find-db records and 77,575 perf-db records: 475 records gain a CK entry, one record loses the CK entry that returned wrong weight gradients, and nothing else moves. ## Why #13697 removed the CK backward-weights entry from 838 find-db keys and 829 perf-db records in this database. The configs behind them are all `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted` instances that CK rejects after #13373, which fixed a wave32 support check that read `GridwiseGemm64::BlockwiseGemmPipe::PrefetchStages` where it should read `GridwiseGemm32`. Before that fix those kernels could return without writing the weight gradient, which is the all-zero `dw` in #13305. Removing the entries was correct and had to ship, but it left those problems on whatever solver was next in the record. Measured on a TheRock build, that fallback is 2.64x slower in geomean over the 475 restored keys and up to 50x slower on the worst shape. 835 of the 838 keys went through the re-tune; the candidate list was built before the last three were spotted, so they were never tuned. 475 came back with a config that is both correct and at least 5 percent faster than the fallback. The other 363 keys, the 360 the re-tune did not bring back plus those 3, are left exactly as #13697 left them. ## Why every split_k is pinned The tuner can return `split_k = -1`, which asks CK to derive `k_batch` at run time from an occupancy query. That query is answered by the device and the compiled kernel, so two builds of the same source can disagree, and `TestPerfCfgParams` then rejects an entry that the tuning build accepted. This is not hypothetical. The first candidate database for this change carried 209 auto split-K entries. It passed `StaticFDBSync/10` on the from-source build it was tuned on and failed on a TheRock build with Invalid perf config found fdb-key:1024-6-6-3x3-1024-14-14-2-0x0-2x2-1x1-0-NCHW-FP16-W_g32, Solver: ConvHipImplicitGemmGroupWrwXdlops: DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted<64, 32, 128, 32, Default, 8, 1, 4, 8, 8, 8, 8, 1, 1, 2>+-1 which is the same class of nightly failure #13697 had to fix. So every one of those 209 keys was re-measured at `split_k` 1, 2, 4, 8, 16, 32, 64 and 128 on the TheRock build, and the fastest value that verifies was taken. Pinning costs nothing measurable: over the 208 keys where both forms run, the auto time divided by the pinned time has geomean 1.031 and median 1.001, and 147 of 208 are within 2 percent. 187 of the 209 chose `split_k = 1`. ## How the entries were produced Three builds are used. - TheRock CI run 38988799797 (linux), which pins rocm-libraries 8a29d2e. That commit is 18 commits past the #13373 merge. This is the build the tuning numbers below were measured on, because it is the build shape nightly CI runs. - A from-source build of develop 6bb613a for gfx1200 inside `rocm/dev-ubuntu-22.04@sha256:80de159c`, CK with `MIOPEN_REQ_LIBS_ONLY`, then MIOpen, `MIOpenCKGroupedConv_gfx1200`, `test_db_sync` and `MIOpenDriver`. - TheRock main CI run 39445395419 (linux), which pins rocm-libraries cb26189. This is the build every check in "How it was checked" was re-run on. The gfx120020 database files are the same git blobs at 8585227, at 6bb613a and at develop 03f8851 (`aff3af95` for the find-db, `2c9fa01a` for the perf-db), and both TheRock builds ship that pair unchanged. That pair is the baseline every check below starts from. Per key, every run is a fresh process with the gfx120020 find-db and perf-db hidden unless stated: 1. Fallback. Shipped post-#13697 database, `MIOPEN_FIND_MODE=2`, empty user database. On a find-db hit modes 2, 3 and 5 all take the recorded entry, and every one of these keys is a hit, so this is what a user gets today. 2. Bench. The candidate perf config written into a fresh user perf-db, `MIOPEN_FIND_MODE=1` restricted to `ConvHipImplicitGemmGroupWrwXdlops`. MIOpen benchmarks that one solution and writes the user find-db record, which carries the workspace size and the algorithm. 3. Replay. A fresh process reading both user database files back, so only the recorded winner runs. A Find benchmarks every candidate into the real output buffer, so a Find pass can report a verification the winner alone would not reproduce. Only the replay proves the recorded winner is the one that works. Every driver invocation used `-V 1 -N 1 -t 1 -i 10`. `-N 1` fills the weight-gradient buffer with NaNs first, which is what makes a kernel that writes nothing fail instead of inheriting a previous result. Any run that logged `Invalid config loaded from Perf Db`, or that loaded a config other than the one asked for, was discarded. The find-db and perf-db record strings come out of MIOpen's own `.ufdb.txt` and `.udb.txt`, never hand-built. ## Inclusion rule A key is added back only if all of these hold, otherwise the record is left exactly as #13697 left it: 1. the replay returned `ConvHipImplicitGemmGroupWrwXdlops`; 2. the replay verifies against the reference; 3. the find-db key MIOpen wrote equals the key being targeted; 4. the perf-db key MIOpen wrote equals the one derived from the find-db key; 5. the replay is at least 5 percent faster than the measured fallback. Where the gain landed between 0 and 10 percent, both sides were re-measured as the median of 3 runs. Results on the TheRock build: 475 accepted, 8 rejected, all 8 on the 5 percent bar. Speedup over the fallback for the accepted set: geomean 2.645, median 2.214, p10 1.240, p90 8.687, min 1.054, max 50.23. ## Recorded times A record's times were measured on the machine that originally tuned it, so an absolute millisecond figure measured here would not be on the same scale. The fallback is measured both ways, here and in the record, so its ratio carries this node's measurement onto the record's scale: t_insert = t_fallback_recorded * (t_replay_here / t_fallback_here) Because the fallback is by construction the fastest entry left in each record, and the replay beats it, the inserted time is the new minimum. MIOpen sorts find-db entries by time through `SolutionTimeComparator` before choosing, so CK ranks first. This is asserted per record at build time and confirmed on the written file for all 475. ## Perf-db key collisions 9 pairs among the 835 re-tuned keys share a perf-db key, 18 keys mapping to 826 distinct perf-db keys, because the perf-db key carries the output spatial dims and a 13x13 and a 14x14 input with 3x3, pad 0, stride 2 both produce 6x6. A perf-db record holds one config per solver. In this set the colliding keys agreed on the config, so none had to be dropped. 5 of those pairs are in the accepted set, which is why 475 restored find-db entries land in 470 perf-db records. ## What could break - The 475 configs are validated on one gfx1200 part, which is the part `gfx120020` exists for. They are not validated on any other gfx1200 SKU. - Timings come from one machine. The ordering inside a record is only as good as that measurement, so a record whose top two entries are within a few percent could pick differently elsewhere. The 5 percent bar exists to keep those out. - The entries are new instances, not the ones that were removed. `DeviceGroupedConvBwdWeight_Wmma_CShuffleV3`, `DeviceGroupedConvBwdWeight_Xdl_CShuffle` and `DeviceGroupedConvBwdWeight_Explicit_Xdl` now appear where the removed ones were all `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3_WmmaPorted`. Any CK change to those families moves these records. - Workspace sizes change for the affected problems, since CK entries carry a workspace and several of the displaced GEMM and Winograd entries carry zero. One related defect, fixed by the parent commits rather than by this one: the exhaustive search produced a wrong result on 4 bf16 keys, all `384-18-18-3x3` shapes, all landing on `DeviceGroupedConvBwdWeight_Wmma_CShuffleV3<96, 96, 96, 48, ...>` with relative error 0.24 to 0.29. Those 4 are excluded by rule 2. The post-#13697 gfx120020 perf-db also used that exact instance in two bf16 entries: `96-18-18-3x3-96-35-35-256` (pad 1, stride 2, split_k 8) and `192-23-1-3x1-192-25-1-1` (stride 1, split_k 1). Both were run. For `96-18-18-3x3-96-35-35-256` the solver that actually ran by default was `ConvHipImplicitGemmGroupWrwXdlops` with that config, in `MIOPEN_FIND_MODE=2` and with the mode unset, three runs each, and it failed verification every time with relative error about 0.306. For `192-23-1-3x1-192-25-1-1` the solver that actually ran is `GemmWrwUniversal`, which verifies; forcing the CK entry on that key verifies as well, so that config is correct for that shape. All 765 records in that database where `ConvHipImplicitGemmGroupWrwXdlops` is the fastest entry were checked the same way on both builds, and `96-18-18-3x3-96-35-35-256` was the only failure. #14037 removes it, and this branch is stacked on that change, so the database here no longer serves it. The CK defect is #14039. ## How it was checked Hardware: a gfx1200 machine with two GPUs, which is the part `gfx120020` serves. Not the single-GPU gfx1200 node used earlier, which has been drained since 2026-10-15 with a prolog error. Every check in this section was re-run on TheRock main CI run 39445395419, against the database this branch produces, which is the parent commits plus this one: the decompressed form of the two `.bz2` blobs this branch carries, find-db sha256 `5b7b43c4675921b9ad7f65fd824c1dbf4c2b1db71e5b7e88fe4079b74af795e7`, perf-db sha256 `431b64f7666b87a842cb768d02ba79a99a017ac24e4a65b5fde74a7b1b36fe9b`. `miopen_gtest` sha256 `f377deb71a777a1da5ecba0bd3404d8a77aa19bb8f32dbce28a8e9170feb9cd6`, `MIOpenDriver` sha256 `56c748f2a962390d33cc95af8a6b3a833e027eddc0cfbaa11a6d582d974e4df9`. - `StaticFDBSync/10` passes on the combined database: zero invalid perf configs, zero not-applicable solvers, zero missing perf-db entries, zero invalid solutions. It passes the same way on the database #13697 shipped and on the parent branch's database, both run in the same session, so the pass is not vacuous. - The other eleven `StaticFDBSync` parameters skip on this hardware, the same on all three databases. - All 476 changed keys, a fresh process each, with `-V 1 -N 1`, both in `MIOPEN_FIND_MODE=2` and with the mode unset. The 475 restored keys each select `ConvHipImplicitGemmGroupWrwXdlops` with exactly the config in the database, verify, and log no `Invalid config loaded from Perf Db` warning. `96-18-18-3x3-96-35-35-256-1x1-2x2-1x1-0-NCHW-BF16-W` selects `GemmWrwUniversal` and verifies. 476 of 476 in each mode. - Positive control in the same session: that key on the unmodified post-#13697 database still selects `ConvHipImplicitGemmGroupWrwXdlops` and still fails verification, relative error 0.3056 to 0.3060, three runs in each mode, so the pick check does catch a wrong result. - A database-loading smoke subset (`CPU_DbPaths_NONE`, `CPU_PerfDb_NONE`, `CPU_SysDbSelectionCu_NONE`, `CPU_Cache_NONE`, `CPU_DBSync_NONE.KDBTargetID`), with the pair copied into the build's own database directory: 39 passed, 1 skipped (`KDBTargetID`, the artifact carries no kernel database), identical on this database, on the parent branch's and on develop's. - The two databases differ from develop's only where they should, checked record by record on the decompressed files: the parent commits change one find-db record and remove one perf-db record, this branch appends one entry to 475 find-db records and to 193 perf-db records and recreates 277 perf-db records, and the two sets of keys do not overlap. Every appended find-db entry is the fastest in its record. Earlier, on TheRock CI run 38988799797 and the from-source build of develop 6bb613a, the same `StaticFDBSync/10` and 476-key pick checks passed in both modes, alongside these controls: - Negative control 1: the earlier candidate database, identical except that one entry keeps its auto split-K, fails `StaticFDBSync/10` on the TheRock build with the `Invalid perf config` line quoted above, while passing on the from-source build. - Negative control 2: a copy of this database with one pinned entry changed to a `split_k` CK rejects for that problem fails `StaticFDBSync/10` on both builds. - All 483 fallbacks verified correctly, so no fallback correctness problem was reproduced there. Closes nothing on its own. Follows #13697, depends on #13373, and restores the performance that #13305 and #13373 forced us to give up.
shumway
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Nov 4, 2026
fix(ck): PrefetchStages on host for wave32 instances ## Motivation ISSUE: #13305 ## Technical Details `DeviceGroupedConvBwdWeight_Xdl_CShuffleV3` uses `GridwiseGemm64::BlockwiseGemmPipe::PrefetchStages` on RDNA (wave32) in `IsSupportedArgument`. This is incorrect for pipeline V2 where the PrefetchStages depend on the wave size. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
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.
Motivation
ISSUE: #13305
Technical Details
DeviceGroupedConvBwdWeight_Xdl_CShuffleV3usesGridwiseGemm64::BlockwiseGemmPipe::PrefetchStageson RDNA (wave32) inIsSupportedArgument. This is incorrect for pipeline V2 where the PrefetchStages depend on the wave size.Submission Checklist