[aledudek]

Motivation

Convolution PRs keep needing @ROCm/rocke-platform review even when only conv code changes. Most recently that was #13993. The cause is that the conv C++ mirrors, the Python-to-C++ bridge, the IR goldens, the CTest enrolment and the conv docs all live under rocke/platform/.

This PR moves everything conv-owned that can reasonably move under rocke/library/, which @ROCm/rocke-lib owns. Platform reviews the mechanism once. After that, conv-only changes touch only library/. Generic platform work (IR ops, shared helpers, arch data, lowering) stays in platform and is still reviewed there.

The mechanism is not conv-specific. Attention/SDPA, whose C++ mirrors also live in platform/cpp, can use it next.

Technical Details

The PR has three commits, meant to be reviewed in order.

1. Enablers. No conv code moves in this commit.

• platform/CMakeLists.txt adds ../library/cpp when that directory exists. It builds two targets:

  • rocke_library_core (librocke_library.a);
  • the pybind module rocke_library_engine.

  Both link rocke_core; nothing in platform links back. -DROCKE_BUILD_LIBRARY_CPP=OFF skips them, and a copy of platform/ alone still configures and gates.

• run_diff.py and check_byte_identity.py: library parity emitters compile against library/cpp/include and link librocke_library.a.

  • Families found under library/tests/parity are anchored in library/tests/parity/golden/llvm_gfx_all.json. Their entries were moved verbatim, so the union of the two golden files is unchanged.
  • The gate fails loudly if library C++ exists but its archive was not built.

• The library CTest entries move from platform/CMakeLists.txt into library/tests/ctest_lists.cmake. The installed test list is unchanged.

• rocke.core.backend publishes import_engine(module), lower_family and name_of, so library bridges reuse the same python/cpp/both dispatch.

• New platform/tests/test_library_reach.py freezes platform's imports of library packages. It is shrink-only.

• AGENTS.md, KERNEL_AUTHORING.md, TESTING.md and BUILDING.md document the new layout.

2. C++ mirrors.

• The conv_*, img2col and deep_fused_conv_pool instances and headers, plus wgrad_streamk_host.cpp, move with git mv to library/cpp/{instances,include,tests}. File contents are unchanged and the rocke/ include prefix is kept.
• Their bindings move out of rocke_engine into rocke_library_engine.
• The conv spec-to-dict converters and the lower_conv_*, lower_img2col and lower_deep_fused_conv_pool functions move from rocke.core.backend to library/kernels/conv/engine_bridge.py.

3. Goldens, helpers, tests and docs.

• The conv representative-IR cases and their golden move to library/tests/test_conv_ir_golden.py and library/tests/golden/conv_ir_sha256.json. The entries were moved verbatim, not re-blessed.
• check_ir_validity.py and check_arch_domain.py gain --cases <file>. library/tests/run_all.py runs both over the conv cases, so their coverage is kept.
• The split-K planner, select_split_k_wgrad and ConvKOuterFragmentReader move into library kernels.
• Conv tests move from test_rocke.py, tests/core and test_rocke_numeric.py into library tests, and are enrolled in ctest_lists.cmake.
• convolution.md, ck_tile_merged_groups.md, the conv case studies and run_one_dgrad.py move into library. The pooling docs stay in platform (small_ops.md).

Removed platform API (new homes are listed above):

• the conv entry points of rocke_engine;
• rocke.core.backend.lower_conv_* and *_spec_to_dict;
• rocke.helpers.split_k_planner;
• rocke.helpers.split_k.select_split_k_wgrad;
• rocke.helpers.ConvKOuterFragmentReader.

Open conv branches that touch the moved files need a rebase. Their golden changes must be re-blessed into the new library golden files rather than merged.

Test Plan

Each item was run after every commit and compared against develop at the same base:

• platform/tools/check_byte_identity.py, the byte-identity gate between the Python and C++ engines;
• rocke_ir_parity_harness.py --check (platform) and test_conv_ir_golden.py (library), each on every LLVM flavor;
• check_ir_validity.py and check_arch_domain.py over the platform corpus, and with --cases over the conv corpus;
• ctest for C++ host tests, including the moved rocke_wgrad_streamk_host;
• platform pytest (-m "not integration") and the library conv CPU tests, with failure sets diffed against develop;
• backend="both" lowering of every conv parity config, before and after the move;
• the generated installed-CTest entry list, before and after;
• configure, build and gate of a platform-only copy with no library/ beside it.

Test Result

• Byte-identity: every family GREEN after each commit; the conv families' emitters now link librocke_library.a.
• IR goldens: platform and library goldens pass on every flavor. Regenerating either golden from scratch reproduces the moved entries byte for byte.
• IR validity and arch-domain: green for both the platform corpus and the conv corpus.
• ctest: all tests pass.
• pytest: platform and library failure sets are identical to develop. Every test moved out of platform passes in library.
• Both-engine lowering: results are identical before and after the move.
• Installed CTest list: the same entries, plus the newly enrolled conv files.
• Platform-only copy: configures, builds and gates green.

Submission Checklist

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

[therock-pr-bot]

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

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

🙋 Wish to Override Policy?

[therock-pr-bot]

⏳ Waiting for required checks to start or finish:

• ⏳ pre-commit

[bartekxk]

PR #14422 — "refactor(rocke): move convolution code from platform into library", author aledudek, branch users/aledudek/rocke-move-conv-to-library → develop. As of 2026-10-27: draft, no reviews assigned, all checks green.

What it does: moves convolution implementation (instances, C++ mirrors/headers/bindings, goldens, helpers, tests, docs) out of platform/ and into library/, in 3 commits (enablers → C++ mirrors via git mv → goldens/helpers/tests/docs). Net ~180 files, +14.5k/-13.2k lines, content preserved (pure moves, not rewrites).

