Repository navigation
[MIOpen] Sync hipconv v0.4.0 - #13213
Merged
Merged
Conversation
#9680 made the kernel label ("direct[tile_size_k=64,...]") ConvHipConv's serialized perf config. PerfConfigBase deserializes field by field, splitting on ',', so every config loaded from the perf DB or the perf-config picker catalog came back as "direct[tile_size_k=64", matched no kernel, and fell back to the default config. When execution rebuilds the invoker from the perf DB, it then runs the default config against the workspace Find reported for the tuned config, and throws "ConvHipConv: not enough workspace for direct kernel" when the default needs more (ROCM-32185: SD2.1 fp16 batch 4 on gfx950 under MIGraphX --exhaustive-tune). Find 2.0 also re-ran the full search after every Find. Derive from PerfConfig and serialize the label as one string. Records written since #9680 store the full label and load unchanged.
Arches + files snapshot: see projects/miopen/src/hipconv/HIPCONV_VERSION Co-authored-by: Andrew Lavin <andrew.lavin@amd.com> Co-authored-by: Boris Prokopenko <boris.prokopenko@amd.com> Co-authored-by: Carsten Uphoff <carsten.uphoff@amd.com> Co-authored-by: Jin Zhou <jin.zhou@amd.com> Co-authored-by: Petar Jovanovic <petar.jovanovic@amd.com> Co-authored-by: Yingmao Lu <yingmao.lu@amd.com>
…sion A full hipconv config description selects one config only within a hipconv minor version, so a record written by another version may select a different config or none. Store records as "<config_version>:family[full description]" and match whole strings: a record from another minor version, an unstamped record, and any record in a build whose config_version is "unknown" select nothing, and the solver falls back to search or the default config. Every record written before this change is unstamped and is discarded once. The performance-log kernel name keeps the short label.
✅ 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. |
averinevg
reviewed
Oct 7, 2026
…ipConv.cpp unit_conv_solver_*.cpp files hold RunTest smoke and full cases. Move SerializeRoundTrip and VersionStamp next to the other per-solver perf-config tests, as CPU_PerfConfig_ConvHipConv_NONE and GPU_PerfConfig_ConvHipConv_FP16.
… before enumerating IsValid enumerated and ranked every config for the problem before comparing a record against them. The perf-config picker from #9680 calls it once per candidate in its bucket (up to 168 on gfx950), and no catalog entry carries a version prefix, so a perf-db miss paid one enumeration per candidate before falling back to config 0. Check the prefix first.
clang-tidy (modernize-use-starts-ends-with) rejects substr(0, n) == prefix.
averinevg
reviewed
Oct 8, 2026
averinevg
left a comment
averinevg
left a comment
Contributor
There was a problem hiding this comment.
Thanks for addressing my comments! A few more notes inline.
Closed
4 of 5 tasks
VersionStamp checked only config 0's record, which is also the config a record that matches nothing falls back to, and only that IsValid accepted it. It now runs over every config of a dense 3x3 fprop layer (four on gfx950 and gfx1250): each record must resolve to its own index, and the record under another version stamp, unknown, or no stamp must be rejected.
hipconv v0.4.0 adds a gfx1250 direct_wgrad kernel, which serves the dense and vertical 3x1 Wrw shapes, so SmokeDenseWrw and the Wrw instances of SmokeDense and SmokeDenseVertical no longer need to be gfx950-only.
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.
JIRA ID : ALMIOPEN-2361
Takes the vendored
src/hipconvsnapshot from v0.3.2 to v0.4.0 and fixes ConvHipConv's perf-db round trip, which broke SD2.1 fp16 batch 4 on gfx950 in the 10-02 nightly (ROCM-32185).ConvHipConv perf-db records
#9680 made ConvHipConv's perf config the hipconv kernel label, e.g.
direct[tile_size_k=256,tile_size_n=1,...].PerfConfigBasedeserializes field by field and splits on,, so a stored label reloaded asdirect[tile_size_k=256, matched no kernel, and the solver fell back to config 0. Under MIGraphX's--exhaustive-tune, execution then built config 0 against the workspace Find had reported for the tuned config and threwConvHipConv: not enough workspace for direct kernel(conv_hipconv.cpp:702).fix(miopen): serialize the ConvHipConv perf config whole:PerformanceConfigConvHipConvderives fromPerfConfigand stores the label as one string.fix(miopen): key ConvHipConv perf-db records to the hipconv minor version: a record is<hipconv config_version>:family[full description], e.g.v0.4:direct[tile_size_k=64,...,direction=fprop,type=fp16,...]. hipconv guarantees that a full description selects the same config only within a minor version, so a record from another minor version, an unstamped record, and any record in a build whoseconfig_version()isunknownselect nothing. The solver then falls back as for any invalid record (search, the AI picker, or hipconv's top-ranked config). Every ConvHipConv record written before this change is unstamped and is discarded once. The performance-log kernel name keeps the short label.#9680's perf-config picker catalog (
lgbm_pcfg.bin) holds 166 ConvHipConv labels in the old form. They never resolved, since every one contains a comma, and they will not match stamped records; they need regenerating against v0.4.0.hipconv v0.4.0
Per-version detail is in the shipped
src/hipconv/CHANGELOG.md.describe_config(kernel, full)renders every field, which is what the stamp stores.version()andconfig_version()report the snapshot's version from the shippedVERSIONfile.direct_wgrad: dense backward-weights, fp16/bf16 with an fp32 dW, 2x2, 3x3, 4x4, 5x5, 3x1 and 1x3 filters, stride 1 (hipconv Add thread include #212, [next-cmake] Updates to accommodate rocroller #297).ConvParams::is_valid()rejects a layer whose dilated filter is wider than its padded input (hipconv [Issue]: failure to handle missing python module joblib #288).Data types are unchanged: fp16, bf16 and TF32.
projects/miopen/CHANGELOG.mdgains the gfx1250 dense backward-weights kernel and the perf-db fix under MIOpen 3.6.2.Downstream patches
None. On develop,
projects/miopen/src/hipconvwas byte-identical to hipconv v0.3.2.Validation
test_unit_conv_solver_ConvHipConvon gfx950 (MI355X), MIOpen built with clang 24 (ROCm 10.1) from this branch: 377 pass and 24 skip (the gfx125X-only suites), the same in two repeats. On MI455 A0 (gfx1250), built forgfx1250;gfx1250-strict: 210 pass and 191 skip (the gfx950-only suites), the same in two runs. New:CPU_UnitTestConvSolverConvHipConvPerfConfig_NONE.SerializeRoundTripandGPU_UnitTestConvSolverConvHipConvPerfConfig_FP16.VersionStamp, which pass on both.conv_hipconv.cpp:702. With this branch, every layer passes on a fresh user perf db (one search each), again reloading those records (no search, no warning), and with the records restampedv0.3(each rejected once with the "Invalid config" warning, re-searched, then passing).A find-db record also caches a workspace size and carries no hipconv version; that is unchanged here.