Repository navigation
[MIOpen] Sync hipconv v0.3.1 - #12000
Conversation
Arches + files snapshot: see projects/miopen/src/hipconv/HIPCONV_VERSION ConvHipConv takes hipconv::ConvParams in place of Conv2dParams and accepts a 3D convolution whose depth is unconvolved, folding depth into the batch for NHWC so every call site sees 2D params. unit_conv_solver_ConvHipConv gains SmokeConv3d cases at depth 2 covering that path. Give the wgrad kernel its workspace. No wgrad kernel asked for scratch before, so GetWorkspaceSize reserved only the fp32 buffer that fp16/bf16 dw is staged through, and both launches passed nullptr for it. v0.3.0 adds depthwise_wgrad_hankel, which writes the partials of a split reduction there, so it failed at launch with hipErrorInvalidArgument on any layer it split. Find caught the failure and picked another solver, so results stayed correct and ConvHipConv silently stopped being used. The failure did not stay local either: the next unrelated launch in the process also reported invalid argument. Size the workspace as the staging region plus the largest scratch any config asks for, and pass the region past the staging buffer to the kernel. The fp32 path needs it too: it stages nothing, but a split kernel still needs scratch. 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>
✅ 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. |
|
The job builds so nothing in this PR is compiled by the job. #11998 hits the identical error. The develop branch passes only because Component CI skips MIOpen when no MIOpen path changed. Needs a fix in rocjitsu. |
A fix for that has been merged #11955 |
|
The PR description talks about TF32 but I don't see it hooked into the TF32 support and controls in MIOpen and I don't see any tests that exercise it. In MIOpen it is an option to either use true FP32 or to use TF32. |
ToHipconvParams carries extents only, so a non-packed tensor reached the kernel as if it were packed. Gate IsApplicable on HasNonPackedTensors, as the other conv solvers do. Instantiate the ConvHipConv fwd/bwd/wrw unit tests for tf32 over the existing 2D and 3D smoke shapes. IsApplicable has admitted fp32 with tf32 compute since v0.3.0, but nothing exercised that path.
hipconv serves gfx1250 through its cdna5 backend, but the tests were gated to gfx950 alone, so the cases would not run on a CDNA5 device even once one is available. Inert today: CI builds gfx125X on every PR and executes nothing on it, so the tests skip there until a runner exists. gfx950 results are unchanged.
@BradPepersAMD TF32 tests added: The TF32 wiring itself landed in #9746, not here — If you had a different control in mind, say which and I'll wire it up. |
WrwStagingBytes omitted the filter depth, so a 3D wgrad with non-unit kd under-allocated and the kernel wrote past its workspace. SmokeConv3dDepth covers it; the existing 3D cases leave depth unconvolved and cannot.
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>
…1334) (#12328) JIRA ID : ROCM-31334 ROCm 10.1 is the first release whose clang defaults the AMDGCN LTO mode to full. Under that pipeline the register allocator mis-ranks live ranges and spills into the software-pipelined inner loops of hipconv's `direct`, `direct_l1` and `direct_wgrad` kernels. The code is correct, so CI passes, but performance degrades by up to 8x. `depthwise` and `grouped` are unaffected. On gfx950, `n256 h32 w32 c512 k512 3x3 wgrad bf16`, same compiler, the flag as the only variable, three interleaved repetitions within 0.5%: | | time | TFLOP/s | VGPR | scratch | |---|---|---|---|---| | with flag | 0.737 ms | 1677 | 253 | 0 B | | without | 5.315 ms | 233 | 256 | 356 B | ## Change `-fno-offload-lto` on the `hipconv` static library and on each `hipconv_arch_<name>` object library. `-DHIPCONV_OFFLOAD_LTO=ON` restores the compiler default. Set with `target_compile_options` rather than appended to `CMAKE_HIP_FLAGS`, because a consumer may compile these sources as CXX — MIOpen does — and HIP flags never reach them then. ## Verification Built through MIOpen in a ROCm 10.1 container (clang 24.0.0git `d6f6cb6`) both ways, with the flag as the only variable. Across all 1,543 cdna4 kernels: | | with flag | without | |---|---|---| | total scratch | 25,548 B | 82,936 B | | kernels at 256 VGPRs | 136 | 180 | The flag removes scratch from 175 kernels and adds a small amount to 7. No kernel in the three affected families is worse. Functional results are unchanged — this is invisible to every existing test. ## Notes Upstream report: ROCm/llvm-project#4434. The underlying allocator defect is unfixed; fixes llvm/llvm-project#224031, llvm/llvm-project#224144 and llvm/llvm-project#224467 are all open. This touches two files that #12000 also rewrites, so whichever lands second needs a trivial conflict resolution. ROCm 10.1's release branch carries hipconv v0.2.7 and is affected the same way; a cherry-pick request will follow separately.
JonathanLichtnerAMD
left a comment
There was a problem hiding this comment.
LGTM, but can we add a one-line comment in the miopen changelog, under ### Added in the 3.6.2 section, something like
- [Conv] Extended the
ConvHipConvsolver to 3D convolutions on the NHWC/NDHWC
layout, and added gfx950 depthwise backward-weights (fp16/bf16) and gfx1250
depthwise and tf32 kernels (hipconv v0.3.1).
There are also some merge conflicts to resolve, and we will need to get @mattwill-amd to approve the doc changes.
# Conflicts: # projects/miopen/src/hipconv/CMakeLists.txt # projects/miopen/test/gtest/unit_conv_solver_ConvHipConv.cpp
|
@amd-alavin MIOpen TF32 support is currently gated to only gfx942 and gfx950, regardless of individual kernel support for this data type, so the TF32 support added for gfx1250 in this PR won't be accessible through MIOpen until we make some changes. I'm opening a task for the MIOpen team to investigate and fix this in a way that doesn't gate datatypes by GPU for all solvers and we can have more granular solver-specific checks. |
|
Thanks @JonathanLichtnerAMD.
@mattwill-amd, could you review the doc changes? There are two files: |
TheRock Submodule Bump ActivityNewest first
|
JIRA ID : ALMIOPEN-2361 Takes the vendored `src/hipconv` snapshot from v0.3.1 to v0.3.2. Per-version detail is in the shipped `src/hipconv/CHANGELOG.md`. ## Data types ConvHipConv supports fp16, bf16 and TF32, and not fp32. This is unchanged from v0.2.7 and v0.3.1. hipconv emulates TF32 with three bf16 MFMAs on a (big, small) bf16 split of the stored fp32, so it needs no TF32 hardware, and its gfx950 and gfx1250 kernels both support it. MIOpen decides when a problem is TF32: `MIOPEN_CONVOLUTION_ATTRIB_MATH_TYPE` set to `miopenMathDefault` or `MIOPEN_TF32_OVERRIDE`, on a device `IsTF32Supported` admits. That list is gfx942 and gfx95x today, so through MIOpen, TF32 reaches hipconv on gfx950 and not yet on gfx1250. Additional TF32 support on gfx950 since v0.3.1: - `direct_l1` forward and backward-data, 2x2 to 5x5 filters (hipconv #226); - `depthwise_1d_toeplitz` forward and backward-data (hipconv #223); - `direct_wgrad` backward-weights at 2x2, 3x1, 3x3, 4x4 and 5x5 (hipconv #234). ## Kernels - gfx950 `patch_embed`: patch-embedding convolutions (stride equal to the filter, no padding), forward, backward-data and backward-weights, fp16/bf16. - gfx1250 `depthwise_wgrad_hankel`: depthwise backward-weights, fp16/bf16, 3x3 to 11x11 odd square, stride 1 or 2. - gfx1250 `direct`: fixes a crash when groups x batch tiles exceeded the 65535 grid z limit (ALMIOPEN-2624), and an LDS race between a TDM load and another wave's reads; the race fix lifts the odd-channel backward-data restriction v0.3.1 added. `direct` no longer claims depthwise layers (one channel per group). - gfx950 `depthwise_wgrad_hankel` sums its partitions in a fold kernel instead of with atomics: geometric mean +88% over fifteen shapes (hipconv #227). ## Build `src/CMakeLists.txt` renames `HIPCONV_ENABLE_POINTWISE` to `HIPCONV_ENABLE_EXPLICIT_GEMM`, the name hipconv v0.3.2 gives the option, which now enables explicit-GEMM patch embedding as well as pointwise. CMake ignores the old name without a warning. Both explicit-GEMM paths stay off, since the new option also defaults to OFF. ## Downstream patches On develop, `projects/miopen/src/hipconv` differed from hipconv v0.3.1 only by #12406, which patched `src/arch/cdna4/direct/direct_wgrad/epilogue.h`. hipconv v0.3.2 contains the same fix (hipconv `e1777ac9`), so this sync keeps the file byte-identical. ## Validation - hipconv's internal test layers (549 convolutions, each in all three directions) through MIOpen with ConvHipConv forced, on gfx950 (MI355X), with MIOpen built by clang 24 (ROCm 10.1) against this snapshot: every case ConvHipConv accepted passed, 1476 in fp16, 1476 in bf16 and 1283 in TF32. 18 of the cases it rejected (12 in TF32) are ones hipconv supports on its own. The cause, noted in #12000, is that `GetAdjustedConvolutionStride{H,W}` hands hipconv a stride of 1 on an axis whose output extent is 1. They are 6 grouped stride-2 layers with a degenerate spatial extent, in forward and backward-data, and 2 patch-embedding layers with a 1x1 output, in all three directions. - hipconv's internal test suite on MI455 A0 (gfx1250): 64/64, including about 14,000 cases that run every kernel config against the same test layers. On gfx950, 65/65. Both ran at hipconv `f58e0771`, from which v0.3.2 differs in formatting, its changelog, host warning flags and one constexpr expression. clang-tidy over the consumer TU is clean with AMD LLVM 24. ## ConvHipConv unit tests `unit_conv_solver_ConvHipConv.cpp` gains NHWC tests that cover every hipconv kernel available through ConvHipConv. The layers come from hipconv's internal tests. | List | gfx950 | gfx1250 | |---|---|---| | `GetConvDenseTestCases` | `direct_l1`, `direct` (Fwd, Bwd); `direct_wgrad` (Wrw); tf32 on `direct_l1`, `direct_wgrad` | `direct` (Fwd, Bwd) | | `GetConvDenseVerticalTestCases` | `direct_l1`, `direct_wgrad` | — | | `GetConvPatchEmbedTestCases` | `patch_embed` (Fwd, Bwd, Wrw) | — | | `GetConvOddChannelTestCases` | tf32 only: `direct_l1`, `direct_wgrad` | `direct` (Fwd, Bwd) | | `GetConvPointwiseTestCases` | — | `direct` 1×1 (Fwd, Bwd) | | `GetConvGroupedTestCases` | `grouped`, `grouped_wgrad` at 4, 8, 16 and 32 channels per group; fp16, bf16, tf32 | `grouped_multi_g`, `grouped_multi_g_wgrad` | | `GetConvGroupedStride2TestCases` | `grouped` at 4, 16 and 32 channels per group (Fwd, Bwd) | `grouped_multi_g` (Fwd, Bwd) | | `GetConvGrouped8cStride2TestCases` | `grouped` at 8 channels per group (Fwd) | `grouped_multi_g` (Fwd, Bwd) | | `GetConvDepthwiseTestCases` | `depthwise_1d_toeplitz` (Fwd, Bwd; tf32 too), `depthwise_wgrad_hankel` (Wrw) | the same, fp16 and bf16 | Each list is instantiated only on the architectures and directions whose kernels support all of its layers, since the harness fails an inapplicable case. The new lists are NHWC only: the NCHW staging does not depend on the kernel, and the existing smoke lists cover it. **`Tunable(0)`.** `GenericSearch` shuffles every valid perf config and benchmarks the first `Tunable(N)`, so `Tunable(5)` verifies the fastest of five random configs. On a layer several kernels support, which kernel that is varies from run to run. The grouped stride-1 layers are such layers (the direct kernels support them too), as are the tf32 depthwise layers. Their suites use `Tunable(0)`: the sample is empty, so the search runs only the default config, which is hipconv's top-ranked, and the grouped or depthwise kernel is the one verified. Every other list keeps `Tunable(5)`, since one kernel supports each of its layers and the sampling still exercises its non-default configs. **Testing.** On gfx950 (MI355X): 399 tests, 375 pass and 24 skip (the gfx125X-only suites), the same in each of five repeats. With `MIOPEN_LOG_LEVEL=5`, each new case's final `GetSolution` names the kernel it was meant to probe. On MI455 A0 (gfx1250): 208 pass and 191 skip (the gfx950-only suites), each case on its intended kernel; CI skips the gfx1250 cases, and hipconv's own suite also passed on MI455 A0 (Validation above). --------- 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>
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>
JIRA ID : ALMIOPEN-2361
Takes the vendored
src/hipconvsnapshot from v0.2.8 to v0.3.1, and brings theConvHipConv solver up with it. Per-version detail is in the shipped
src/hipconv/CHANGELOG.md, which is new in this snapshot.Kernels
v0.3.0 adds
depthwise_wgrad_hankel(gfx950 depthwise conv2d wgrad, fp16/bf16,3x3 to 11x11 odd square, stride 1 or 2) and
depthwise_1d_toeplitzfor gfx1250(fprop, dgrad, fp16/bf16/tf32). It also adds tf32 across the gfx1250 kernels,
which MIOpen cannot select yet because
IsTF32Supportedadmits only gfx942 andgfx95x; conv3d descriptors with the unfolding that reduces a degenerate 3D convolution to
conv2d, and removes 1x1 from the gfx950
directfilter set.v0.3.1 adds no kernel: a gfx1250 dgrad dispatch guard,
-fno-offload-ltoonhipconv's device targets, and the removal of a
static_assertdiscussed below.Solver
IsSupportedProblemnow admits 3D on the NHWC/NDHWC path, and rejects non-packedtensors for every layout rather than only NCHW. 3D stays out of the NCHW staging
path:
MakeTransposePlanreads four extents positionally, so a 5D tensor wouldsilently lose one.
TransposeSolutionDefault2Ndhwcis what a channels-first 3Dpath would be built on, if wanted later.
The merge of #11845 drops
WrwStagingBytes/WrwScratchOffset/WrwScratch, andwith them the fix in 652084e --
GetWgradCastSizesizes fromGetWeights().GetElementSize(), which counts depth inherently. That commit isstill in this branch's history; its subject describes a function the merged tree
no longer has.
Downstream patches this reconciles
#12085 added
gfx1250-strictto the vendoredsrc/arch/cdna5/CMakeLists.txt.The publish rsyncs
src/with--delete, so the next sync would have revertedit. It is now upstream in hipconv, along with the
GPU_TARGETSfilter change thatretains the
-strictsuffix -- the filter previously dropped it as a non-offloadtarget, so the registry row would have had no device code behind it. The vendored
file is byte-identical to what #12085 produced.
#11989 removed the
rounds_are_rungs()static_assertfromlds_layout.hfor GCC 14+ headers. hipconv removed the same assert upstream and moved the check
to a test that does not ship; this snapshot also removes the now-dead function the
downstream patch left behind. Nothing referenced it.
#12328 and #12406, merged from develop after this snapshot was taken.
#12328 set
-fno-offload-ltoon the vendored device targets; v0.3.1 carries thesame change upstream, so the merge keeps the snapshot text. #12406 fixed
direct_wgrad LDS ordering in
epilogue.h(hipconve1777ac9, after v0.3.1); itmerges cleanly over this snapshot, and
HIPCONV_VERSIONkeeps its out-of-snapshotnote until a sync at or after
e1777ac9replaces it.Validation
gfx950:
(ROCm 10.1), including 60 conv3d cases across fwd/bwd/wrw x fp16/bf16/tf32 and
[MIOpen] Order hipconv direct_wgrad's staging reuse at waves_q == 1 #12406's 12 dense wgrad cases.
454 pass, 0 fail, 9 no-match. Three no-matches are 1x1, which is pointwise and
off by default in a MIOpen build. The other six are grouped stride-2 layers with
a degenerate spatial extent, where
GetAdjustedConvolutionStrideWrewrites thestride to 1 on the axis whose output extent is 1 and hands hipconv a mixed
(2,1) stride that its uniform-stride config table does not cover. Pre-existing,
unrelated to this sync.
corpus.
gfx1250-strict was verified on ROCm 10.1, where the target is supported: MIOpen
configures with it,
hipconv_arch_cdna5compiles for bothgfx1250andgfx1250-strict, and the generated registry carries both rows.