feat(miopen): Enable channel last Gemm convolution

[sbalint98]

Motivation

This PR enables the channel last Gemm convolutions within MIOpen

Closes #11192

Technical Details

In order to achieve this new im2col col2im kernels were added. Furthermore the gemm variables were modified to fit the new layout.

Test Plan

Extended the already existing the unit test suit with relevant shapes.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

co-authored by @sikba, @Snektron

[github-actions]

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

• Add new commits
• Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

[github-actions]

This pull request has been automatically closed due to inactivity (30 days with no updates).

If you'd like to continue working on this, feel free to reopen the PR or create a new one.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[Copilot]

Pull request overview

Enables NHWC/NDHWC GEMM convolutions across forward, backward-data, and backward-weight paths.

Changes:

• Adds channel-last Im2Col/Col2Im kernels.
• Makes GEMM descriptors layout- and group-aware.
• Extends solver applicability and convolution tests.

Reviewed changes

Copilot reviewed 16 out of 16 changed files in this pull request and generated 5 comments.

Show a summary per file

File	Description
projects/miopen/test/gtest/unit_conv_solver.hpp	Declares layout-aware test cases.
projects/miopen/test/gtest/unit_conv_solver.cpp	Implements layout-aware test construction.
projects/miopen/test/gtest/unit_conv_solver_GemmWrwUniversal.cpp	Adjusts numerical tolerances.
projects/miopen/test/gtest/unit_conv_solver_GemmWrw1x1_stride1.cpp	Tests NHWC solver selection.
projects/miopen/test/gtest/unit_conv_solver_GemmBwdRest.cpp	Adds channel-last test shapes.
projects/miopen/src/solver/conv/gemm.cpp	Enables channel-last forward GEMM.
projects/miopen/src/solver/conv/gemm_wrw.cpp	Enables channel-last weight gradients.
projects/miopen/src/solver/conv/gemm_bwd.cpp	Enables channel-last data gradients.
projects/miopen/src/kernels/MIOpenIm3d2Col.cpp	Adds NDHWC Im2Col.
projects/miopen/src/kernels/MIOpenIm2d2Col.cpp	Adds optimized NHWC Im2Col variants.
projects/miopen/src/kernels/MIOpenCol2Im3d.cpp	Adds NDHWC Col2Im.
projects/miopen/src/kernels/MIOpenCol2Im2d.cpp	Adds NHWC Col2Im.
projects/miopen/src/include/miopen/util.hpp	Extends utility layout/group APIs.
projects/miopen/src/include/miopen/gemm_v2.hpp	Updates convolution GEMM interfaces.
projects/miopen/src/hip/utilocl.cpp	Configures and launches new kernels.
projects/miopen/src/gemm_v2.cpp	Adds layout-aware GEMM descriptors.

Suppressed comments (3)

projects/miopen/src/solver/conv/gemm.cpp:85

• This removes every layout restriction from the shared forward applicability check. Packed layouts such as NCHWc4 can still report packed tensors, then reach GEMM/Im2Col with IsLayoutNHWC()==false and be interpreted as ordinary NCHW. Limit the newly supported layouts to coherent default or channel-last descriptors, as the backward bases do.

    return problem.IsDirectionForward() &&
           !(gemm::IsAnyBufferBf16(xDesc, yDesc, wDesc) && !gemm::IsBf16Supported) &&
           !(gemm::IsAnyBufferFp16(xDesc, yDesc, wDesc) && !gemm::IsFp16Supported);

projects/miopen/src/hip/utilocl.cpp:245

• The new Im2Col execution paths are not covered by a channel-last non-point convolution test. Existing NHWC/NDHWC cases in the forward and WRW solver tests are point-output shapes that bypass Im2Col, while the newly expanded backward-data suite exercises Col2Im only. Add non-point NHWC/grouped-NHWC and NDHWC execution cases for GemmFwdRest or GemmWrwUniversal so these kernel variants and GEMM layouts are validated.

        if(layoutNHWC)
        {
            const bool use_aligned       = num_groups == 1;
            const bool use_channel_based = !use_aligned;

projects/miopen/src/kernels/MIOpenCol2Im2d.cpp:123

• Although col_idx is wide, the nested index expression is calculated in 32-bit arithmetic before assignment, so the 64-bit kernel variant reads wrapped locations in large NHWC column buffers. Cast the first operand to index_t to keep the full calculation wide.

            size_t col_idx =
                (((((cy * col_w + cx) * wei_h + fy) * wei_w + fx) * channels_per_group) +
                 channel_in_group) +
                size_of_group * current_group;

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[brentmaas]

LGTM AFAICT. I do have three comments, but I'd consider them more nice-to-haves rather than hard requirements.

[astrelsky]

Unfortunately I don't understand enough to leave an honest or beneficial review, nor do I have the free time to build and try out the changes at the moment.

I might have the time come this weekend though.

[JonathanLichtnerAMD]

Can we also add some NHWC FWD and WRW tests. Also claude found a number of cases where we might hit INT_MAX overflow issues. Can we add one or two tests for this as well, and maybe run these under ASAN?

We need to be careful when adding large tensor tests that might cause an overflow, since they can be very slow. In that case we would not want them running in the quick or standard categories but they could either go in the nightly category comprehensive or full.