Introduces a new pattern: a librocke_library.a + rocke_library_engine pybind binding alongside the existing librocke_core.a/rocke_engine, so library families (starting with conv) own both their Python builder (library/kernels/conv/) and C++ mirror (library/cpp/), instead of living in platform/.

Why: conv PRs kept needing @ROCm/rocke-platform review even for conv-only changes; this lets conv owners (@ROCm/rocke-lib) merge independently. Author explicitly frames this as a template for future families (Attention/SDPA named directly).

How to apply: this directly extends the rocke/CLAUDE.md library/platform layering rules (one-way library → platform dependency, no cycles). If/when this lands, conv code paths referenced in older notes or code (e.g. platform/python/rocke/helpers/ for conv, platform/tests/instances/ for conv) will have moved to library/kernels/conv/ and library/tests/. Since the PR is still draft/unmerged, verify current location before assuming the move happened — check git log/file existence, don't trust this snapshot. Also relevant if doing SDPA/attention work later, since the author intends the same split to be copied there.

Code review (done 2026-10-27, against actual diff — fetched PR branch via git fetch origin pull/14422/head, diffed against the correct merge-base 2eee8a7e567/9850374..., not the stale local develop)

Verified the move itself is clean and mechanically correct: no conv golden entries lost or duplicated between platform/tests/golden/rocke_representative_ir_sha256.json and the new library/tests/golden/conv_ir_sha256.json (133–135 cases per LLVM flavor, zero overlap); the C++ binding extraction (rocke_engine.cpp → rocke_library_engine_conv.cpp/_img2col.cpp) preserves every reg3/m.def registration 1:1; the Python bridge (backend.py's _lower_family/_name_of → public lower_family/name_of/import_engine, kept as back-compat aliases since still used by the surviving gemm functions) is correctly rewired so library conv bridges (library/kernels/conv/engine_bridge.py) import rocke_library_engine instead of rocke_engine, and no stale references to the removed rocke.core.backend.lower_conv_* functions remain anywhere in the tree. check_byte_identity.py/run_diff.py correctly grew a second archive (librocke_library.a) and golden file (library/tests/parity/golden/llvm_gfx_all.json) path, with a golden-conflict guard if a family is accidentally blessed in both files.

Two real findings reported (both architecture/quality, not correctness bugs — see PR comments):

1. library/kernels/common/conv_split_k.py imports underscore-prefixed "private" _ARCH_NUM_CUS/_DEFAULT_NUM_CUS straight from rocke.helpers.split_k in platform — a fragile cross-boundary coupling on an implementation detail, working against the PR's own "library owns both halves" goal.
2. The new platform/tests/test_library_reach.py ratchet test (good design — AST-walks platform files, freezes-but-never-raises the count of platform→library imports) reveals platform already violates the documented one-way library → platform dependency in several files, worst being tests/instances/rocke_ir_parity_harness.py with 19 such imports. This PR doesn't introduce it, just formalizes tracking it — worth checking if there's a follow-up plan to burn it down, since CLAUDE.md states the invariant as unconditional.

[aledudek]

PR #14422 — "refactor(rocke): move convolution code from platform into library", author aledudek, branch users/aledudek/rocke-move-conv-to-library → develop. As of 2026-10-27: draft, no reviews assigned, all checks green.

What it does: moves convolution implementation (instances, C++ mirrors/headers/bindings, goldens, helpers, tests, docs) out of platform/ and into library/, in 3 commits (enablers → C++ mirrors via git mv → goldens/helpers/tests/docs). Net ~180 files, +14.5k/-13.2k lines, content preserved (pure moves, not rewrites).

Introduces a new pattern: a librocke_library.a + rocke_library_engine pybind binding alongside the existing librocke_core.a/rocke_engine, so library families (starting with conv) own both their Python builder (library/kernels/conv/) and C++ mirror (library/cpp/), instead of living in platform/.

Why: conv PRs kept needing @ROCm/rocke-platform review even for conv-only changes; this lets conv owners (@ROCm/rocke-lib) merge independently. Author explicitly frames this as a template for future families (Attention/SDPA named directly).

How to apply: this directly extends the rocke/CLAUDE.md library/platform layering rules (one-way library → platform dependency, no cycles). If/when this lands, conv code paths referenced in older notes or code (e.g. platform/python/rocke/helpers/ for conv, platform/tests/instances/ for conv) will have moved to library/kernels/conv/ and library/tests/. Since the PR is still draft/unmerged, verify current location before assuming the move happened — check git log/file existence, don't trust this snapshot. Also relevant if doing SDPA/attention work later, since the author intends the same split to be copied there.

Code review (done 2026-10-27, against actual diff — fetched PR branch via git fetch origin pull/14422/head, diffed against the correct merge-base 2eee8a7e567/9850374..., not the stale local develop)

Verified the move itself is clean and mechanically correct: no conv golden entries lost or duplicated between platform/tests/golden/rocke_representative_ir_sha256.json and the new library/tests/golden/conv_ir_sha256.json (133–135 cases per LLVM flavor, zero overlap); the C++ binding extraction (rocke_engine.cpp → rocke_library_engine_conv.cpp/_img2col.cpp) preserves every reg3/m.def registration 1:1; the Python bridge (backend.py's _lower_family/_name_of → public lower_family/name_of/import_engine, kept as back-compat aliases since still used by the surviving gemm functions) is correctly rewired so library conv bridges (library/kernels/conv/engine_bridge.py) import rocke_library_engine instead of rocke_engine, and no stale references to the removed rocke.core.backend.lower_conv_* functions remain anywhere in the tree. check_byte_identity.py/run_diff.py correctly grew a second archive (librocke_library.a) and golden file (library/tests/parity/golden/llvm_gfx_all.json) path, with a golden-conflict guard if a family is accidentally blessed in both files.

Two real findings reported (both architecture/quality, not correctness bugs — see PR comments):

1. library/kernels/common/conv_split_k.py imports underscore-prefixed "private" _ARCH_NUM_CUS/_DEFAULT_NUM_CUS straight from rocke.helpers.split_k in platform — a fragile cross-boundary coupling on an implementation detail, working against the PR's own "library owns both halves" goal.
2. The new platform/tests/test_library_reach.py ratchet test (good design — AST-walks platform files, freezes-but-never-raises the count of platform→library imports) reveals platform already violates the documented one-way library → platform dependency in several files, worst being tests/instances/rocke_ir_parity_harness.py with 19 such imports. This PR doesn't introduce it, just formalizes tracking it — worth checking if there's a follow-up plan to burn it down, since CLAUDE.md states the invariant as unconditional.

Thanks for the review, both points addressed, and the branch is rebased onto current develop.

1. Private _ARCH_NUM_CUS / _DEFAULT_NUM_CUS imports. Agreed, and the coupling was wider than conv_split_k.py: six library files reached into those names (conv split-K, conv wgrad stream-K pool, grouped-conv dispatch, attention dispatch, the split-K cost-fit benchmark and a test). rocke.helpers.split_k now exports a public, read-only ARCH_NUM_CUS, DEFAULT_NUM_CUS and num_cus(arch), and every library caller uses those. num_cus is the same plain lookup as before (no target-id stripping), so the stream-K default pool still matches its C++ mirror. While rebasing I also finished moving the conv split-K code: select_split_k_wgrad, select_split_k_fwd and WgradSplitPolicy now live in library/kernels/common/conv_split_k.py. Platform keeps only the generic GEMM selector, SplitKDecision and the CU table.

2. Platform → library back-edges. Right, these predate this PR; test_library_reach.py exists to stop them growing and to force each count down as code moves. This PR already burns a large part down: platform→library imports go from 145 in 24 files on develop to 50 in 18 files. rocke_ir_parity_harness.py goes from 72 to 19, and the 19 left are attention, KDA and lightning-indexer cases, with no conv. The plan is the same mechanism for attention/SDPA (named in the description): moving those cases next to their kernels empties most of the harness's remaining count. So the rule doesn't read as unconditional while exceptions exist, CLAUDE.md/AGENTS.md now point at the ratchet as the list of frozen pre-existing back-edges that may only shrink.

Rebased onto develop (961dd76). Develop had restructured several things this PR touches, so the rebase re-applies the move to them rather than just resolving conflicts:

• the representative-IR golden is now per-family shards, so the conv shards were git mv'd to library/tests/golden/conv_representative_ir/, with no re-bless;
• the differential golden moved to golden_store schema 2, and the platform/library split now sits on top of it;
• the new gfx1250 WMMA/Winograd conv C++ and bindings moved with the rest;
• the conv C++ host tests moved out of the new platform/tests/cmake fragments.

Byte-identity is GREEN, and every golden is OK on llvm20/22/23/24 with the same blessed and unblessed counts as develop. Commit structure is unchanged; commits 1 and 2 each build on their own.

[jakpiase]

Review: blockers only

Verdict: approve with fixes. At its current base the move is clean. The blocker is the rebase onto develop. Paths below are relative to dnn-providers/hip-kernel-provider/rocke/.

At the current base (961dd76) the move is clean

Checked on CPU by building and running the gates:

• Pure move: git sees 162 renames, 14 adds and 68 modifies, and no file is deleted without a counterpart. The renamed C++ differs only in comments. All 77 test functions removed from platform exist in library tests.
• Python bindings: the 120 registered names are all still present (now 122; the new ones are build_id and engine_version).
• Golden split: across ir_canonical and all four ll flavours, the platform and library golden files together exactly equal the old platform file. Nothing is lost, duplicated or changed.
• Build: works with ROCKE_BUILD_PYBIND OFF and ON. A copy of platform/ without library/ next to it still configures and builds.
• Tests and gates: the CTest list is identical to base (46 tests, all pass). check_byte_identity is GREEN on the same 95 families. 337 library and 428 platform CPU pytests pass.

Blocker: develop has added new conv code to platform/ since this PR's base

Seven files conflict, and that part is visible. The danger is that the new conv files merge with no conflict and stay in platform/. A rebase that just resolves the conflicts in the PR's favour then breaks the build or quietly loses test coverage:

1. The rocke_core build breaks (reproduced with a trial merge).

   • Develop adds 8 conv files under platform/cpp/:
     • instances/common/conv_direct_dw_fwd_wmma.cpp, conv_direct_fwd_gstream_wmma.cpp, conv_direct_fwd_tapouter_wmma.cpp, conv_fwd_wmma_tdm.cpp
     • include/rocke/instance_conv_direct_dw_fwd_wmma.h, instance_conv_direct_fwd_gstream_wmma.h, instance_conv_direct_fwd_tapouter_wmma.h, instance_conv_fwd_wmma_tdm_internal.h
   • platform/CMakeLists.txt:160 compiles every cpp/*.cpp it finds, but the headers these files include now exist only under library/cpp/include:

     instance_conv_direct_fwd_gstream_wmma.h:40:10: fatal error: rocke/instance_conv_direct_grouped.h: No such file or directory
     

     dw_fwd_wmma and tapouter_wmma fail the same way.
   • Fix: move these 8 files into library/cpp/{instances/common,include/rocke}.

2. Four test files can drop out of CTest without any error.

   • Develop adds test_direct_dw_fwd_wmma.py, test_direct_fwd_gstream_wmma.py, test_direct_fwd_tapouter_wmma.py and test_conv_fill.py to the CTest lists in platform/CMakeLists.txt, inside the conflicting blocks.
   • This PR moved those lists to library/tests/ctest_lists.cmake, so keeping the PR side drops all four.
   • Fix: add them to library/tests/ctest_lists.cmake in both the CPU and GPU lanes, then diff the installed CTestTestfile.cmake against develop.

3. Three new kernel families can lose golden gating while CI still passes.

   • Develop adds golden entries to platform/tests/instances/differential/golden/llvm_gfx_all.json for conv_direct_dw_fwd_wmma, conv_direct_fwd_gstream_wmma and conv_direct_fwd_tapouter_wmma. It also re-blesses the llvm24 hashes for conv_direct_fwd_wmma and conv_fwd_small_wmma.
   • After this PR, those families are checked against library/tests/parity/golden/llvm_gfx_all.json, which git won't update because the PR creates it.
   • Configs with no blessed golden are reported as information only (NEW (unblessed) … GOLDEN OK). If the entries are dropped from the platform file and never re-blessed into the library file, the three families are no longer gated and the check still passes.
   • Develop also adds:
     • 3 golden files under platform/tests/golden/rocke_representative_ir_sha256/conv_direct_*_wmma.json;
     • about 250 lines of conv cases in rocke_ir_parity_harness.py.
   • Fix: move the golden entries into the library golden file, the 3 golden files into library/tests/golden/conv_representative_ir/, and the parity cases into library/tests/test_conv_ir_golden.py.

4. Develop's new conv Python and binding code must move too.

   • Develop adds conv_direct_{dw_fwd,fwd_gstream,fwd_tapouter}_wmma_spec_to_dict and lower_conv_direct_* to platform/python/rocke/core/backend.py.
   • It also adds the matching reg3 / *_describe bindings to platform/cpp/bindings/rocke_engine.cpp. Those bindings include the headers moved in item 1.
   • Left in place, ROCKE_BUILD_PYBIND=ON breaks, and conv code ends up in platform again.
   • Fix: move them into library/kernels/conv/engine_bridge.py and library/cpp/bindings/rocke_library_engine_conv.cpp.

Merge condition: after the rebase, the heavy rocke CI (byte-identity and goldens on all LLVM flavours) must pass on the new head. Only the 5 lightweight checks have run so far.

Not verified

• The rebased head (it doesn't exist yet).
• GPU lanes.
• The full TheRock superbuild and wheel packaging.
• IR parity at llvm20/22/24 with the library tree in place (llvm23 was checked).

[Copilot]

🟡 Changes recommended

Unresolved critical installation regressions leave convolution headers and symbols unavailable to C/C++ consumers and break installed test tooling.

8 open issues

• Install and export the moved convolution API or provide compatibility · New
• Restrict the source-tree audit to source checkouts · New
• Stage _corpus.py for installed arch-domain tests · New
• Include the library corpus in the top-level validation gate · New
• Preserve the source package in the subprocess environment · New
• Skip compilation only when COMGR is unavailable · New
• Document PoolOp as a string-valued alias · New
• Match the pooling API's field names and string operation type · New

What changed in this PR

This PR separates convolution ownership from rocKE’s platform layer by relocating convolution C++ mirrors, tests, goldens and documentation into library/, while retaining shared platform infrastructure.

Changes:

• Adds a library C++ archive and companion rocke_library_engine bindings.
• Moves convolution validation assets into library-owned suites and extends IR gates to accept separate corpora.
• Updates helper imports, documentation references and platform dependency auditing.

File	Description
dnn-providers/​hip-kernel-provider/​rocke/​TESTING.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tools/​check_ir_validity.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tools/​check_byte_identity.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tools/​check_arch_domain.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tools/​_corpus.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​test_rocke_ci_static.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​test_library_reach.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​test_binding_compiler_selection.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​test_rocke_numeric.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​differential/​test_check_byte_identity_progress.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​differential/​golden/​README.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​core/​test_pipeline_basic.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​core/​test_ir_serialize.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​core/​test_arch_domain_lane.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​cmake/​InstanceTests.cmake	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​helpers/​__init__.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​examples/​gfx1201/​README.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​examples/​gfx1201/​ALGORITHM.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​reference/​env_flags.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​README.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​optimization/​measured_results.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​instances/​small_ops.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​development/​verify_dsl_docs.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​architecture/​index.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​bindings/​prove_parity_binding.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​bindings/​CMakeLists.txt	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​AGENTS.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_verify_dsl_docs_conv.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_fwd_small_wmma.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_direct_fwd_wmma.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_direct_dgrad_wmma.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_dgrad_stem_wmma.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_winograd_gen_cpu.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_wgrad_streamk.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_split_k.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_split_k_wgrad.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_split_k_planner.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_implicit_gemm_vector_width.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_implicit_gemm_numeric.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_fwd_group_merge_gate.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_direct_grouped_backend_parity.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_conv_dgrad_correctness.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_ceil_div_callers.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​run_all.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​golden/​conv_representative_ir/​conv_winograd_gemm_wmma.json	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​golden/​conv_representative_ir/​conv_winograd_fused_1d_wmma.json	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​golden/​conv_representative_ir/​conv_wgrad_two_stage.json	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​golden/​conv_representative_ir/​conv_fwd_split_k_reduce.json	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​golden/​conv_representative_ir/​conv_direct_fwd_smallk_wmma.json	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​golden/​conv_representative_ir/​conv_direct_fwd_smallc_wmma.json	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​golden/​conv_representative_ir/​conv_dgrad_smallc_wmma.json	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​golden/​conv_representative_ir/​conv_dgrad_patch_wmma.json	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​conv/​docs/​ck_tile_merged_groups.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​conv/​__init__.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​README_conv_implicit_gemm_wgrad.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​README_conv_implicit_gemm_dgrad.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​README_conv_direct_nongrouped.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​README_conv_direct_grouped.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​conv_implicit_gemm_wgrad.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​conv_implicit_gemm_dgrad.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​conv_fwd_wmma_gate.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​conv_direct_grouped.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​dispatch/​grouped_convolution.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​dispatch/​attention/​common.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​tests/​wgrad_streamk_host.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​tests/​CMakeLists.txt	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​instances/​gfx1201/​deep_fused_conv_pool_maxpool_gates_emitters.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​instances/​gfx1201/​deep_fused_conv_pool_conv1_gemm_lds_staging.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​instances/​common/​conv_implicit_gemm_wgrad.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​instances/​common/​conv_implicit_gemm_conv_compute_phase.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​instances/​common/​conv_fwd_wmma_gate.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​include/​rocke/​instance_conv_winograd_fwd_rt.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​include/​rocke/​instance_conv_fwd_wmma_gate_internal.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​include/​rocke/​instance_conv_fwd_split_k_reduce.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​include/​rocke/​instance_conv_direct_wgrad_gather.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​include/​rocke/​instance_conv_direct_dw_wgrad_mfma.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​include/​rocke/​instance_conv_abi.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​CMakeLists.txt	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​bindings/​winograd_glue.hpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​bindings/​rocke_library_engine.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​bindings/​rocke_library_engine_winograd_gen.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​bindings/​rocke_library_engine_winograd_gemm.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​bindings/​rocke_library_engine_winograd_fwd.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​bindings/​rocke_library_engine_winograd_fused1d.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​cpp/​bindings/​rocke_library_engine_img2col.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_wgrad/​wgrad_split_k_planner_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_wgrad/​wgrad_merged_groups_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_wgrad/​wgrad_lds_layout_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_fwd/​fwd_merged_groups_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_dgrad/​stride1_igemm_dgrad_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_dgrad/​README.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_dgrad/​grouped_generic_row_stream_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_dgrad/​dgrad_lds_layout_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_dgrad/​dgrad_fused_weights_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_dgrad/​dgrad_4c_bf16_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx950/​conv_dgrad/​depthwise_windowed_dgrad_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​builders/​gfx1250/​conv_kouter/​gfx1250_transpose_read_case_study.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​benchmarks/​gfx950/​conv_dgrad/​run_one_dgrad.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​benchmarks/​common/​split_k_cost_fit.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​benchmarks/​common/​benchmark_implicit_gemm_conv.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​KERNEL_AUTHORING.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​BUILDING.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​AGENTS.md	Updated as part of this pull request.

_{Review effort: Max}

[aledudek]

Review: blockers only

Verdict: approve with fixes. At its current base the move is clean. The blocker is the rebase onto develop. Paths below are relative to dnn-providers/hip-kernel-provider/rocke/.

At the current base (961dd76) the move is clean

Checked on CPU by building and running the gates:

• Pure move: git sees 162 renames, 14 adds and 68 modifies, and no file is deleted without a counterpart. The renamed C++ differs only in comments. All 77 test functions removed from platform exist in library tests.
• Python bindings: the 120 registered names are all still present (now 122; the new ones are build_id and engine_version).
• Golden split: across ir_canonical and all four ll flavours, the platform and library golden files together exactly equal the old platform file. Nothing is lost, duplicated or changed.
• Build: works with ROCKE_BUILD_PYBIND OFF and ON. A copy of platform/ without library/ next to it still configures and builds.
• Tests and gates: the CTest list is identical to base (46 tests, all pass). check_byte_identity is GREEN on the same 95 families. 337 library and 428 platform CPU pytests pass.

Blocker: develop has added new conv code to platform/ since this PR's base

Seven files conflict, and that part is visible. The danger is that the new conv files merge with no conflict and stay in platform/. A rebase that just resolves the conflicts in the PR's favour then breaks the build or quietly loses test coverage:

1. The rocke_core build breaks (reproduced with a trial merge).

   • Develop adds 8 conv files under platform/cpp/:

     • instances/common/conv_direct_dw_fwd_wmma.cpp, conv_direct_fwd_gstream_wmma.cpp, conv_direct_fwd_tapouter_wmma.cpp, conv_fwd_wmma_tdm.cpp
     • include/rocke/instance_conv_direct_dw_fwd_wmma.h, instance_conv_direct_fwd_gstream_wmma.h, instance_conv_direct_fwd_tapouter_wmma.h, instance_conv_fwd_wmma_tdm_internal.h

   • platform/CMakeLists.txt:160 compiles every cpp/*.cpp it finds, but the headers these files include now exist only under library/cpp/include:

     instance_conv_direct_fwd_gstream_wmma.h:40:10: fatal error: rocke/instance_conv_direct_grouped.h: No such file or directory
     

     dw_fwd_wmma and tapouter_wmma fail the same way.

   • Fix: move these 8 files into library/cpp/{instances/common,include/rocke}.

2. Four test files can drop out of CTest without any error.

   • Develop adds test_direct_dw_fwd_wmma.py, test_direct_fwd_gstream_wmma.py, test_direct_fwd_tapouter_wmma.py and test_conv_fill.py to the CTest lists in platform/CMakeLists.txt, inside the conflicting blocks.
   • This PR moved those lists to library/tests/ctest_lists.cmake, so keeping the PR side drops all four.
   • Fix: add them to library/tests/ctest_lists.cmake in both the CPU and GPU lanes, then diff the installed CTestTestfile.cmake against develop.

3. Three new kernel families can lose golden gating while CI still passes.

   • Develop adds golden entries to platform/tests/instances/differential/golden/llvm_gfx_all.json for conv_direct_dw_fwd_wmma, conv_direct_fwd_gstream_wmma and conv_direct_fwd_tapouter_wmma. It also re-blesses the llvm24 hashes for conv_direct_fwd_wmma and conv_fwd_small_wmma.

   • After this PR, those families are checked against library/tests/parity/golden/llvm_gfx_all.json, which git won't update because the PR creates it.

   • Configs with no blessed golden are reported as information only (NEW (unblessed) … GOLDEN OK). If the entries are dropped from the platform file and never re-blessed into the library file, the three families are no longer gated and the check still passes.

   • Develop also adds:

     • 3 golden files under platform/tests/golden/rocke_representative_ir_sha256/conv_direct_*_wmma.json;
     • about 250 lines of conv cases in rocke_ir_parity_harness.py.

   • Fix: move the golden entries into the library golden file, the 3 golden files into library/tests/golden/conv_representative_ir/, and the parity cases into library/tests/test_conv_ir_golden.py.

4. Develop's new conv Python and binding code must move too.

   • Develop adds conv_direct_{dw_fwd,fwd_gstream,fwd_tapouter}_wmma_spec_to_dict and lower_conv_direct_* to platform/python/rocke/core/backend.py.
   • It also adds the matching reg3 / *_describe bindings to platform/cpp/bindings/rocke_engine.cpp. Those bindings include the headers moved in item 1.
   • Left in place, ROCKE_BUILD_PYBIND=ON breaks, and conv code ends up in platform again.
   • Fix: move them into library/kernels/conv/engine_bridge.py and library/cpp/bindings/rocke_library_engine_conv.cpp.

Merge condition: after the rebase, the heavy rocke CI (byte-identity and goldens on all LLVM flavours) must pass on the new head. Only the 5 lightweight checks have run so far.

Not verified

• The rebased head (it doesn't exist yet).
• GPU lanes.
• The full TheRock superbuild and wheel packaging.
• IR parity at llvm20/22/24 with the library tree in place (llvm23 was checked).

Thanks @jakpiase, this is exactly the trap: the new conv files merged cleanly and would have stayed in platform/. Rebased onto develop 0c4056e and moved everything #14835 added:

1. rocke_core build: the 8 files (conv_direct_{dw_fwd,fwd_gstream,fwd_tapouter}_wmma.cpp, conv_fwd_wmma_tdm.cpp and their 4 headers) are git mv'd to library/cpp/{instances/common,include/rocke}; no conv source or header is left under platform/cpp. Builds with ROCKE_BUILD_PYBIND OFF and ON, and a platform/-only copy still configures and builds.
2. CTest: test_direct_dw_fwd_wmma.py, test_direct_fwd_gstream_wmma.py, test_direct_fwd_tapouter_wmma.py (CPU and GPU lanes) and test_conv_fill.py (CPU lane) are in library/tests/ctest_lists.cmake. The installed CTestTestfile.cmake has the same 64 tests as develop; the only differences are the conv test files this PR moved out of platform, which the library lanes now also list.
3. Golden gating: I re-split develop's llvm_gfx_all.json instead of merging it, so the library file now carries the 3 new families and the conv_direct_fwd_wmma / conv_fwd_small_wmma additions, and nothing is unblessed that develop blesses: platform + library = develop's 1725 blessed / 349 unblessed ir_canonical configs exactly, with 0 changed. The 3 representative-IR files moved to library/tests/golden/conv_representative_ir/, and the new harness cases to test_conv_ir_golden.py (develop's 395 cases = 190 platform + 205 library, no overlap; the NVFP4 GEMM cases that landed in the same block stay in platform). Both IR goldens pass at llvm20/22/23/24.
4. Bridge and bindings: the new spec_to_dict, lower_conv_direct and the reg3 / *_describe bindings are in engine_bridge.py and rocke_library_engine_conv.cpp. All 206 entry points develop registers are present across the two modules, and none of the conv ones remain in rocke_engine. Develop's new conv tests now import the library engine and bridge.

Gates on the rebased head: byte-identity GREEN (same as develop), ctest 46/46, golden and differential checks OK on all four flavors, and commits 1 and 2 each build on their own. The library CPU suites fail only the same gfx1250 Winograd build tests as develop (host toolchain). GPU lanes and the full superbuild still need CI.

[cgmillette]

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max

PR #14422 — refactor(rocke): move convolution code from platform into library (aledudek) — 🟠 First-round: rebase must carry develop's work

Prior item	Status
jakpiase (11-04) 4 blockers: build, CTest, goldens, bindings	🟢 FIXED 31b2d97842 • all four hold at head for 14835's code: build, CTest, goldens, bindings (ran) • same trap again: 14832, 15024, 14259 (N1–N3); 14833 next (N4) • merge condition open: no heavy CI on any head since 10-26
Copilot (11-05) 8 threads: install, audit, gates, docs	🟢 FIXED 31b2d97842 • ckc::rocke_library, archive, headers install; _corpus.py imports; audit skips installed (ran) • other four fixed (read); no test links ckc::rocke_library yet • small_ops.md:508 fixed but its thread is open: resolve it
bartekxk (10-27) private CU names; platform→library back-edges	🟢 FIXED 4190bb0b5d • library imports public ARCH_NUM_CUS, num_cus; no private name left (read) • back-edges frozen per file, counts honest (ran); no CI lane runs the ratchet: N8
Bots / CI	🟠 Not run • fails: therock-pr-bot — GitHub API 500 inside its policy script (infra) • no pull_request workflow since 10-26: each push already conflicted • 2 ✓ (labeler, freshness); merge: dirty, 5 files
Our checks	🟠 Dirty • merge vs c3fcdb9e6d: 5 conflicts; PR-side trial fails goldens and tests: N1–N3 (ran) • head: ctest 47 (46 ✓, 1 skip, as base); byte identity GREEN on base's families, 4 flavors (ran) • goldens OK at 4 flavors; golden split exact (1,750 / 336 as base); 64 installed entries as base (ran) • installed changed tests 1,042 ✓, 2 ✗ also at base; moved files add about a fifth to the conv entry, already over its Windows limit (15035) (ran) • rules 1–3: one new PR number in a comment, test_conv_platform_regressions.py:435
Code quality	🟠 Should-fix • bindings and golden writer copied, not shared: N6 (ran) • .orig leftovers: N5; aliases the API standard forbids: N7 (read) • ratchet counts honest (ran); no CI lane runs it: N8 • AGENTS.md:49, CMakeLists.txt:6-8 still call library/ Python-only (read)

#	Sev	file:line	Finding
N1	🟠	platform/cpp/bindings/rocke_engine.cpp:3399	Three conflicts hold 14832's conv changes; a PR-side resolution drops them • trial, PR side (ran): 2 vector-cap tests (15 subtests) and 3 fp16 split-K tests fail; conv golden misses 4 cases • port rocke_engine.cpp:3399-3419 → rocke_library_engine_conv.cpp:1608; split_k.py:136-169,461-465 → conv_split_k.py:207-219,384 • port harness :199,232,2895-2962,3249-3269 → test_conv_ir_golden.py:153,1754,2040; with these, all pass (ran)
N2	🟠	platform/CMakeLists.txt:749	A PR-side CMake resolution drops 7 CTest items develop added • 4 attention CPU files, 1 GPU file, 2 cache-key entries (15024, 14259) • impact: they leave CI with no error; ctest_lists.cmake names none of them (ran) • fix: move develop's CMakeLists.txt:749-754,776-795,1047 to ctest_lists.cmake:90,122,309
N3	🟠	platform/tests/instances/differential/golden/llvm_gfx_all.json:1	Golden conflict: develop changed 1,191 entries of 9 library-owned families • trial, PR side (ran): library file 200 changed, 1 gone (red); unblessed 235 → 264 (silent) • from 14832 (conv wgrad, dgrad), 15024 and 14259 (7 attention families) • fix: re-split develop's file as at 8b6ada3, not hunk by hunk; block_scale_gemm (14843) stays (ran)
N4	🟠	14833 platform/cpp/instances/common/conv_direct_grouped_build_wgrad_image.cpp:38	Land order with 14833: its new conv sources would stay in platform • 3 conv sources and a host test under platform/ include headers this PR moves • impact: fatal error: rocke/instance_conv_direct_grouped.h: No such file or directory (ran) • fix: the second to land moves them into library/cpp; the merge plan lists 14833 first
N5	🟠	library/kernels/conv/engine_bridge.py.orig:1	Three pre-rebase .orig copies are committed (7,426 lines) • also rocke_library_engine_conv.cpp.orig, test_conv_ir_golden.py.orig • impact: two ship in the installed test tree (ran); greps and scans hit dead copies • fix: git rm the three
N6	🟠	library/cpp/bindings/rocke_library_engine_conv.cpp:63	The move copies platform machinery instead of sharing it • six helpers repeat rocke_engine.cpp:149-660 verbatim (ran); library/cpp/CMakeLists.txt:70-72 says "shared, not copied" • test_conv_ir_golden.py repeats the harness's golden writer and CLI: two homes (ran) • fix: helpers into family_glue.hpp; writer into rocke.core.ir_golden
N7	🟠	platform/python/rocke/core/backend.py:998	The new public bridge keeps its old names as aliases • _lower_family, _name_of, _import_engine (:998, :1006, :530); API_DESIGN_STANDARD.md Part 1 rule 5 (read) • lower_family never reads spec; nine positional parameters; not exported from rocke.core • fix: rename the callers and the test patch, delete the aliases; keyword-only parameters
N8	🟠	platform/tests/test_library_reach.py:94	The layering ratchet runs in no CI lane • installed rocke_pytest skips it (ran); AGENTS.md:62-67 cites it as the rule's enforcement • fix: run it from source in CI, e.g. a local pre-commit hook; pytest --noconftest on it passes from the repo root (ran)
N9	🟡	library/cpp/instances/common/conv_implicit_gemm_wgrad.cpp:170	Overlap: 14464 copies this CU table as ck_gemm_arch_num_cus • copies: here (develop's, moved); 14464 platform/cpp/instances/common/gemm_spec.cpp:1021, same table (ran) • owner 14464, natural home: platform cannot reach library, so the C++ table goes in platform • this PR keeps its copy; the second to land points it at that helper

Utility: 🟢 Positive — an exact move at its base; library owns conv end to end

---

Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

First-round review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: trial merges with develop, head build + ctest, byte identity and goldens at 4 flavors, installed CTest lists and tests, code quality. Flags whether a closer look is needed; not an approval or rejection.

[cgmillette]

Automated overlap notice (rocKE merge monitor)

Hi @aledudek, this PR and #14464 both keep a C++ copy of the per-arch CU table and both add a public Python lookup for it to rocke.helpers.split_k: here develop's wgrad_arch_num_cus, moved into library/cpp, and num_cus; in #14464 the platform helper rocke_arch_num_cus and arch_num_cus. Row N9 of our 6 November first-round review of this PR flagged the C++ copy. This notice records the owner, adds the Python side and gives rebase steps checked on the merge.

The copies, here at 31b2d97842 and in #14464 at 61501dd46f (paths under dnn-providers/hip-kernel-provider/rocke/):

• Here: develop's static wgrad_arch_num_cus, moved unchanged to library/cpp/instances/common/conv_implicit_gemm_wgrad.cpp:167-186 and called once, in rocke_wgrad_conv_streamk_default_pool (:195). In Python, platform/python/rocke/helpers/split_k.py:54-69 makes the table public and read-only (ARCH_NUM_CUS, a MappingProxyType, and DEFAULT_NUM_CUS) and adds num_cus(arch), which three library modules import as arch_num_cus (library/dispatch/grouped_convolution.py:120, library/kernels/common/conv_split_k.py:66, library/kernels/common/conv_implicit_gemm_wgrad.py:1474).
• feat(hipTensor/rocKE): Add a new rocKE backend to hipTensor and add optimized kernel support #14464: rocke_arch_num_cus() in platform/cpp/helpers/arch_num_cus.cpp:21-40, declared in platform/cpp/include/rocke/arch_num_cus.h:20 with ROCKE_DEFAULT_NUM_CUS (:17), called from platform/cpp/instances/common/gemm_cf32.cpp:217; platform/tests/instances/test_arch_num_cus.py:36-41 compares it with the Python table. A second, static copy, ck_gemm_arch_num_cus, sits in platform/cpp/instances/common/gemm_spec.cpp:1021 (called at :2180). In Python, arch_num_cus(arch) (split_k.py:55) reads the private _ARCH_NUM_CUS (:47); five other platform modules of feat(hipTensor/rocKE): Add a new rocKE backend to hipTensor and add optimized kernel support #14464 call it.

All three C++ tables hold the same three rows and the same default of 256.

Owner: #14464, as our review proposed and the monitor recorded. Platform code cannot use library code, so the one C++ table has to live in platform, and #14464 already puts it there: platform/cpp/helpers/arch_num_cus.cpp, the C mirror of rocke.helpers.split_k, which this PR's library code can include. The Python lookup goes with it: both PRs add it to the same module, and #14464's arch_num_cus came first (25 October) and is the function its C helper mirrors. This PR's C++ copy is develop's own, moved. Resolution: this PR keeps the moved develop copy until #14464's helper lands, then points rocke_wgrad_conv_streamk_default_pool at rocke_arch_num_cus and deletes wgrad_arch_num_cus. In Python it keeps its public, read-only ARCH_NUM_CUS and DEFAULT_NUM_CUS, which arch_num_cus then reads, and takes arch_num_cus in place of num_cus. On #14464's side the static copy in gemm_spec.cpp becomes a call to rocke_arch_num_cus, so one C++ table remains.

Order: the merge plan on develop c3fcdb9e6d lists neither PR among its 35 near-merge PRs. This PR conflicts with develop in 5 files (last push 5 November); #14464 conflicts in 24 (last push 1 November).

• If feat(hipTensor/rocKE): Add a new rocKE backend to hipTensor and add optimized kernel support #14464 lands first, this PR makes the steps below at its rebase.
• If this PR lands first, feat(hipTensor/rocKE): Add a new rocKE backend to hipTensor and add optimized kernel support #14464 makes the same switch at its rebase: arch_num_cus reads this PR's public table and replaces num_cus in the three library modules above; its test, its planner_ref.py and its helper's comments move to the public names (steps 1-2); and rocke_wgrad_conv_streamk_default_pool, then in library/cpp, calls rocke_arch_num_cus (step 3).

Rebase onto develop with #14464:

1. platform/python/rocke/helpers/split_k.py is where the two PRs meet on this piece (this PR already conflicts there with develop: N1 of our review). Keep this PR's public, read-only ARCH_NUM_CUS and DEFAULT_NUM_CUS; keep feat(hipTensor/rocKE): Add a new rocKE backend to hipTensor and add optimized kernel support #14464's arch_num_cus, reading them; drop num_cus. The XCD table and WgradSplitPolicy stay moved to library/kernels/common/conv_split_k.py. In the three library modules above, import arch_num_cus directly.
2. Four files feat(hipTensor/rocKE): Add a new rocKE backend to hipTensor and add optimized kernel support #14464 adds merge without a conflict but use the private names. platform/tests/instances/test_arch_num_cus.py:37 and :41 read split_k._ARCH_NUM_CUS and split_k._DEFAULT_NUM_CUS: use the public names. platform/tests/dispatch/dispatch_tests/tensor_gett/planner_ref.py:747 assigns into the table (SK._ARCH_NUM_CUS[arch] = int(cus)), which the read-only mapping refuses: rebind it instead (SK.ARCH_NUM_CUS = MappingProxyType({**SK.ARCH_NUM_CUS, arch: int(cus)}), with from types import MappingProxyType), since arch_num_cus reads the module's mapping at each call. None of the other split_k names feat(hipTensor/rocKE): Add a new rocKE backend to hipTensor and add optimized kernel support #14464 imports (arch_num_cus, select_split_k, TARGET_SLICE_K, _parse_env_override) is one this PR moves out. The comments at platform/cpp/helpers/arch_num_cus.cpp:5, 20 and platform/cpp/include/rocke/arch_num_cus.h:5, 16 name the private table too.
3. C++: in library/cpp/instances/common/conv_implicit_gemm_wgrad.cpp, include rocke/arch_num_cus.h beside the other platform helper headers (:39-41), call rocke_arch_num_cus(arch ? arch : "gfx950") at :195, and delete wgrad_arch_num_cus with its comment (:167-186); the comment at library/cpp/include/rocke/instance_conv_implicit_gemm_wgrad.h:266 then names rocke_arch_num_cus. The rows and the default are the same, so the pool does not change.
4. Check: the rocke_wgrad_conv_streamk_default_pool cases in library/cpp/tests/wgrad_streamk_host.cpp (CTest rocke_wgrad_streamk_host), and from platform/: PYTHONPATH=python python -m pytest tests/instances/test_arch_num_cus.py.

Both PRs also change 9 other files (among them platform/CMakeLists.txt, platform/cpp/bindings/rocke_engine.cpp, platform/python/rocke/core/backend.py and the differential golden); this notice covers the CU table only.

Checked with git merge-tree on develop c3fcdb9e6d (no build): this PR conflicts in 5 files and #14464 in 24, and #14464's split_k.py merges cleanly; a three-way merge of the two trial merges' split_k.py stops in one region, the CU block (with the XCD table and WgradSplitPolicy this PR moves out). The C++ tables sit in files only one PR changes, so that part has no conflict. On this PR's head, the read-only ARCH_NUM_CUS refuses item assignment, and the lookup returns a count from a rebound mapping. The copies, callers and line numbers come from both heads' sources.

If you think this PR should own it instead, reply here and on #14464, and the reviews will follow the new decision.

Automated overlap check: one owner per shared piece of code, coordinated on the PRs. Nothing else changed.

[CongMa13]

I reviewed this PR again. The following review threads are still unresolved, or comments have no author reply:

• refactor(rocke): move convolution code from platform into library #14422 (comment)
• refactor(rocke): move convolution code from platform into library #14422 (comment)
• refactor(rocke): move convolution code from platform into library #14422 (comment)

For reassessment, the PR author should reply to top-level feedback or state that thread feedback is addressed. Before approval, top-level comments need explicit resolution from their author or that author's approval of the PR.