Repository navigation
fix(miopen): remove gfx803-tuned solvers and stale kernel sources - #11761
Merged
averinevg merged 7 commits intoSep 21, 2026
Merged
Conversation
Removes the ConvBinWinogradRxSFused solver and gfx803-specific code paths/kernel sources from ConvBinWinograd3x3U and ConvBinWinogradRxS, along with unused gfx803 find-database files. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…ove-gfx803-tuned-solvers-and-kernels
This env var gated the ConvBinWinogradRxSFused solver, which was already removed from the codebase. Co-Authored-By: Claude Sonnet 5 <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. |
…ove-gfx803-tuned-solvers-and-kernels
…ove-gfx803-tuned-solvers-and-kernels
The gfx803 solver/kernel removal is unreleased work for the upcoming 10.2.0 release, not the already-cut 10.1.0 one, so move the three Removed bullets accordingly and align them with the [Conv] prefix convention used elsewhere in the changelog. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
iq136boy
approved these changes
Sep 18, 2026
averinevg
deleted the
users/evaverin/remove-gfx803-tuned-solvers-and-kernels
branch
September 21, 2026 11:30
Contributor
TheRock Submodule Bump ActivityNewest first
|
mattwill-amd
pushed a commit
that referenced
this pull request
Oct 14, 2026
JIRA ID: ALMIOPEN-2913 ## Motivation Only 5 of the 30 commits that landed against `projects/miopen` since the 3.6.2 version bump (`2d20a8f775e`, #12211) were reflected in `CHANGELOG.md`. The remaining 25 were reviewed and the 15 externally-visible ones documented, so the ROCm 10.2.0 section accurately reflects what shipped. 14 entries added (two of the fixes share one entry), plus a duplicate-heading fix. ### Added * [Conv] NHWC (channel-last) support in the GEMM convolution solvers, with new im2col/col2im kernels (#8693) * [Conv] Cross-architecture AI heuristics for immediate mode (`MIOPEN_FIND_MODE=2`), extending data-driven selection to Navi (#9680) * [Conv] Missing 3x3 stride-1 pad-1 tile-28 CK depthwise backward-data instance — ALMIOPEN-2423 (#11271) ### Changed * [Conv] `ConvHipConv` held out of the AI perf-config picker until its model is retrained — ALMIOPEN-2494 (#13337) * [Conv] CK byte-span gate lifted for forward and backward-data, now that CK fixes the 32-bit offset overflow itself — ALMIOPEN-2645 (#12396) ### Optimized * [Conv] `ConvDirectNaiveConv`: per-direction kernel files, float accumulators, spatial tiling, bfloat16 path — ALMIOPEN-2485/2486 (#7529) * Kernel cache reuses cached SQLite connections instead of reopening per load/store — ALMIOPEN-2598 (#11933) ### Resolved Issues * [Conv] Incorrect backward-weights results from the gfx950 direct kernel at `waves_q == 1` (missing cross-lane ordering on LDS staging reuse) — ROCM-31508 (#12406) * [Conv] Up to 8x slower hipconv `direct`/`direct_l1`/`direct_wgrad` under ROCm 10.1's full AMDGCN LTO default; fixed with `-fno-offload-lto` — ROCM-31334 (#12328) * [Conv] gfx942 multi-node training throughput regression from backward GEMM on small-batch 1x1 NHWC shapes — ROCM-31155 (#13232) * [Conv] Invalid CK perf-config entries in the gfx1201 and gfx1100 system databases — ROCM-29026 (#12088, #12117) * [BatchNorm] JIT compilation failures on non-GFX9 targets; Welford DPP assembly restricted to GFX9 — ALMIOPEN-2838 (#12889) * Out-of-bounds write in `miopenLoadSolution` when a serialized solution carries more than three work dimensions — ROCM-31082 (#12242) * `uint64_t` typedef conflict between `miopen_cstdint.hpp` and system `<stdint.h>` on LP64 Linux (#12623) ## Technical Details Single-file change to `projects/miopen/CHANGELOG.md`. The 10.2.0 section had **two** `### Resolved Issues` headings. The two entries under the trailing one (`ConvHipConv` gfx1250 crash/LDS race, and the perf-db reload config loss) were merged into the first, and the trailing heading removed — this accounts for the 2 removed bullets in the diff. Section order is now `Resolved Issues → Added → Changed → Optimized → Removed`, matching the 3.6.1 section. `Optimized` is new to this release section. Already documented, so untouched: #12874 (reduce int32 rejection), #12000 / #13016 / #13213 (hipconv v0.3.1 / v0.3.2 / v0.4.0), #11761 (gfx803 removal). ## Notes Deliberately excluded as not externally visible (10 commits): * **hipDNN convolution forwarding (#12276)** — `MIOPEN_ENABLE_HIPDNN_WRAPPER` defaults to `OFF` (`projects/miopen/CMakeLists.txt:162`), so no shipped build compiles the forwarding code and the runtime `MIOPEN_HIPDNN_FORWARDING` flag has nothing to enable. * **Test-only**: #13289 (gfx1250 depthwise test applicability, ALMIOPEN-2841), #11610 (hipDNN parity harness), #12660 (Jenkins status post). * **Internal kernel-math cleanup, generated ISA unchanged**: #13312 (drop direct `__ocml` calls), #13190 (remove broken single-argument `__half` overloads). * **Add/revert pairs, net no change**: #12466 + #13066 (gfx1101 Navi system DBs), #12496 + #12800 (gfx1250 in the dev image arch lists). ## Risk Assessment Documentation-only change to a single Markdown file; no code, build, or runtime impact. **Risk Level: 🟢 Low** ## Test Plan None required — Markdown-only. Verified that the entry topics are not already documented elsewhere in the file, that no content was lost when merging the duplicate heading, and that every new entry maps to a commit in `2d20a8f775e..origin/develop -- projects/miopen`. ## Test Result 16 bullets added, 2 relocated — 14 net-new entries covering 15 commits. Duplicate heading removed; section ordering matches 3.6.1. Co-authored-by: Claude <noreply@anthropic.com>
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.
Summary
This PR removes the
ConvBinWinogradRxSFusedsolver entirely, along with the gfx803-specific code paths inConvBinWinograd3x3UandConvBinWinogradRxS, plus the kernel sources and find-db files that only those paths used.JIRA ID : ALMIOPEN-2580
Risk Assessment
Risk level: 1/5. This is a deletion of old code and unused kernel/find-db files - nothing new is added, and nothing changes for any other architecture.
ASIC Coverage
Everything removed here was only reachable through gfx803-specific
IsApplicable()checks in the three solver files touched (conv_bin_wino3x3U.cpp,conv_bin_winoRxS.cpp, and the deletedconv_bin_winoRxS_fused.cpp). No other architecture's solver selection, kernels, or find-db entries were touched, so passing PR CI is enough to confirm no regression elsewhere.Testing Summary
Testing Checklist
Technical Changes
ConvBinWinogradRxSFusedsolver (solver/conv_bin_winoRxS_fused.cpp) and its registration (include/miopen/fusion/solvers.hpp,fusion.cpp.ConvBinWinograd3x3U(solver/conv/conv_bin_wino3x3U.cpp) andConvBinWinogradRxS(solver/conv/conv_bin_winoRxS.cpp).conv_3x3_wheel_alpha_v3_0b*,conv_3x3_wheel_alpha_v9_2_7*, and the gfx8 includes/variants ofconv_3x3_wheel_alpha_v9_0_15*; updated the remainingv9_0_15*.sfiles to drop gfx803 references.gfx803_36.HIP.fdb.txt.bz2,gfx803_64.HIP.fdb.txt.bz2) and gfx803 branches inlegacy_ck_common.hppandCMakeLists.txt.test/gtest/cba_find2_infer.cpp,cba_infer.cpp, andnaive_timeout.cppthat only existed to work around the removed solvers.MIOPEN_DEBUG_AMD_FUSED_WINOGRADreferences fromdocs/how-to/debug-log.rstanddocs/reference/env_variables.rst.