[aledudek]

Motivation

Add cardinality-grouped (groups>1) forward convolution coverage to the implicit-GEMM conv, including small per-group channel counts.

Technical Details

• grid-per-group scheme (group on blockIdx.z); groups==1 stays byte-identical
• make_a_descriptor: grouped channel decode (unmerge k->y,x,c_in_group over cpg, then embed((group,c_in_group)->c))
• make_b_descriptor: channel extent C->cpg
• build body + epilogues: thread group index into A (group coord) and fold into B/D as k_out = g*kpg + n
• implicit_gemm_conv_grid() -> (gn, gm, groups); group-consistency in validate()
• dispatch: drop the groups==1 gate (validate C/K % G), add bf16, grid z=groups, relax no-pad divisibility to M-only for grouped (N/K are tail-safe)

Test Plan

test_grouped_conv_numeric.py, dispatch_tests/conv/test_conv.py

Passed. groups==1 byte-identity gate GREEN.

Submission Checklist

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

JIRA ID: AICK-1752
Changes are behind the feature flag HIPKERNELPROVIDER_ENABLE_ROCKE which is defaulted to OFF.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*). Current: code file(s) changed: dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/helper_rocke.instances.common.conv_implicit_gemm.h, dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_conv_implicit_gemm_internal.h, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_implicit_gemm.cpp, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_implicit_gemm_conv_build_glue.cpp, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_implicit_gemm_conv_epilogue.cpp (+7 more); no test file found
🔎 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?

[Copilot]

Pull request overview

Adds grouped (groups>1) forward Conv2d support to the implicit-GEMM convolution path in Rocke, extending descriptor construction, kernel indexing, epilogues, and dispatcher behavior so grouped and cardinality-grouped shapes are supported while preserving groups==1 byte identity.

Changes:

• Extend implicit-GEMM conv descriptors and kernel body/epilogues to handle per-group channel decoding and blockIdx.z as the group index.
• Update conv dispatcher validation/selection to accept grouped conv (C/K divisible by G), bf16, and to use a shared (gn, gm, groups) launch grid; add gfx1250 WMMA candidate and tests.
• Add GPU numeric correctness and dispatch+launch tests for grouped conv, including small per-group channel “hero” cases and bf16.

Reviewed changes

Copilot reviewed 4 out of 4 changed files in this pull request and generated 2 comments.

File	Description
dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/test_grouped_conv_numeric.py	New on-GPU numeric correctness + dispatch/launch tests covering grouped forward conv.
dnn-providers/hip-kernel-provider/rocke/platform/tests/dispatch/dispatch_tests/conv/test_conv.py	Dispatcher unit tests updated for bf16, grouped conv selection, and gfx1250 WMMA behavior.
dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/instances/common/conv_implicit_gemm.py	Core grouped-conv enablement: grid shape, A/B descriptor changes, group-aware indexing and epilogues, WMMA allow-list update.
dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/dispatch/families/conv.py	Dispatcher now builds dtype spec, validates grouped constraints, relaxes divisibility gates for grouped, and adds gfx1250 WMMA candidate + grouped grid.

[therock-pr-bot]

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