Repository navigation
DO NOT MERGE: CI-only rerun of reverted perf-db cleaning content (flake vs regression check) - #11143
Closed
yassinsolim wants to merge 17 commits into
Closed
yassinsolim wants to merge 17 commits into
yassinsolim wants to merge 17 commits into
Conversation
ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC (the ASM-GTC forward solver from the unmaintained MISA project) indexes global tensor memory with 32-bit element indices and does not implement large-tensor support. Its IsApplicable() relied on AllTensorsDimsFitIntoInt(), which only checks that each individual length/stride fits in int32 -- it never bounds the flattened element count. A tensor whose total element count exceeds INT_MAX therefore overflows the kernel's 32-bit indexing and is silently computed incorrectly (observed forward verification error ~0.37 on gfx942 and gfx950). Gate the solver not-applicable when any tensor's element count exceeds INT_MAX, so a large-tensor-capable solver (e.g. grouped CK xdlops) is selected instead. This is a no-op for every int32-safe shape, so existing workloads are unaffected. Adds a CPU device-applicability unit test that shares geometry across a batch-size boundary so the element count crosses INT_MAX, asserting the solver stays applicable just below the limit and is gated off just above. Related to ROCM-27526.
The preceding commit gates ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC not-applicable when any tensor's element count exceeds INT_MAX. The shipped find-dbs still list that solver for such shapes -- it had been recorded there (in some entries as the fastest kernel) even though its 32-bit element indexing computes those shapes incorrectly. The db_sync FDB consistency test asserts every solver listed in a find-db entry is still IsApplicable(), so the now-gated entries fail Dbsync on gfx908/gfx90a/gfx942/gfx950. Prune the solver from exactly the find-db entries whose input or output element count exceeds INT_MAX, mirroring the gate and the MIOPEN_DBSYNC_CLEAN removal path (cf. #3725, "Clean db files to pass db_sync"). Entries at or below INT_MAX are left untouched; lines where it was the only listed solver are dropped. gfx90878.HIP.fdb.txt : 43 gfx90a68.HIP.fdb.txt : 102 gfx90a6e.HIP.fdb.txt : 102 gfx942130.HIP.fdb.txt : 268 gfx942e4.HIP.fdb.txt : 206 gfx950100.HIP.fdb.txt : 309 Related to ROCM-27526.
ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC shares the MISA 32-bit element indexing of the forward solver already gated in this PR. AllTensorsDimsFitIntoInt() bounds each length and stride but not the flattened element count, so a packed NHWC tensor with a large batch passes the check while the total overflows and the kernel silently returns wrong results. Apply the same guard.
ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC shares the MISA 32-bit element indexing of the forward solver already gated in this PR. AllTensorsDimsFitIntoInt() bounds each length and stride but not the flattened element count, so a packed NHWC tensor with a large batch passes the check while the total overflows and the kernel silently returns wrong results. Apply the same guard.
Mirrors the forward test: same geometry, N=140 vs N=141 brackets the INT_MAX element-count boundary, so the int32-safe case must stay applicable and the over-threshold case must be applicable on no device.
Mirrors the forward test: same geometry, N=140 vs N=141 brackets the INT_MAX element-count boundary, so the int32-safe case must stay applicable and the over-threshold case must be applicable on no device.
Matches PR #10020, already merged to release/rocm-rel-7.0.2.1, which carries the same bwd/wrw gate with a single test file. Taking that file verbatim keeps the two branches identical and defines the four shared helpers once instead of three times.
…ded by the consolidated file
…ded by the consolidated file
…ded by the consolidated file
The gates added here compared each tensor's element count against INT_MAX, but these assembly kernels address global memory with 32-bit byte offsets. For fp16 that lets shapes up to twice the safe size through, and fp32 up to four times, where they return silently wrong results. Compare GetNumBytes() instead, matching the byte-based form already used for the equivalent CK solver guard in conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp. The unit test's boundary cases move with the threshold: for the shared 1024x162x92 fp16 geometry the effective element ceiling halves, so the in-range and gated cases become N=70 (2,136,637,440 bytes) and N=71 (2,167,160,832 bytes) rather than N=140 and N=141.
…erflow gate The previous commit changed the ASM-GTC fwd/bwd/wrw NHWC solvers' overflow guard from an element-count comparison to a byte-count comparison, because these MISA-derived assembly kernels address global memory with 32-bit byte offsets, not 32-bit element offsets. develop ships no gate at all on these solvers, so its find-db files contain tuned rows that name one of these solvers as the winner for shapes whose input, output, or weight tensor now exceeds INT_MAX bytes. Those rows were already silently computing wrong results before this PR; the new gate only makes StaticFDBSync notice. Removed the exact (key, solver) pairs identified by a static byte-accurate parser that mirrors ParseProblemKey()/IsApplicable() in db_sync.cpp and the three solver files, cross-validated key-for-key (zero difference either direction) against IsApplicable() run live on gfx942/MI300X-304CU hardware in CI. Re-parsing every cleaned file with the same tool confirms zero stale rows remain. A key keeps every entry that is still valid; a line is dropped outright only when removing its stale entries would leave it empty. Per-file counts (stale pairs removed / find-db lines edited in place / lines dropped entirely): gfx90878 1210 / 1209 / 1 gfx90a68 1227 / 1202 / 25 gfx90a6e 1244 / 1228 / 16 gfx942130 3206 / 3200 / 6 gfx942e4 2654 / 2654 / 0 gfx950100 3366 / 3333 / 33 Perf-db cleanup for the same shapes is a separate commit.
…NT_MAX bytes Redo of the perf-db half of this PR's database cleanup. The prior attempt at this commit (local SHA 3dd0181, never merged, force-pushed away) was reverted after two apparently unrelated tests failed in CI: Full/GPU_LstmMSRnn_FP32.TestFloat/1357 (gfx942) and Smoke/GPU_UnitTestConvSolverTransposedWinograd3x3MultipassF1x1x7x3Wrw_FP32 .../0 (gfx90a). This commit re-derives the exact same perf-db content with a strictly stronger invariant and states the root-cause investigation for the two failures. Root cause of the LSTM/Winograd failures, with evidence: Neither failure traces to this commit's database content. Three independent checks confirm it: 1. A byte-for-byte row/config diff of every one of the 6 perf-db files between this PR's find-db-only commit (03ef09d, verified, pushed, untouched) and the reverted perf-db commit (3dd0181) shows every single row added, removed, or edited names one of the three gated ASM-GTC solvers. Zero rows for any other solver were touched in any of the 6 files, and neither "rnn"/"lstm" nor "winograd3x3multipass"-family solver names appear anywhere in any of the 6 shipped perf-db files, before or after, on any architecture (grepped directly, both the 5 text files and the gfx90a6e SQLite perf_db table). 2. The failing Winograd test calls solver::conv::SolverBaseTunable:: FindSolutionSimple(), whose doc comment in src/include/miopen/solver.hpp (line ~172) states plainly: "does not obey search parameters from the Context and does not use the database. Intended to be used in unit tests." That test cannot observe perf-db content by construction, so it cannot have been broken by a database change. 3. Full/GPU_LstmMSRnn_FP32.TestFloat is a documented flaky suite on gfx942: PR #7220 disabled a different index of this exact test (1095) in May for observed intermittent failure on gfx942 specifically (confirmed passing on gfx90a in that same investigation). Index 1357 is not covered by that disablement, so the same known-flaky suite failing on the same hardware under a different parameter index is consistent with pre-existing flakiness rather than a new regression. Both "Fp32 Hip Install All gfx942" and "...gfx90a" jobs passed cleanly on 03ef09d and only failed on 3dd0181 (confirmed via the GitHub check-runs API against both SHAs), so something about that CI run did fail. Given (1)-(3) above, the evidence points at ordinary flake landing on an unlucky run rather than a data-driven regression, but this is stated as the best evidence-supported conclusion available offline, not as a certainty beyond doubt. The corrected invariant: The prior attempt's invariant checked only "does every surviving find-db entry naming one of the three gated solvers still have a matching perf-db row." That is necessary but not sufficient: it never checked whether a perf-db row about to be deleted might still be needed by some OTHER live find-db entry, of any solver, that happens to derive the identical perf-db key (the perf-db key is a positional field dump that drops out_h/out_w, so two structurally different find-db keys can legitimately collide onto one perf-db row; the prior attempt already found and protected 82 such rows for backward-stride convolutions). This commit runs the broadened check: before deleting any perf-db (key, solver) row, confirm no surviving find-db entry for ANY solver still maps to that exact perf-db row under the same key-derivation logic, not just entries naming one of the three gated solvers. Concretely, this means building the live (perfkey, solver) claim set from a full parse of every find-db line's every solver (not just the three gated ones, and across both 2D and 3D key shapes), then checking each row slated for deletion against that full set rather than a target-solver-only subset. Run as the two required offline passes, both against all six shipped per-architecture combinations (5 text find/perf-db pairs plus gfx90a6e's SQLite perf_db): BASELINE (before any deletion, every solver, current develop-derived content at 03ef09d): 2,089,857 tunable-solver (find-db key, solver) pairs checked, 2 pre-existing missing pairs, both on gfx90a68, both solver ConvMlirIgemmFwdXdlops, unrelated to ASM-GTC and pre-existing in develop content, not introduced by this change or the find-db commit. AFTER (post corrected cleaning, every solver): identical scope, checked against the newly-cleaned perf-db content. Zero NEW missing pairs on any architecture; the same 2 gfx90a68/ConvMlirIgemmFwdXdlops pairs remain the only gap, unchanged from baseline. The broadened, all-solver reverse invariant protects exactly the same 82 rows the narrower one protected (13 gfx90878, 0 gfx90a68, 4 gfx90a6e, 21 gfx942130, 22 gfx942e4, 22 gfx950100), and the resulting cleaned perf-db content is byte-for-byte identical, per architecture, to the content in the reverted commit 3dd0181 (verified directly: sorted text diff empty on all 5 text files, and identical (solver, config, params) row sets on the gfx90a6e SQLite perf_db table). This is independent confirmation that the original cleaning was not missing any collision class: the LSTM and Winograd solvers, having zero entries in any of these files, were never candidates for deletion or collision in the first place, under either the narrow or the broadened check. Per-file counts (of N stale target-solver pairs, M protected because a surviving entry of any solver still needs that exact perf-db row, same as last time): gfx90878 1101 removed, 13 protected (of 1210 stale) gfx90a68 1162 removed, 0 protected (of 1227 stale) gfx90a6e 1173 rows deleted (SQLite), 4 protected (of 1244 stale) gfx942130 2887 removed, 21 protected (of 3206 stale) gfx942e4 2492 removed, 22 protected (of 2654 stale) gfx950100 3033 removed, 22 protected (of 3366 stale) Total 11,848 perf-db rows removed, matching the reverted commit exactly. LSTM/Winograd row survival, explicitly checked: grepped all 6 perf-db files (5 text, plus the gfx90a6e SQLite perf_db.solver column) for any solver name matching "Lstm", "Rnn"/"RNN", "MultipassWrW", or "Winograd3x3Multipass" -- zero matches in every file, before and after this change. There is no perf-db row for either failing test's solver to survive or be dropped; both tests are structurally outside this cleaning's scope, confirmed by direct lookup rather than inference. Format preservation: every file decompresses cleanly (bzip2 -t passes on all 6), recompressed with bzip2 -9 (matching the existing "BZh9" header block size on every shipped file), and every non-target-solver row across all 6 files is verified byte-identical / untouched (line-for-line for the 5 text files, full (id, solver, config, params) row-set comparison for the gfx90a6e SQLite perf_db table; the config table's 134,579 rows are untouched, count-for-count and value-for-value). Built on top of 03ef09d (verified, pushed, untouched by this commit).
✅ 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. |
Member
Author
|
CI-only verification complete. Fp32 Hip Install All gfx942 and gfx90a both passed clean against database content byte-identical to the previously reverted 3dd0181 (confirmed at the individual-test level: Full/GPU_LstmMSRnn_FP32.TestFloat/1357 and Smoke/GPU_UnitTestConvSolverTransposedWinograd3x3MultipassF1x1x7x3Wrw_FP32.../0 both OK). Closing without merging, as stated when this was opened. |
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.
JIRA ID: ALMIOPEN-2367
This draft exists only to trigger Jenkins CI against database content byte-identical to the previously reverted commit 3dd0181, to determine whether the two test failures that caused that revert (Full/GPU_LstmMSRnn_FP32.TestFloat/1357 on gfx942, and the transposed Winograd multipass smoke test on gfx90a) were a genuine content-driven regression or CI flake unrelated to the database change. It will be closed without merging once Fp32 Hip Install All gfx942 and Fp32 Hip Install All gfx90a report a result. Not for review. Relates to PR #9568.