refactor(rocke): move convolution code from platform into library

dnn-providers/hip-kernel-provider/rocke/AGENTS.md

@@ -21,7 +21,8 @@ artifact, and treat uncertainty as confidential until a human resolves it:
 ## The #1 invariant: byte-identity
 
 The Python engine (`platform/python/rocke/core/lower_llvm.py`) and the C++ engine
-(`platform/cpp/`) **MUST emit the same LLVM-IR bytes** for every kernel family. Any
+(`platform/cpp/`, plus `library/cpp/` for library-owned families) **MUST emit the
+same LLVM-IR bytes** for every kernel family. Any
 op / instance / atom / fusion / arch change must be made in **both** engines in the
 same change; if output is meant to change, re-bless the golden in the same change.
 
@@ -47,10 +48,23 @@ rocke/
 │   └── tools/ tests/ cmake/
 └── library/         # the SDPA product — build-time-only Python (NOT installed, no wheel)
     ├── kernels/ builders/ dispatch/ benchmarks/ tests/
+    └── cpp/           C++ mirrors of library families (librocke_library.a + rocke_library_engine)
 ```
 
+A library family owns both halves of its byte-identity contract: the Python
+builder in `library/kernels/` and its C++ mirror in `library/cpp/`, with parity
+emitters and their golden in `library/tests/parity/`. `platform/CMakeLists.txt`
+adds `library/cpp` when the sibling tree exists; its targets link `rocke_core`,
+never the reverse, and `tools/check_byte_identity.py` gates them with the
+engine's own families. Library tests enrol in CI through
+`library/tests/ctest_lists.cmake`.
+
 One-way dependency: **`library → platform` only**; platform never imports
-`kernels`/`builders`/`dispatch` and stays standalone-installable.
+`kernels`/`builders`/`dispatch` and stays standalone-installable. The
+back-edges that predate the rule (parity-harness cases, examples, tools) are
+frozen per file in
+[`platform/tests/test_library_reach.py`](platform/tests/test_library_reach.py):
+a count may only go down, and a new platform file may not add one.
 
 **No cycles inside `library/` either.** The packages are layered lowest → highest
 and may only import *downward*:
@@ -145,7 +159,8 @@ Hard rule from the runbook: **never report speed without correctness** — run t
 relevant parity/verify harness (attention → `library/builders/**/parity_*.py`;
 non-attention → `platform/python/rocke/examples/common/*parity*.py`) and stay within
 the runbook §1.4 tolerances before claiming any win. Leave a replayable case study in
-the relevant `examples/<arch>/<workload>/` folder.
+the relevant `examples/<arch>/<workload>/` folder (`library/builders/<arch>/<workload>/`
+for a library family).
 
 ## Deeper material
 

dnn-providers/hip-kernel-provider/rocke/BUILDING.md

@@ -136,4 +136,11 @@ Override the roots for out-of-tree consumption with `ROCKE_PLATFORM_ROOT`,
 
 The platform C++ engine (`platform/cpp/` → `librocke_core.a`) and its
 `rocke_engine` pybind binding are built by CMake, not pip, and are intentionally
-not part of any wheel. See `platform/BUILD.md`.
+not part of any wheel. See `platform/BUILD.md`. The same configure also builds
+the library families' C++ mirrors (`library/cpp/` → `librocke_library.a`, beside
+`librocke_core.a`) and, with `ROCKE_BUILD_PYBIND=ON`, their `rocke_library_engine`
+binding beside `rocke_engine`; `-DROCKE_BUILD_LIBRARY_CPP=OFF` skips them.
+`cmake --install` puts the library archive and headers into the same `ckc`
+package: a C/C++ consumer of a library family (e.g. the conv API in
+`rocke/instance_conv_*.h`) links `ckc::rocke_library`, which brings
+`ckc::rocke_core` with it.

dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md

@@ -67,7 +67,13 @@ The two invariants that gate every process below:
    2D/3D bodies) exactly like every other family, and additionally carries
    `.py`/`.c` emitter pairs under
    [library/tests/parity/](library/tests/parity/) that are byte-compared per
-   config.
+   config. Convolution (`conv_*`, `img2col`, `deep_fused_conv_pool`) keeps its
+   C++ mirror in the library instead: [library/cpp/](library/cpp/)
+   (`rocke_library_core`, bound into `rocke_library_engine`, bridged by
+   `kernels/conv/engine_bridge.py`), gated by the same `check_byte_identity.py`.
+   New library families put theirs there too, so the library owns both halves.
+   Library families' byte-identity golden is
+   [library/tests/parity/golden/](library/tests/parity/golden/).
 2. **One-way dependency.** `library → platform` only. `platform` never imports
    `kernels` / `builders` / `dispatch`, and stays standalone-installable.
    Anything reusable across families belongs in platform (see
@@ -531,7 +537,12 @@ Python emitter and the compiled C emitter for each config index and diffs the
    [rocke_representative_ir_sha256/](platform/tests/golden/rocke_representative_ir_sha256/),
    one `<family>.json` shard per family; `--only <family>[,...]` rewrites just
    those shards, so unrelated changes do not conflict in the golden);
-   library goldens re-bless with `--write` on their own `library/tests/test_*_golden.py`.
+   library goldens re-bless with `--write` on their own `library/tests/test_*_golden.py`
+   (the convolution families -- implicit-GEMM fwd/wgrad/dgrad, deep fused conv,
+   direct conv, Winograd -- live in
+   [test_conv_ir_golden.py](library/tests/test_conv_ir_golden.py) and its
+   [conv_representative_ir/](library/tests/golden/conv_representative_ir/) shards,
+   not the platform harness).
 3. Re-run the gate GREEN at **every** flavor in `LLVM_FLAVORS` before considering it done.
 4. A default-OFF, unwired knob must produce **no** golden diff — a clean gate
    with no re-bless is the proof it is inert.

dnn-providers/hip-kernel-provider/rocke/TESTING.md

@@ -336,7 +336,11 @@ and by both engines is green in both. The **emitted-IR validity gate**,
 question: it lowers every case in the representative corpus
 ([`rocke_ir_parity_harness.cases()`](platform/tests/instances/rocke_ir_parity_harness.py),
 the same corpus §4.2 hashes) and pushes each module through `clang` to a linked
-hsaco.
+hsaco. Library families keep their representative cases beside their own golden
+(convolution: [`library/tests/test_conv_ir_golden.py`](library/tests/test_conv_ir_golden.py));
+`--cases <file>` points this gate, and the arch-domain gate below, at such a
+corpus, and [`library/tests/run_all.py`](library/tests/run_all.py) runs both over
+every library corpus.
 
 Three design points, each of them load-bearing:
 
@@ -513,8 +517,9 @@ and selected library host/GPU suites. Build-tree C++ test registration is
 orchestrated by
 [`platform/tests/CMakeLists.txt`](platform/tests/CMakeLists.txt) and owned by its
 [concern-specific fragments](platform/tests/README.md#cmake-registration-ownership)
-under `platform/tests/cmake/`. Installed Python and library CTest registration
-remains in [`platform/CMakeLists.txt`](platform/CMakeLists.txt). The provider's
+under `platform/tests/cmake/`. Installed Python CTest registration remains in
+[`platform/CMakeLists.txt`](platform/CMakeLists.txt), which includes the library
+suites' own list, [`library/tests/ctest_lists.cmake`](library/tests/ctest_lists.cmake). The provider's
 external test-category YAML and TheRock's tier selection determine which
 registered tests actually execute. Check the selected list in the job log or
 with CTest's verbose listing. Registration alone does not establish CI coverage.

dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_implicit_gemm_conv.py

@@ -165,7 +165,7 @@ def _conv_split_k_override() -> int | None:
 # physical unit -- only ratios between degrees are compared -- so a constant's
 # absolute magnitude says nothing about its influence. Scope: gfx950, forward;
 # refit before relying on it elsewhere. Method and fitting tooling are described
-# in examples/gfx950/conv_fwd/fwd_merged_groups_case_study.md.
+# in library/builders/gfx950/conv_fwd/fwd_merged_groups_case_study.md.
 _FWD_MERGE_DEGREES = (1, 2, 4, 8, 16, 32, 64)
 _FWD_MERGE_CONSTS = (
     # Cost of issuing the A load per K step, divided by the achieved vector
@@ -5931,7 +5931,7 @@ def _split_k_degrees(
 
     ``direction="fwd"`` serves the split-K forward binary: the resolver's
     degree is the forward policy's
-    (:func:`rocke.helpers.split_k.select_split_k_fwd`), added only where it
+    (:func:`kernels.common.conv_split_k.select_split_k_fwd`), added only where it
     splits -- where it does not, the plain forward binary is the one to time --
     and a degree is also dropped when its grid exceeds
     ``_FWD_SPLIT_K_MAX_CTAS`` workgroups or its fp32 workspace passes the i32

dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/split_k_cost_fit.py

@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: MIT
 """Measure split-K degree curves and fit the split-K cost model to them.
 
-The constants of :class:`rocke.helpers.split_k_planner.SplitKCostModel` are
+The constants of :class:`kernels.conv.split_k_planner.SplitKCostModel` are
 fitted, not derived. This tool is how they are produced, so they can be
 refreshed after a kernel change and fitted for a new arch:
 
@@ -42,7 +42,7 @@
     ``SplitKCostModel(...)`` literal are then fitted on every shape.
 
 ``margin``
-    Measures :attr:`~rocke.helpers.split_k_planner.SplitKCostModel.ladder_margin`
+    Measures :attr:`~kernels.conv.split_k_planner.SplitKCostModel.ladder_margin`
     on a sweep that also timed degrees off the power-of-two ladder
     (``sweep --degrees 2-160,...``)::
 
@@ -71,7 +71,7 @@
 from pathlib import Path
 
 import numpy as np
-from rocke.helpers.split_k_planner import (
+from kernels.conv.split_k_planner import (
     MAX_GRID_DIM_Z,
     SplitKKernel,
     SplitKProblem,
@@ -463,7 +463,7 @@ def _feasible_times(m: _Kernel, max_degree: int) -> dict:
 
 def _decisions(kernels: list[_Kernel], cost, max_degree: int = 1024) -> list[_Decision]:
     """One :class:`_Decision` per main loop with more than one reduction."""
-    from rocke.helpers.split_k_planner import predict_split_k_cost
+    from kernels.conv.split_k_planner import predict_split_k_cost
 
     groups: dict[tuple, list] = {}
     for k in kernels:
@@ -528,7 +528,7 @@ def _legacy_decision_regret(decisions_kernels, max_degree: int = 1024) -> np.nda
     """The previous rule: the packed atomic whenever built, else the workspace,
     at the CK degree -- timed at the best binary of that kind. A CK degree of 1
     is the unsplit kernel; a decision without a timed unsplit one is skipped."""
-    from rocke.helpers.split_k import select_split_k_wgrad
+    from kernels.common.conv_split_k import select_split_k_wgrad
 
     r = []
     for members in decisions_kernels:
@@ -601,7 +601,7 @@ def _sig(v: float) -> float:
 
 
 def _model_of(x: dict, max_degree: int = 1024):
-    from rocke.helpers.split_k_planner import SplitKCostModel
+    from kernels.conv.split_k_planner import SplitKCostModel
 
     return SplitKCostModel(
         num_cus=int(x["num_cus"]),
@@ -620,8 +620,9 @@ def _model_of(x: dict, max_degree: int = 1024):
 
 
 def cmd_fit(args: argparse.Namespace) -> int:
-    from rocke.helpers.split_k import _ARCH_NUM_CUS, select_split_k_wgrad
-    from rocke.helpers.split_k_planner import predict_split_k_cost, split_k_cost_model
+    from kernels.common.conv_split_k import select_split_k_wgrad
+    from kernels.conv.split_k_planner import predict_split_k_cost, split_k_cost_model
+    from rocke.helpers.split_k import ARCH_NUM_CUS
 
     kernels = _load(args.logs, args.fill_logs)
     if not kernels:
@@ -633,7 +634,7 @@ def cmd_fit(args: argparse.Namespace) -> int:
     # but its hardware facts are its own.
     start = split_k_cost_model(args.arch)
     if start is None:
-        num_cus = args.num_cus or _ARCH_NUM_CUS.get(args.arch.split(":", 1)[0])
+        num_cus = args.num_cus or ARCH_NUM_CUS.get(args.arch.split(":", 1)[0])
         if num_cus is None:
             print(
                 f"unknown CU count for {args.arch!r}: pass --num-cus", file=sys.stderr
@@ -793,7 +794,7 @@ def cmd_margin(args: argparse.Namespace) -> int:
     best binary at its chosen degree). The recommendation is the margin with
     the best shape geomean whose worst shape loses at most ``--max-worse``.
     """
-    from rocke.helpers.split_k_planner import select_split_k_plan, split_k_cost_model
+    from kernels.conv.split_k_planner import select_split_k_plan, split_k_cost_model
 
     cost = split_k_cost_model(args.arch)
     if cost is None:

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx950/conv_dgrad/run_one_dgrad.py → dnn-providers/hip-kernel-provider/rocke/library/benchmarks/gfx950/conv_dgrad/run_one_dgrad.py

@@ -18,16 +18,17 @@
 ``--print-name-only`` builds and compiles but never touches the GPU queue, which
 is how you get the ``--kernel-regex`` for a capture without perturbing it.
 
-Usage:
-    python3 run_one_dgrad.py --kouter on --stride 2
-    python3 run_one_dgrad.py --tile-m 128 --print-name-only
-    python3 run_one_dgrad.py --tap-outer off   # the flat k_dg loop at stride 1
+Usage (from ``rocke/``, with ``PYTHONPATH=$(pwd)/platform/python:$(pwd)/library``):
+    python3 -m benchmarks.gfx950.conv_dgrad.run_one_dgrad --kouter on --stride 2
+    python3 -m benchmarks.gfx950.conv_dgrad.run_one_dgrad --tile-m 128 --print-name-only
+    python3 -m benchmarks.gfx950.conv_dgrad.run_one_dgrad --tap-outer off   # the flat k_dg loop at stride 1
 
     ROCPROF_TRACE_DECODER_LIB=<dir with librocprof-trace-decoder.so> \\
-    python3 ../../../../../dsl_docs/optimization/utilities/tools/wavescope/\\
+    python3 platform/dsl_docs/optimization/utilities/tools/wavescope/\\
 capture_wavescope_trace.py --output-dir ./att_out \\
         --kernel-regex "rocke_trace_dgrad.*kouter" \\
-        -- python3 run_one_dgrad.py --kouter on --stride 2 --warmup 1 --iters 2
+        -- python3 -m benchmarks.gfx950.conv_dgrad.run_one_dgrad \\
+           --kouter on --stride 2 --warmup 1 --iters 2
 """
 
 import argparse

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx1250/conv_kouter/gfx1250_transpose_read_case_study.md → dnn-providers/hip-kernel-provider/rocke/library/builders/gfx1250/conv_kouter/gfx1250_transpose_read_case_study.md

@@ -81,7 +81,7 @@ formula and the gate stayed GREEN throughout.
 What settled it was a **per-lane hardware probe**: fill LDS with an encoded
 `(k, col)` value, issue the transpose read, and decode what each lane actually
 received. One run answers the question that no amount of source reading did.
-`examples/gfx1250/wmma_probe.py` is the existing precedent for this technique —
+`platform/python/rocke/examples/gfx1250/wmma_probe.py` is the existing precedent for this technique —
 note its own docstring calls the gfx1250 lane maps a "hypothesis verified
 empirically", so the layout this builds on carries the same caveat.
 
@@ -116,22 +116,22 @@ nothing*, which is why the K-outer sweeps now carry an `_assert_ran()` backstop.
 
 ## Replay
 
-From `platform/`, with `PYTHONPATH=$(pwd)/python`. Items marked **[gfx1250]**
+From `platform/`, with `PYTHONPATH=$(pwd)/python:$(pwd)/../library`. Items marked **[gfx1250]**
 need the hardware; the rest are CPU-only and run anywhere.
 
 Correctness **[gfx1250]** — arch-aware, so on a wave32 box these select the
 `16x16x32` atom and the `default` epilogue automatically:
 
 ```bash
-python3 -m pytest tests/instances/test_conv_wgrad_correctness.py -q -rs -k k_outer
-python3 -m pytest tests/instances/test_conv_dgrad_correctness.py -q -rs -k KOuter
+python3 -m pytest ../library/tests/test_conv_wgrad_correctness.py -q -rs -k k_outer
+python3 -m pytest ../library/tests/test_conv_dgrad_correctness.py -q -rs -k KOuter
 ```
 
 The newly-reachable 16-bit-dW path from Finding 3.2 — read the `-rs` skip list
 and confirm `epilogue=default` subTests **run** rather than skip:
 
 ```bash
-python3 -m pytest tests/instances/test_conv_wgrad_correctness.py -q -rs \
+python3 -m pytest ../library/tests/test_conv_wgrad_correctness.py -q -rs \
     -k "pipeline_mem or pipeline_compv3 or pipeline_compv4 or pipeline_basic"
 ```
 
@@ -140,7 +140,7 @@ emits `ds_load_tr16_b128` with the right operand shapes, and that the wavelet
 combination is rejected:
 
 ```bash
-python3 -m pytest tests/instances/test_conv_dgrad_correctness.py -q \
+python3 -m pytest ../library/tests/test_conv_dgrad_correctness.py -q \
     -k Gfx1250Emit
 ```
 
@@ -155,7 +155,7 @@ ROCKE_LLVM_FLAVOR=llvm22 python3 tools/check_byte_identity.py --only conv
 Step 0 lever sweep **[gfx1250]**:
 
 ```bash
-python3 python/rocke/benchmark/benchmark_implicit_gemm_conv.py \
+python3 ../library/benchmarks/common/benchmark_implicit_gemm_conv.py \
     --direction wgrad --arch gfx1250 --dtype bf16 \
     --miopen-cmd "./MIOpenDriver convbfp16 -n 8 -c 128 -H 32 -W 32 -k 128 \
         -y 3 -x 3 -p 1 -q 1 -u 1 -v 1 -l 1 -j 1 -g 1 -F 4 -in_layout=NHWC" \

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx950/conv_dgrad/README.md → dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/conv_dgrad/README.md

@@ -83,5 +83,5 @@ python3 -m benchmarks.common.benchmark_direct_conv \
 `--variants` takes a comma-separated list of `DIRECT_CAPABILITIES` keys
 (`direct_depthwise_dgrad_win`, `direct_grouped_dgrad_4c`,
 `direct_grouped_dgrad_fused`, ...). The standalone implicit-GEMM driver for
-traces is `run_one_dgrad.py` in this folder (`--kouter`, `--tap-outer`,
+traces is [`run_one_dgrad.py`](../../../benchmarks/gfx950/conv_dgrad/run_one_dgrad.py) (`--kouter`, `--tap-outer`,
 `--uniform-k` default to what the AOT cache builds for the shape).

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/examples/gfx950/conv_dgrad/depthwise_windowed_dgrad_case_study.md → dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/conv_dgrad/depthwise_windowed_dgrad_case_study.md

@@ -9,18 +9,18 @@ Measured numbers are deliberately absent — see `platform/AGENTS.md` §Complian
 What is recorded here is mechanism, instruction-count evidence, the levers that
 were tried (kept and reverted), what the AOT port changed, and the replay path.
 The knob reference is in
-[`dsl_docs/instances/convolution.md`](../../../../../dsl_docs/instances/convolution.md),
+[`dsl_docs/instances/convolution.md`](../../../kernels/conv/docs/convolution.md),
 section "Windowed Depthwise Dgrad Kernel".
 
 ## Kernel and where it lives
 
 `DirectDepthwiseDgradWindowedSpec` / `build_direct_depthwise_dgrad_windowed`
 (VALU and `dot2` forms) and `_build_dw_dgrad_win_toeplitz` (Toeplitz MFMA form) in
-[`conv_direct_grouped.py`](../../../../../../library/kernels/common/conv_direct_grouped.py);
-C++ mirror [`conv_direct_grouped_build_depthwise_dgrad_win.cpp`](../../../../../cpp/instances/common/conv_direct_grouped_build_depthwise_dgrad_win.cpp).
+[`conv_direct_grouped.py`](../../../kernels/common/conv_direct_grouped.py);
+C++ mirror [`conv_direct_grouped_build_depthwise_dgrad_win.cpp`](../../../cpp/instances/common/conv_direct_grouped_build_depthwise_dgrad_win.cpp).
 AOT cache variant `direct_depthwise_dgrad_win` (`DW_WIN_*`,
 `_dw_win_policy_reason` in
-[`direct_aot_jobs.py`](../../../../../../library/kernels/conv/direct_aot_jobs.py));
+[`direct_aot_jobs.py`](../../../kernels/conv/direct_aot_jobs.py));
 not dispatched yet (the grouped-convolution dispatcher is being reworked
 separately), so it is reached through the cache (`--run-from-cache`).
 Scope: `cpg = kpg = 1`, stride 1, fp16 / bf16 (implicit GEMM declines