Repository navigation
feat(miopen): gfx950 7x7 depthwise fwd + backward-data conv (fp16/bf16) - #9168
Merged
yassinsolim merged 2 commits intoJul 22, 2026
Merged
Conversation
✅ 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. |
yassinsolim
force-pushed
the
users/yassinsolim/miopen-gfx950-depthwise-7x7
branch
3 times, most recently
from
July 7, 2026 21:37
525e4d6 to
642e9f7
Compare
yassinsolim
changed the title
yassinsolim
changed the title
yassinsolim
force-pushed
the
users/yassinsolim/miopen-gfx950-depthwise-7x7
branch
from
July 7, 2026 23:52
642e9f7 to
a38e6f5
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
Adds fast CK-backed depthwise convolution coverage for ConvNeXt-style 7x7 (stride-1, pad-3) on gfx950 by expanding the CK instance set and introducing a new backward-data solver that reuses the existing CK forward depthwise kernel with a compile-time 180° filter flip. This extends auto-selection and tuning to both fp16 and bf16 without introducing new CK kernels.
Changes:
- Extend CK depthwise forward support to include 7x7 stride-1 pad-3 instances (and broaden dtype handling to fp16 + bf16).
- Add a new depthwise backward-data solver (
ConvDepthwiseBwdData2D, solver id 186) implemented via forward-kernel reuse withFlipFilter. - Wire up CK loader/interface plumbing for the new solver and add/extend unit tests for forward/bwd-data depthwise cases (including bf16).
Reviewed changes
Copilot reviewed 15 out of 15 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| projects/miopen/test/gtest/unit_conv_solver_ConvDepthwiseBwdData2D_test.cpp | Adds unit coverage for the new depthwise backward-data solver (fp16/bf16, multiple filter sizes/tiles). |
| projects/miopen/test/gtest/unit_conv_solver_ConvCkGroupedConvFwd.cpp | Extends depthwise forward unit cases to include 7x7 convnext shapes and adds bf16 test instantiations. |
| projects/miopen/test/gtest/CMakeLists.txt | Skips the new CK depthwise bwd-data unit test under ASAN packaging (consistent with existing CK test skips). |
| projects/miopen/src/solver/conv/conv_ck_grouped_conv_fwd.cpp | Enables bf16 for depthwise fwd perf config/applicability and introduces the new ConvDepthwiseBwdData2D solver implementation and guards. |
| projects/miopen/src/solver/ck_impl_lib_loader.cpp | Binds CK dynamic library symbols for the new DepthwiseBwdData solver type. |
| projects/miopen/src/solver.cpp | Registers the new solver in the global solver registry. |
| projects/miopen/src/mlo_dir_conv.cpp | Adds the new depthwise bwd-data solver to the direct solver list. |
| projects/miopen/src/include/miopen/solver/ck_impl_lib_loader.hpp | Extends CKSolverType enum with DepthwiseBwdData and updates Count. |
| projects/miopen/src/include/miopen/solver/ck_impl_interface.hpp | Adds the CK interface API surface for depthwise backward-data (fill/applicable/args/workspace/solution). |
| projects/miopen/src/include/miopen/conv/solvers.hpp | Declares the new tunable solver and performance config types for depthwise backward-data. |
| projects/miopen/src/include/miopen/conv/device_grouped_conv_fwd.hpp | Adds bf16 inner_product support and the FlipFilter compile-time flag plumbing into the gridwise/device wrappers and kernel id string. |
| projects/miopen/src/fin/fin_interface.cpp | Maps solver id 186 to ConvDepthwiseBwdData2D for FIN dispatch. |
| projects/miopen/src/ck_impl/CMakeLists.txt | Adds the new CK depthwise backward-data implementation file to the CK-impl build. |
| projects/miopen/src/ck_impl/ck_depthwise_fwd_impl.cpp | Templates the forward factory by dtype, adds 7x7 instances, and strengthens instance-shape matching/selection logic. |
| projects/miopen/src/ck_impl/ck_depthwise_bwd_data_impl.cpp | Implements CK depthwise backward-data via forward-kernel reuse (FlipFilter=true) for fp16/bf16 with shape guards and solution building. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
yassinsolim
force-pushed
the
users/yassinsolim/miopen-gfx950-depthwise-7x7
branch
from
July 13, 2026 18:10
c681238 to
353a367
Compare
JonathanLichtnerAMD
approved these changes
Jul 14, 2026
Contributor
There was a problem hiding this comment.
This LGTM. Is there a separate activity to update the system DBs?
Also, I don't think this is specific to gfx950, but also applies to other CDNA arches (although it might not be quite as fast). Should we test and measure some of these shapes on gfx942?
We also probably will want something similar for gfx1250... + @EnricoDeg for awareness
Adds FilterSize=7 stride-1 pad-3 CK depthwise instances (forward) and a new ConvDepthwiseBwdData2D solver (backward-data via the forward kernel with a 180-degree-rotated filter), so convnext 7x7 depthwise no longer falls back to naive/im2col on gfx950. Also adds a 3x3 stride-1 tile-112 instance, bf16 support (dtype-templated factories + gfx950 hardware bf16 dot), and hardens the dgrad solver (IsArgsSupported guard + InstanceShapeMatchesProblem). Reuses the existing CK depthwise gridwise kernel; no new CK kernels. Adds unit tests: a new ConvDepthwiseBwdData2D solver test (fp16+bf16, 7x7/3x3/5x5 stride-1) and 7x7 + bf16 cases for the existing ConvDepthwiseFwd2D test. Validated on MI355X: 40/40 correctness (fwd+dgrad x fp16+bf16 x 3x3/5x5/7x7), auto-selected by the default find, no regressions; convnext_large fp16 training +50%.
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
yassinsolim
force-pushed
the
users/yassinsolim/miopen-gfx950-depthwise-7x7
branch
from
July 21, 2026 20:01
353a367 to
53ff08c
Compare
yassinsolim
enabled auto-merge (squash)
July 21, 2026 22:05
1 task done
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Aug 13, 2026
JIRA ID: ALMIOPEN-2364 ## Motivation ROCm 7.15 has been renamed to ROCm 10.0, and MIOpen 3.6.0 is the corresponding release. `projects/miopen/CHANGELOG.md` needs a new version section documenting what changed for downstream users, release-notes tooling, and packagers between MIOpen 3.5.2 (ROCm 7.14.0) and MIOpen 3.6.0 (ROCm 10.0.0), consistent with the changelog's existing per-release format. ## Technical Details - Inserted a new `## MIOpen 3.6.0 for ROCm 10.0.0` section at the top of `projects/miopen/CHANGELOG.md`, directly above the existing `## MIOpen 3.5.2 for ROCm 7.14.0` section, using the same `### Added` / `### Changed` / `### Resolved Issues` / `### Removed` subheading structure as every prior release section in the file. - Entries were compiled by diffing `origin/release/therock-10.0` against `origin/release/therock-7.14` for `projects/miopen` and hand-selecting changelog-worthy commits, excluding: test-only commits, CI-only commits, build/version-bump commits, doc-only commits, reverted commit/revert pairs, and commits already surfaced in the 3.5.2/7.14.0 section via cherry-pick. - Resulting entries added: - **Added:** gfx950 (MI350X/MI355X) 7x7 depthwise forward and backward-data convolution support, fp16/bf16 (#9168) — fixes a slow fallback-to-naive-kernel regression in ConvNeXt-style depthwise convolutions. - **Changed:** - Restored gfx12x support in the Winograd Rage solver (#6439), recovering performance lost when earlier gfx12 support was reverted. - Refreshed the gfx1100, gfx1102, and gfx1201 (Navi) SystemDBs (#9529) with updated tuned find/perf-database entries. - Refreshed the gfx950 SystemDB (#9111) with additional tuned entries. - **Resolved Issues:** - Fixed an off-by-stride indexing bug in BatchNorm backward `CalcStats` mean/variance remainder loop causing a ~1% systematic bias in NCHW batchnorm backward results (#9031). - Fixed a tensor-op integer overflow on gfx1250 for large allocations that could cause memory access faults (#9049). - Fixed naive convolution solvers silently failing to launch (and leaving a stale HIP error visible after Find returned success) when a global work size reached 2^32 or more work-items; such launches are now rejected up front (#10188). - **Removed:** the OpenCL (OCL) backend (#8433); MIOpen now supports the HIP backend only. - No code, build, or test files were touched — this is a pure documentation addition; the underlying functional changes were already merged to `release/therock-10.0` in their own PRs (referenced above) and are only being documented here. ## Risk Assessment Single file (`projects/miopen/CHANGELOG.md`), 17 lines added, 0 removed, no code or build-system changes. **Risk Level: 🟢 Low** ### Impacted Components - `projects/miopen/CHANGELOG.md` only — no source, kernel, solver, build, or test files are touched. ### Potential Side Effects - None functionally. The only observable effect is the new changelog text itself; no build artifacts, APIs, or runtime behavior change. - Cosmetic risk only: if any release-notes/packaging tooling parses this file programmatically, it should be checked for compatibility with the new section, though the added section strictly follows the existing heading/bullet conventions used by all prior sections in the file. ### Mitigation Steps - Followed the exact existing section format (`##` version header, `###` category subheadings, `* ` bullet entries, `[Component]` prefix tags) used by the adjacent 3.5.2/7.14.0 and 3.5.1/7.13.0 sections, minimizing formatting risk. - Each entry was cross-referenced against a specific merged PR number (#9168, #6439, #9529, #9111, #9031, #9049, #10188, #8433) rather than paraphrased from memory, so reviewers can verify each line against its source commit. ## Test Plan No functional tests apply — this is a documentation-only change to a Markdown changelog file. Verification consisted of manually reviewing the diff against the existing file's structure and confirming the new section's heading hierarchy and bullet formatting match every other version section in `CHANGELOG.md`. ## Test Result - `git diff origin/develop...HEAD --stat` confirms the change is scoped to exactly one file (`projects/miopen/CHANGELOG.md`, +17/-0). - Visual/manual review confirms the new `## MIOpen 3.6.0 for ROCm 10.0.0` section renders as valid Markdown consistent with surrounding sections. No build or CI run is applicable to a docs-only change. --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
1 task done
yassinsolim
added a commit
that referenced
this pull request
Oct 3, 2026
…stance (#11271) JIRA ID: ALMIOPEN-2423 ## Motivation The CK depthwise backward-data instance table added in #9168 covers three filter families, and the 3x3 stride-1 pad-1 family is the only one missing its tile-28 entry: | family | tiles present | |---|---| | 7x7 stride-1 pad-3 | 7, 14, **28**, 56 | | **3x3 stride-1 pad-1** | 7, 14, **—**, 56 | | 5x5 stride-1 pad-2 | 7, 14, **28** | The gap is recorded in the repository itself: the coverage comment in `unit_conv_solver_ConvDepthwiseBwdData2D_test.cpp` reads `3x3 pad-1 (tiles 7/14/56)` while listing 28 for both sibling families. 3x3 is the most common depthwise filter in the efficientnet/mobilenet family the surrounding table comment names. A dgrad problem landing on 3x3 stride-1 pad-1 at BlockTileSize 28 matches no instance, so `ConvDepthwiseBwdData2D` is not applicable and the problem falls back to the naive backward solver. **Scope, stated plainly.** This is a symmetry fix to an existing table, not a strategy for depthwise dispatch coverage generally. ALMIOPEN-2423 argues - with a fleet coverage curve - that enumerating instances does not converge against arbitrary production tile sizes, and recommends a generic-spatial fallback instead. I agree with that ticket and am not proposing enumeration as the fix. This one instance is worth landing on its own terms only because it completes a progression the table already defines, using parameters that already ship; it is independent of, and does not compete with, the generic-fallback work or the `naive-bailout` direction discussed on ROCM-28401. ## Technical Details One instance, in `projects/miopen/src/ck_impl/ck_depthwise_bwd_data_impl.cpp`: ```cpp + DwBwd<DType, 3, 1, 28, 32, 4, 4, 4, 4>, ``` No new kernel, no new solver, no new source file. `DwBwd` is the existing alias over `DeviceGroupedConvFwd` used by all ten current entries, and the parameters (`NBatch=32, SubH=4, SubW=4, VecIn=4, VecOut=4`) are copied verbatim from the two sibling families that already ship tile 28 - `DwBwd<DType, 7, 3, 28, 32, 4, 4, 4, 4>` and `DwBwd<DType, 5, 2, 28, 32, 4, 4, 4, 4>`. Only `FilterSize` and `Pad` differ, and both values already appear in this table for tiles 7, 14 and 56. Inserted in ascending tile order within its family, so the block reads 7 / 14 / 28 / 56 like the 7x7 block above it. Second commit extends `test/gtest/unit_conv_solver_ConvDepthwiseBwdData2D_test.cpp` with the matching dgrad case and corrects that file's coverage comment to `3x3 pad-1 (tiles 7/14/28/56)`. Extending the existing case list rather than adding a new test file, since #9168 created this file for exactly these instances. Arch-general, consistent with #9168: no gfx950-only gate, builds per-arch. Verified absent from `develop` (HEAD `20aff86` at the time of writing); no open pull request touches either file. ## Test Plan - Static check that the instance is absent upstream and that no in-flight PR collides. - Differential syntax check on gfx950 in a ROCm 7.14 container: compile the modified translation unit with and without the added line and compare diagnostics. - Add the corresponding case to the existing `ConvDepthwiseBwdData2D` gtest so the new instance is exercised by CI alongside the ones already covered. ## Test Result - **Absent upstream:** confirmed against a fresh `develop` clone - 10 instances present, none matching 3x3 pad-1 tile 28. - **Differential compile:** identical diagnostic set with and without the line, so the addition introduces no new diagnostics. A fully clean local compile needs a complete CMake configure, which I could not complete in my environment, so **CI is the authoritative build gate here** rather than a local pass I am claiming. - **gtest:** case added and covered by the existing suite; not executed locally for the same configure reason. CI runs it. - **Not claimed:** no standalone performance measurement is attributed to this single instance. Fleet measurement of depthwise dispatch coverage on MI355X is reported in ALMIOPEN-2423; this change is one entry inside that picture, and I would rather point at the ticket than imply a per-instance delta I have not isolated. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: yassinsolim <yassinsolim@users.noreply.github.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: ROCM-27769
Motivation
On gfx950 (MI350/MI355), convnext depthwise convolutions fall back to slow kernels in both
directions, which is the milabench
convnext_largeregression. The CK depthwise instancefactory only had FilterSize 3 and 5, so the 7x7 depthwise matched no instance and used the
naive forward solver, and there was no CK depthwise backward-data path at all, so dgrad always
used the naive backward solver (the single most expensive op). A complete/latest CK package
does not fix this on its own: it just moves 7x7 depthwise from naive to im2col-GEMM, which is
about as slow (measured ~2.8 ms fwd / ~7.4 ms dgrad at 192@56 on ROCm 7.14, vs ~0.11 ms with
this change). The goal of this PR is to make convnext 7x7 depthwise fast and auto-selected on
gfx950, in fp16 and bf16, with no new CK kernels.
Technical Details
Reuses the existing CK depthwise gridwise kernel. This PR adds instances plus a thin solver:
the convnext depthwise stages), plus a 3x3 stride-1 tile-112 instance, so
ConvDepthwiseFwd2Dbecomes applicable for 7x7.
ConvDepthwiseBwdData2D, id 186): for stride-1 symmetric filters,dx = conv_fwd(dy, rot180(w)). Implemented with a compile-timeFlipFilterflag on the forwardgridwise kernel (reads the filter 180-degree rotated, no new kernel, no workspace, no extra
launch) plus the solver and its registration.
gfx950 hardware bf16 dot (
v_dot2_f32_bf16via__builtin_amdgcn_fdot2_f32_bf16), guarded to__gfx950__with a convert-to-float fallback on other arches.GetSolutionguards onIsArgsSupportedbefore calling the CK lib(a stale SystemDB
kernel_idfails gracefully instead of aborting), andInstanceShapeMatchesProblemis applied in the factory lookups so a wrong-filter instance cannot be selected at a shared tile.
The change is arch-general (no gfx950-only gate on the instances); it builds per-arch and would
also help gfx942 once built there. Complementary to #8653 (ck-depthwise fixes: it adds correctness
guards for the existing depthwise path); this PR is independent of it and applies cleanly on
develop.Files: 12 changed, +967/-124 (1 new file,
ck_impl/ck_depthwise_bwd_data_impl.cpp).Test Plan
Measured on MI355X (gfx950), fp16 and bf16, N=128, via MIOpenDriver (forced solver, verified vs
GPU reference) and PyTorch end-to-end:
instance shares the tile.
Test Result
ConvDepthwiseFwd2D/ConvDepthwiseBwdData2Dfor fp16 andbf16 and verifies, including the shared-tile 3x3@56 case.
resolves, all verify. Inapplicable shapes report "no suitable algorithm" and fall back gracefully.
4.39 -> 0.11 ms).
efficientnet_b4 fp16 +20%.
Submission Checklist