feat(rocke): split-K forward implicit-GEMM convolution on gfx1250 by jakpiase · Pull Request #14393 · ROCm/rocm-libraries

jakpiase · 2026-10-26T15:24:26Z

Summary

Small-M forward layers leave the WMMA implicit GEMM on gfx1250 with too few workgroups to fill the device.

This PR adds split-K to the forward implicit-GEMM convolution, in both engines and in the AOT cache.

JIRA ID: AICK-2778

Changes

Two-stage split-K

Each workgroup stores an fp32 partial tile into its workspace slab.
The new conv_fwd_split_k_reduce family folds the slabs.
The degree is a kernarg. No workspace clearing is needed, and the result is deterministic.

Single-launch split-K (split_k_cluster): the slices of one output tile form a gfx1250 workgroup cluster and fold through a cluster barrier.

Auto-deduced degree (split_k = -1)

resolve_conv_split_k("fwd", ...) resolves the forward degree, as it does for wgrad and dgrad. The policy is rocke.helpers.split_k.select_split_k_fwd:
- it splits only an unsplit grid under half of 8 workgroups per CU, over a reduction at least 512 deep;
- the degree is a power of two toward that target, at most 32, with at least 2 K tiles per slice;
- the result is held to the gridDim.z budget and the i32 workspace.
It resolves to 1 on targets other than gfx1250, on pointwise and 3-D problems, and for specs the split form cannot take (a fused bias / scale / ReLU / clamp epilogue, among others), so those still build the plain kernel. The AOT cache shares the arch list.
The builder resolves -1 before validation. lower_conv_implicit_gemm also resolves it before choosing an engine, because the C++ port takes explicit degrees only, as its wgrad port does.
The partials are fp32 slabs, so no 16-bit accuracy cap applies.

ABI, AOT and runtime

The forward ABI and ConvArgs gain the split-K variant.
The AOT grid builds the split-K and cluster binaries.
KernelIdentity gains split_k_cluster; it is hash-neutral at 0.
KernelIdentity.needs_stage2 (two_stage and not split_k_cluster) is the single predicate for launching Stage 2; two_stage means "takes the split-K workspace ABI", which both forward forms need.
The runtime keeps the cluster shape on pre-packed launches. _enqueue and prepare_launch share one cluster launch helper, Runtime._launch_cluster.

Benchmark (benchmark_implicit_gemm_conv --run-from-cache)

--split-k means the same for forward as for wgrad:
- -1 (default) launches a cached split-K binary at the resolved degree, and skips it where the policy does not split;
- 0 sweeps --split-k-degrees plus that degree;
- n pins n.
Binds both split-K stages once per timed loop.
Runs patch-embedding and 2x2 stride-2 forwards width-folded.

split_k = 1 emits the same kernels as before.

Both engines

The split-K implicit GEMM, its cluster form, the forward two-stage ABI and conv_fwd_split_k_reduce are mirrored in the C++ engine.
Two existing deviations in the C++ fp32 CShuffle store are fixed so it matches Python.
The C++ backend now forwards num_load_waves.

Parity: configs cover:

two-stage split-K: mem / wavelet, cshuffle / default, grouped;
cluster split-K with S = 4 / 8 / 12 / 15;
the ABI lists and the reduce kernel.

The representative IR golden only gains entries.

Review follow-ups

Split-K rejects non-identity epilogues (bias / ReLU / clamp), split_k <= 0 and a NULL dtype_d in both engines.
Cluster split-K requires K and kpg to be multiples of 4.
Split-K forward AOT jobs are gfx1250-only; cluster jobs also need llvm23.
The forward split-K degree is auto-deduced (see above).

Develop merges: earlier conflicts were in the representative IR golden, re-blessed with rocke_ir_parity_harness.py --write on each merge result (develop's cases plus this PR's split-K cases).
The latest merge collided with develop's new conv_implicit_gemm parity configs 36-39, so this PR's split-K parity configs are now 40-53 (both emitters). The differential golden llvm_gfx_all.json was re-recorded for conv_implicit_gemm only: develop's 36-39 plus this PR's hashes, unchanged, moved from 36-44 to 40-48.

The Winograd forward merge conflicted in platform/CMakeLists.txt, where both test lists are kept, and in the representative IR golden, which was re-blessed on the merge. Every case hash matches one side.

The split-K wgrad / LLVM24-lane merge conflicted in the conv_split_k.py module docstring, where both sides' bullets are kept, and in the representative IR golden, again re-blessed on the merge; every case hash matches one side.

The wave32 WMMA direct and small-channel forward conv merge (#14392) conflicted in platform/CMakeLists.txt and rocke_engine.cpp (both sides' test and header kept) and in the representative IR golden, re-blessed on the merge (311 ok cases and 19 expected failures at llvm20 / 22, 330 ok at llvm23); every case hash matches one side. This PR adds no llvm23-only byte-identity family: its cluster configs live in conv_implicit_gemm, which every flavor runs.

The strided, dilated and packed-group WMMA direct dgrad merge (#14385) conflicted only in the representative IR golden, re-blessed on the merge (311 ok cases and 25 expected failures at llvm20 / 22, 336 ok at llvm23); cases and expected failures are the union of both sides, and every case hash matches one side. platform/CMakeLists.txt, rocke_engine.cpp, SUPPORT_MATRIX.md and the parity harness merged cleanly with both sides kept.

The newest develop merge (7e3d78a3d1, which adds gfx942 f64 GEMM among other changes) conflicted only in the representative IR golden's summary counters, re-blessed on the merge (313 ok cases and 25 expected failures at llvm20 / 22, 338 ok at llvm23); every case hash matches one side.

The wave32 WMMA stem dgrad merge (#14387, develop 311cb32338) conflicted only in the representative IR golden (expected failures and summary counters), re-blessed on the merge (313 ok cases and 29 expected failures at llvm20 / 22, 342 ok at llvm23); every case hash matches one side and nothing is dropped. rocke_engine.cpp, backend.py, platform/CMakeLists.txt, SUPPORT_MATRIX.md, the parity harness and the differential golden merged cleanly with both sides kept, and check_byte_identity.py is develop's.

Docs: support matrix, instance index, convolution docs and the knob catalog.

Testing

check_byte_identity.py is GREEN at llvm20, llvm22 and llvm23. Every commit passes the gate.
CPU tests for validity, naming, ABI, AOT grid, launch gating, and the auto-degree policy and resolver: test_conv_fwd_split_k.py, registered in CTest, plus test_cluster_launch.py.
gfx1250 GPU numerics for both forms, in bf16 and fp16, and for an auto spec at its resolved degree.
After the latest develop merge: rocke_ir_parity_harness.py --check passes at all three flavors (and the llvm24 alias).
Also after that merge, run_diff.py --python-only --check-golden (ll at all three flavors, and ir_canonical) gives GOLDEN OK.
Width-folded --run-from-cache --verify on a 16x16 stride-16 and a 2x2 stride-2 forward.

Small-M forward layers (a few thousand output pixels over a reduction of 2-17k) leave the WMMA implicit GEMM with a handful of workgroups, each walking the whole reduction. Add a split-K forward: - ImplicitGemmConvSpec.split_k > 1 builds the split-K binary. The degree and slice width are the ks_count / ks kernargs (as in wgrad), so one binary serves every degree. Each workgroup reduces one slice and stores its fp32 partial tile into slab `slice` of a [ks_count][M][K] workspace; every slab element is written once, so the workspace needs no clearing and the result is deterministic. WMMA only (mem / wavelet pipelines); async_dma, unroll_k, chiplet swizzle, group merge, pointwise and 3-D are rejected by the validator. - conv_fwd_split_k_reduce: Stage 2 folds the slabs and casts into D, vectorised, with the slab loop unrolled for the degree it is built for (a runtime loop stays available). - The fwd ABI gains the split-K variant (two_stage on fwd): ws_ptr, ws_bytes, ks, ks_count after the problem block. ConvArgs fills them. - AOT grid: split-K fwd kernels on small WMMA tiles (CACHE_FWD_SPLIT_K_*), recorded as split_k = 2 with two_stage. - benchmark_implicit_gemm_conv --run-from-cache times split-K binaries at degrees 2..32 (capped by the K tiles, the i32 workspace and a 16k workgroup grid), Stage 1 + Stage 2 together. Both engines. The C++ engine mirrors all of it: split_k on rocke_implicit_gemm_conv_spec (name tag, admission gate rocke_conv_fwd_split_k_available, fp32 C tile in the LDS budget, the trailing kernargs, the z = group * ks_count + slice decode, slice-bounded mem / wavelet K loops and the fp32 slab epilogue), the fwd two-stage list in the C++ conv_abi, and a new conv_fwd_split_k_reduce instance, header and binding. Two C++ deviations in the fp32 CShuffle store are fixed to match Python: the C tile is read with the generic vector LDS load, and the 4-element fp32 store cap is applied. The binding and the backend spec dict forward split_k and num_load_waves. Parity: conv_implicit_gemm_emit 36-39 (wavelet / mem cshuffle / mem default / grouped), conv_abi_emit 13 (now the split-K list), 17 (3-D) and 18 (dgrad two_stage, rejected), and a new conv_fwd_split_k_reduce_emit pair (0-9). Goldens: new representative cases for the split-K forward and the reduce kernel; llvm_gfx_all gains the new entries and re-blesses conv_abi 13. No existing hash changes: split_k = 1 emits the same kernel as before. test_conv_fwd_split_k.py joins the conv CTest lane; support matrix, instance index, convolution.md and the runbook knob catalog document the knob and the Stage 2 family. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

KernelLauncher.bind() packs a launch once for timed loops, but Runtime.prepare_launch had no cluster path, so a launcher compiled with cluster_dims was bound as a plain hipModuleLaunchKernel. The kernel then ran as 1x1x1 clusters and its cluster barrier synchronised nothing. prepare_launch takes the cluster shape, validates it against the grid up front, and enqueues through hipDrvLaunchKernelEx like launch() does; bind() passes the launcher's resolved cluster. Plain launches are unchanged. The bind test's fake runtime takes the new argument and checks both an unclustered and a clustered launch. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The split-K forward run called the Stage 1 and Stage 2 launchers with their argument dicts, repacking the kernargs on every timed iteration. For the small layers split-K serves that host work is longer than the kernels, so the benchmark timed the host. Bind both stages once, as the single-kernel path already does. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

ImplicitGemmConvSpec.split_k_cluster = S (2..15, mem pipeline) puts the S K-slices of one output tile into a (1, 1, S) gfx1250 workgroup cluster. Each slice stores its fp32 partial tile into its workspace slab as before, then a cluster barrier makes the slabs visible and each slice folds 1/S of the tile's rows across the S slabs and stores them cast into D. The Stage 2 launch goes away. The cluster barrier's fences are cluster scope: s_wait_storecnt and a global_inv at SE scope, with no L2 writeback, unlike the agent-scope release that made an in-launch fold with arrival counters impractical. The degree is baked (one binary per S) and the default (0) emits none of this. Both engines: the C++ engine gains split_k_cluster (skc<S> name tag, admission checks, the cluster_dims attribute and the cluster fold), and the binding / backend spec dict forward it. Parity: conv_implicit_gemm_emit 40-44 (S = 4 / 8 / 15 on cshuffle, S = 4 grouped, S = 12 on the default epilogue; compared in the llvm23 lane, both engines reject them at llvm20 / llvm22) and 45 (split-K on MFMA, rejected). The representative golden gains a cluster case. Support matrix, convolution.md and the runbook knob catalog document the knob. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The single-launch cluster split-K forward (split_k_cluster) bakes its degree as the (1, 1, S) cluster shape, so the AOT cache gets one binary per degree. The identity gains split_k_cluster (left out of the hash at 0, so existing entries keep their names) and the label a skc<S> tag. The forward grid builds S in {4, 8, 12, 15} on the mem / cshuffle 16x32, 32x32, 64x32, 64x64 and 128x128 tiles (tile_k 64, 16-byte vectors). --run-from-cache launches them with their cluster through the pre-packed launch, with no Stage 2, and only offers them to problems with fewer than 256 output tiles, a reduction of at least 1500 and 4-channel output quads, where the single launch pays. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

fold_width folds f adjacent W positions into the channels when f divides the filter width, the W stride, the input width and the W padding. The identity does not depend on the direction, so a forward conv on the folded problem computes the original output from the same tensors, with no copy. For a patch embedding (stride equal to the filter, no padding) the implicit GEMM A loads go from C-element gathers (C = 3, vector width 1) to X * C contiguous channels (48 for a 16x16 patch, vector width 8). The 2x2 stride-2 downsample folds as well. The benchmark's --run-from-cache now launches forward problems folded, as it already did for dgrad. References and FLOP counts stay the original problem's. The dispatcher's _gemm_problem still folds dgrad only. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The GPU numerics of test_conv_fwd_split_k copy through pinned host buffers: a pageable copy registers the host range with the driver, which must not overlap a kernel in flight. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

therock-pr-bot · 2026-10-26T15:30:38Z

✅ 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?

Copilot

🟡 Changes recommended

Unresolved buffer-boundary and reduction-semantics defects can cause memory corruption or incorrect convolution results.

9 open issues

What changed in this PR

Adds split-K forward implicit-GEMM convolution to rocKE for small-output workloads on gfx1250, with Python/C++ implementations and ahead-of-time cache integration.

Changes:

Adds two-stage workspace reduction and single-launch cluster reduction.
Extends kernel arguments, cache identities and prepared launches; updates cached benchmarking and forward width folding.
Adds numerical, admission, parity and golden coverage, plus support documentation.

Paths are relative to dnn-providers/hip-kernel-provider/rocke/.

File	Description
`platform/tests/test_rocke.py`	Checks cluster propagation through bound launches.
`platform/tests/runtime/test_cluster_launch.py`	Tests prepared cluster launches and validation.
`platform/tests/instances/rocke_ir_parity_harness.py`	Adds representative split-K build cases.
`platform/tests/instances/differential/golden/llvm_gfx_all.json`	Updates ABI and split-K differential baselines.
`platform/tests/golden/rocke_representative_ir_sha256.json`	Adds split-K hashes and expected flavor failures.
`platform/python/rocke/runtime/launcher.py`	Preserves cluster shape when binding launches.
`platform/python/rocke/runtime/hip_module.py`	Supports prepared cluster launches.
`platform/python/rocke/instances/SUPPORT_MATRIX.md`	Documents split-K support and limitations.
`platform/python/rocke/core/backend.py`	Forwards split-K and loader-wave fields.
`platform/dsl_docs/optimization/optimization_runbook.md`	Documents split-K tuning knobs.
`platform/dsl_docs/instances/index.md`	Registers the reducer and forward ABI extension.
`platform/dsl_docs/instances/convolution.md`	Explains workspace and cluster reduction.
`platform/cpp/instances/common/conv_implicit_gemm_conv_spec_descriptors.cpp`	Adds defaults, naming, admission and workspace LDS sizing.
`platform/cpp/instances/common/conv_implicit_gemm_conv_kloop_drivers.cpp`	Uses split-specific K-loop bounds.
`platform/cpp/instances/common/conv_implicit_gemm_conv_epilogue.cpp`	Stores partial slabs and implements cluster folding.
`platform/cpp/instances/common/conv_implicit_gemm_conv_build_glue.cpp`	Builds split arguments, slice bounds and workspace addressing.
`platform/cpp/instances/common/conv_fwd_split_k_reduce.cpp`	Implements the C++ second-stage reducer.
`platform/cpp/instances/common/conv_abi.cpp`	Extends forward argument lists.
`platform/cpp/include/rocke/instance_conv_implicit_gemm.h`	Exposes split-K spec fields and admission API.
`platform/cpp/include/rocke/instance_conv_implicit_gemm_internal.h`	Adds internal split-K state and epilogue parameters.
`platform/cpp/include/rocke/instance_conv_fwd_split_k_reduce.h`	Declares the reducer spec and public API.
`platform/cpp/include/rocke/instance_conv_abi.h`	Documents the extended forward ABI.
`platform/cpp/helpers/epilogues.cpp`	Aligns fp32 CShuffle load emission with Python.
`platform/cpp/bindings/rocke_engine.cpp`	Parses new fields and exposes reducer bindings.
`platform/CMakeLists.txt`	Registers split-K tests for installation.
`library/tests/test_conv_wgrad_prefetch_addr.py`	Preserves historical identity-hash coverage.
`library/tests/test_conv_fwd_split_k.py`	Adds admission, cache, launch and GPU numerical tests.
`library/tests/test_conv_dgrad_spec_policy.py`	Accounts for the hash-neutral cluster field.
`library/tests/test_conv_dgrad_patch.py`	Tests forward width-fold equivalence.
`library/tests/test_builder_signature_contract.py`	Registers the reducer signature contract.
`library/tests/parity/conv_implicit_gemm_emit.py`	Adds Python split-K parity configurations.
`library/tests/parity/conv_implicit_gemm_emit.c`	Mirrors split-K parity configurations in C.
`library/tests/parity/conv_fwd_split_k_reduce_emit.py`	Adds the Python reducer parity emitter.
`library/tests/parity/conv_fwd_split_k_reduce_emit.c`	Adds the C reducer parity emitter.
`library/tests/parity/conv_abi_emit.py`	Covers forward split-K argument lists.
`library/tests/parity/conv_abi_emit.c`	Mirrors forward ABI parity cases.
`library/kernels/conv/kernel_label.py`	Adds the cluster-degree label tag.
`library/kernels/conv/kernel_identity.py`	Adds a hash-neutral default cluster field.
`library/kernels/conv/aot_jobs.py`	Enumerates and counts split-K cache jobs.
`library/kernels/common/conv_implicit_gemm.py`	Implements Python split-K construction and cluster folding.
`library/kernels/common/conv_fwd_split_k_reduce.py`	Adds the Python reducer and workspace helpers.
`library/kernels/common/conv_args.py`	Packs forward split-K slice arguments.
`library/kernels/common/conv_abi.py`	Defines the extended forward ABI.
`library/benchmarks/common/kernel_sweep.py`	Supplies cached split-K launch arguments.
`library/benchmarks/common/benchmark_implicit_gemm_conv.py`	Adds split-K timing, workspace handling and forward folding.

_{Review effort: Balanced}

therock-pr-bot · 2026-10-26T15:45:46Z

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

Address review findings on the split-K forward: - Reject a non-identity acc_epilogue with split-K. Stage 1 stores raw partials and the reduction only adds and casts them, so a bias was added once per slice and ReLU / clamp ran before the sum. This matches how the wgrad stream-K path treats partial sums. - Validate every non-default split_k (split_k != 1), so 0 and negative values reach the ">= 1" rejection instead of building an unsplit kernel. - The cluster fold reads 4-channel quads bounded by their first channel. Require K and kpg to be multiples of 4 on the build problem and on every launched problem (is_valid_spec_for_problem, identity_supports_problem, the cached-launch gate), through one shared predicate. - Reject the cluster form on targets without workgroup-cluster launch. - Build split-K forward AOT jobs only for gfx1250, and the cluster jobs only at the LLVM flavor that lowers clusters; count_jobs agrees. - --run-from-cache --split-k 0 sweeps --split-k-degrees for the forward binary as it does for wgrad; -1 keeps the forward ladder. - C++ Stage 2 reduce: reject a NULL dtype_d like the Python validator instead of defaulting it to bf16, and zero the builder before any early return so a name that does not fit is reported, not read uninitialised. C++ mirror updated alongside; parity configs 46-49 (conv_implicit_gemm) and 10 (conv_fwd_split_k_reduce) are rejected by both engines. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…es once - conv_implicit_gemm_spec_to_dict reads split_k, split_k_cluster and num_load_waves from the spec instead of defaulting missing ones, so a malformed spec fails at the Python-to-C++ boundary rather than building an unsplit, cluster-free kernel. - Runtime.prepare_launch validates the cluster shape, queries the device and builds the hipDrvLaunchKernelEx config once; the returned callable only enqueues it. Runtime._enqueue shares the same config builder. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…reduce dtype - conv_implicit_gemm_spec_to_dict forwards the accumulator epilogue and the rocke_engine conv_implicit_gemm binding parses it. Before, the C++ path always built the identity epilogue, and so admitted split-K specs with a bias / ReLU / clamp epilogue that Python rejects. - The conv_fwd_split_k_reduce binding turns an explicit dtype_d=None into a NULL dtype, which the C++ validator rejects like Python, instead of falling back to bf16. Only an absent key keeps the default. The shared dict parser is unchanged. Engine-agreement tests reach both engines through the backend. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

bartekxk · 2026-10-26T23:02:56Z

Review: PR #14393 — split-K forward implicit-GEMM convolution on gfx1250

#14393

Title: feat(rocke): split-K forward implicit-GEMM convolution on gfx1250
Author: jakpiase
State: Open
Scope: 45 files, ~3,688 insertions, scoped to dnn-providers/hip-kernel-provider/rocke/

Summary

Adds a split-K variant of the WMMA forward implicit-GEMM conv for gfx1250,
for small-M/deep-K layers that otherwise underfill the device. Two forms:

Two-stage: Stage 1 (the plain WMMA K-loop run over one K-slice) stores
an fp32 partial tile into slab slice of a [ks_count][M][K] workspace
(every element written exactly once, no clearing, no atomics); Stage 2
(conv_fwd_split_k_reduce.py / the mirrored .cpp) sums the slabs and
casts into D.
Cluster (split_k_cluster): the S K-slices of one output tile form a
(1, 1, S) gfx1250 workgroup cluster; after cluster_barrier(), each slice
folds 1/S of the tile's rows by reading all S slabs and storing straight
into D — no Stage 2 launch.

Both variants are built symmetrically in the Python and C++ engines (new
conv_fwd_split_k_reduce.py/.cpp, and matching additions to
conv_implicit_gemm.py/conv_implicit_gemm_conv_*.cpp), with admission gates
(fwd_split_k_reason / rocke_conv_fwd_split_k_available,
fwd_split_k_cluster_problem_reason / rocke_conv_fwd_split_k_cluster_problem_ok)
duplicated line-for-line between engines, AOT cache entries
(CACHE_FWD_SPLIT_K_*), and GPU correctness tests for both variants across
several shapes (plain, grouped, strided, k-tail-not-a-multiple-of-tile).

I verified the core math (slab addressing, z = group*ks_count + slice
decoding, row-ownership in the cluster fold, K/kpg % 4 admission) is
consistent between the Python and C++ emitters, and that the two engines'
golden-IR / cache-key bookkeeping line up. One design wart stood out as a
latent correctness trap (not an active bug today).

Findings

1. Cluster split-K forward's `KernelIdentity.two_stage=True` is a landmine gated only by an incidental `not split_k_cluster` check

Files: library/kernels/conv/aot_jobs.py:775 (cluster BuildJob's
KernelIdentity(..., two_stage=True, **cfg) with cfg["split_k_cluster"]
set), library/benchmarks/common/benchmark_implicit_gemm_conv.py:5754

The cluster binary performs its own in-kernel fold and explicitly has "no
Stage 2 launch" (per the PR's own docstrings in
conv_implicit_gemm.py:633-642 and aot_jobs.py's CACHE_FWD_SPLIT_K_*
comment block). Yet its KernelIdentity sets two_stage=True — the same
flag the two-stage (Stage-1-then-Stage-2) variant uses, and the same flag
kernel_identity.py/wgrad's dispatch code elsewhere treats as "this binary's
reduction isn't finished until a second kernel runs."

The only place that currently prevents a double Stage-2 launch for a cluster
identity is one local guard in the benchmark harness:

if direction == "fwd" and ident.two_stage and not ident.split_k_cluster:
    ...  # launch Stage 2

(benchmark_implicit_gemm_conv.py:5754). Every other consumer of
ident.two_stage for fwd (_signature_for_identity, _extras_and_grid,
_grid_for) happens to want the same answer for both variants (both need
the split-K ABI/workspace/grid-z scaling), so today nothing breaks. But the
field itself no longer means "this binary needs a second kernel launched" —
it now means "this binary needs the split-K ABI," and only one call site
remembers to re-derive the first meaning by also checking
split_k_cluster. A future dispatch/benchmark code path (or a port of this
to dispatch/grouped_convolution.py, which this PR does not yet touch) that
checks ident.two_stage to decide whether to enqueue a reduce kernel, without
independently excluding split_k_cluster, will silently double-process a
cluster binary's already-correct output — launching Stage 2 against a
workspace/ABI the cluster kernel didn't populate the same way (Stage 2 reads
ws_ptr/ks_count and writes D; the cluster kernel already wrote D
itself), corrupting the output without any test catching it unless that new
call site happens to include the same incidental exclusion.

Fix: stop overloading two_stage for the cluster variant. Either give it
its own flag (e.g. needs_stage2 = two_stage and not split_k_cluster,
computed once and reused everywhere two_stage currently gates reduce-kernel
behavior) or keep two_stage=False on the cluster KernelIdentity and
introduce a separate uses_split_k_abi predicate (split_k > 1, covering
both forms) for the ABI/workspace/grid-z call sites that currently keyed off
two_stage.

2. Minor: `prepare_launch`'s new cluster branch duplicates `_enqueue`'s cluster logic instead of reusing it

File: platform/python/rocke/runtime/hip_module.py:871-888 vs. the
existing _enqueue static method (hip_module.py:663-692)

_enqueue already has a cluster is not None branch that builds the config
via _cluster_launch_config and calls _hipDrvLaunchKernelEx. The new
prepare_launch cluster path re-implements the same two steps inline
(cfg, attr = self._cluster_launch_config(...) then its own
_hipDrvLaunchKernelEx call) rather than routing through _enqueue, which
launch_blocking/launch_kernelparams do use. I checked the argument order
and params=None convention and they're consistent with _enqueue's cluster
branch today, so this isn't a bug, but it's a second copy of launch-call
plumbing that can drift the way this repo's docs warn about (see
pr14114_review.md finding #5) — e.g. if _enqueue's cluster branch later
gains a retry-on-error path or a different where string format, nothing
forces prepare_launch's copy to follow.

Fix: have prepare_launch's cluster path call _enqueue once to get the
(cfg, attr) keep-alive tuple, the way the non-cluster path already shares
call_args/_hipModuleLaunchKernel, instead of re-deriving the
hipDrvLaunchKernelEx call locally.

Overall assessment

No correctness blockers found. The slab addressing, group/slice z-decoding,
cluster row-ownership math, and the K % 4 / kpg % 4 cluster-fold admission
gate are all mirrored correctly and consistently between the Python and C++
engines, and are exercised by GPU numeric tests for both the two-stage and
cluster forms across plain/grouped/strided/k-tail shapes. Finding #1 is a
real but currently-inert landmine worth fixing before this flag's meaning is
relied on elsewhere (e.g. once forward split-K reaches
dispatch/grouped_convolution.py, which this PR does not touch); finding #2
is a cosmetic duplication.

cgmillette · 2026-10-27T01:16:11Z

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

PR #14393 — split-K forward implicit-GEMM convolution on gfx1250 — 🟢 First-round: no closer look needed — byte-identical in both engines, every shipped gfx1250 split-K job compiles with ROCm clang 24 without spills, and Copilot's nine items are fixed; a golden-only rebase is due, and bartekxk's notes are open

Prior item	Status
Copilot (10-26) ×9: split-K with bias / ReLU / clamp epilogues; a channel-quad overrun; jobs beyond gfx1250 / llvm23; `--split-k-degrees` ignored; `split_k <= 0`; a NULL `dtype_d`; a builder freed before init; defaulted spec fields; cluster setup inside the timed loop	🟢 FIXED in `24177571` / `906cbfc5` / `57ce3193`, per the replies (read). Checked here: split-K forward jobs exist only for gfx1250, and cluster jobs only at llvm23 (ran), and both engines agree byte for byte
bartekxk (10-26): the cluster identity reuses `two_stage=True`, which only one benchmark guard (`two_stage and not split_k_cluster`) separates from "launch Stage 2"; `prepare_launch` duplicates `_enqueue`'s cluster branch	🟡 OPEN, no reply yet. We confirm the first (read): that guard is the only site that excludes cluster binaries, and `kernel_sweep` labels cluster jobs "two-stage". No dispatch path launches Stage 2 for forward today, so it is latent, as bartekxk says
bartekxk (10-26): what does split-K buy for forward, and is it used for auto-deduce?	🟡 OPEN: a question for the author
Bots / CI	🟡 The PR conflicts with develop in the representative golden only, so `57ce3193` got no full CI run (read). Regenerated on a scratch merge, the golden differs from develop's only in the PR's own split-K cases: 4 at llvm20 / 22, 5 at llvm23 (ran)
Our checks	🟢 • On that scratch merge: the golden gate passes at all three flavors, the native build and 32 ctests pass, and byte identity is GREEN at llvm20 / 22 / 23 (`conv_fwd_split_k_reduce` 11 configs, `conv_implicit_gemm` 50, 0 bad) (ran). • Independent compile check, ROCm clang 24 for gfx1250 at -O3: all 1,860 shipped split-K forward jobs compile with no scratch spills (ran). That is 910 two-stage and 20 cluster jobs each for bf16 and fp16. • Tests: at llvm23, the PR's and the touched test files fail only develop's pre-existing `test_off_keeps_the_historical_hash` (ran); N1 covers llvm22. • `rule_scan` is clean (ran)

#	Sev	file:line	Finding
N1	🟡	`library/tests/test_conv_fwd_split_k.py` (`test_grid_has_one_binary_per_degree`)	The test depends on the host's LLVM flavor. Cluster jobs exist only at llvm23, and the test enumerates the grid without pinning a flavor, so on an llvm22 host it finds 0 of the 20 cluster jobs and fails. It passes at llvm23 (ran); the worker caught this → pin `ROCKE_LLVM_FLAVOR=llvm23` in the test, or skip below llvm23.

The worker's other items are dropped. test_off_keeps_the_historical_hash already fails on develop, not because of this PR (ran). Its two_stage item repeats bartekxk's.

Utility: 🟢 Positive — deterministic split-K for small-M gfx1250 forward convolutions, in two-stage and single-launch cluster forms, mirrored in both engines and served from the AOT cache. split_k = 1 emits the same kernels as before.

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

First-round review: a local qwen3.6 model made the independent first pass. The orchestrator (Claude Opus 5.5) verified it. The checks were the golden conflict, regenerated on a scratch merge and compared case by case with develop, and on that merge the golden gate, native ctests and byte identity at three flavors. They also covered the PR's tests against develop at two flavors, the AOT grid per arch and flavor, an independent compile of every gfx1250 split-K job with ROCm clang 24, and an audit of the two_stage consumers. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

…ke/gfx1250-fwd-splitk Conflicts in generated goldens only, re-blessed on the merged tree: - rocke_representative_ir_sha256.json: develop's cases plus this PR's split-K cases; every hash matches one side. - differential llvm_gfx_all.json: develop's per-flavor layout, with conv_abi, conv_fwd_split_k_reduce and conv_implicit_gemm re-recorded (--only); all other families are develop's. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…luster launch - KernelIdentity.needs_stage2 (two_stage and not split_k_cluster) is the only predicate for "launch the Stage 2 reduce"; two_stage keeps meaning "takes the split-K workspace ABI", which the cluster forward also needs. The cached benchmark's fwd and wgrad Stage 2 launches use it, and kernel_sweep reports cluster jobs as "cluster". Hashes are unchanged. - Runtime._launch_cluster is the single hipDrvLaunchKernelEx call site, used by _enqueue and by prepare_launch's callable; the prepared config is still built once. - test_grid_has_one_binary_per_degree pins ROCKE_LLVM_FLAVOR=llvm23: the cluster jobs exist only there, so it failed on llvm22 hosts. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-10-27T11:12:38Z

@bartekxk, thanks. Both findings are addressed in fe16a31.

two_stage on the cluster identity.
- KernelIdentity.needs_stage2 (two_stage and not split_k_cluster) is now the one predicate for "launch Stage 2", and it is documented as such.
- Both Stage 2 launch sites in the cached benchmark (fwd and wgrad) use it.
- kernel_sweep now reports cluster jobs as cluster instead of two-stage.
- two_stage keeps its meaning of "takes the split-K workspace ABI", which both forward forms need, so identity hashes and cache names are unchanged.
- Test: test_only_the_two_stage_form_needs_stage2.
prepare_launch duplication.
- The hipDrvLaunchKernelEx call now lives in one helper, Runtime._launch_cluster, used by both _enqueue and the prepared callable.
- The config is still built once, at prepare time.
- Test: test_launch_and_prepared_launch_share_the_cluster_launch.

@cgmillette, thanks.

N1: fixed in fe16a31. test_grid_has_one_binary_per_degree now pins ROCKE_LLVM_FLAVOR=llvm23 for the enumeration. It failed on llvm22 before the fix and now passes at llvm22 and llvm23.
Golden conflict: resolved by merging develop (a185387), with both goldens re-blessed on the merge.
- Representative IR: develop's cases plus this PR's split-K cases. Every hash matches one side.
- Differential: develop's new per-flavor layout, with only conv_abi, conv_fwd_split_k_reduce and conv_implicit_gemm re-recorded. The diff against develop is exactly this PR's entries.

…ke/gfx1250-fwd-splitk Conflict only in platform/tests/golden/rocke_representative_ir_sha256.json; re-blessed with rocke_ir_parity_harness.py --write on the merge result. The golden is develop's cases plus this PR's split-K cases; every hash matches one side. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Review: bartekxk asked to investigate auto-deduction for the forward split-K. The forward split now resolves split_k=-1 the same way wgrad and dgrad do, through the shared conv split-K resolver. - rocke.helpers.split_k.select_split_k_fwd: the degree policy. - It splits only when the unsplit grid is under half of a target of 8 workgroups per CU, and the reduction is at least TARGET_SLICE_K deep. - The degree is the power of two toward that target, at most 32, with at least 2 K tiles per slice. - The constants are fitted on gfx1250 split-K forward degree curves. - The partials meet in fp32 slabs, so there is no 16-bit accuracy cap. - kernels.common.conv_split_k: resolve_conv_split_k("fwd", ...) now runs this policy on FWD_SPLIT_K_ARCHS (gfx1250). - The result is held to the gridDim.z budget and the i32 workspace. - It resolves to 1 elsewhere, and on pointwise and 3-D problems. - The override and explicit paths match dgrad's. - allow_unsplit=False raises instead of returning degree 1. - The AOT cache takes its forward split-K arch list from the same tuple. - The builder resolves split_k=-1 before validation, as wgrad and dgrad do (resolve_fwd_split_k_spec). lower_conv_implicit_gemm also resolves it before choosing an engine, because the C++ port, like its wgrad port, takes explicit degrees only. Emitted IR is unchanged. - benchmark --run-from-cache gives forward the same --split-k meaning as wgrad: - -1 launches a cached split-K binary at the resolved degree, and skips it where the policy does not split; - 0 sweeps --split-k-degrees plus that degree; - n pins n. The fixed forward ladder is gone. - CPU tests cover the policy, the resolver, the builder and backend resolution, and the benchmark modes. A GPU test launches the binary that an auto spec builds at the resolved degree. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette · 2026-10-29T08:41:10Z

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

PR #14393 — split-K forward implicit-GEMM convolution on gfx1250 (jakpiase, AICK-2778) — 🟢 Follow-up: all prior items addressed — the auto-deduced degree bartekxk asked for (`77f43d41`) checks out on a clean merge with develop; one 🟡 remains (auto with a fused epilogue raises), and a code-owner approval is still missing

Prior item	Status
Copilot (10-26) ×9: split-K with bias / ReLU / clamp epilogues; a channel-quad overrun; jobs beyond gfx1250 / llvm23; `--split-k-degrees` ignored; `split_k <= 0`; a NULL `dtype_d`; a builder freed before init; defaulted spec fields; cluster setup inside the timed loop	🟢 FIXED `24177571` / `906cbfc5` / `57ce3193`, still fixed at head. On the merge, `test_conv_fwd_split_k.py` passes with this tree's C++ engine binding. That includes the four engine-agreement tests CI skips: fused epilogue rejected by both, NULL `dtype_d`, overlong builder name. `--split-k 0` still sweeps `--split-k-degrees` after `77f43d41`, and both engines reject the must-reject parity configs alike (ran). The prepared cluster launch is set up once: `test_cluster_launch.py` passes 21/21 on all four CI shards (read). All 9 threads resolved
bartekxk (10-26) ×2: the cluster identity's `two_stage=True` is a latent "launch Stage 2" trap; `prepare_launch` duplicates `_enqueue`'s cluster launch	🟢 FIXED `fe16a31a`. `KernelIdentity.needs_stage2` (`kernel_identity.py:219`) now gates both Stage 2 launches (`benchmark_implicit_gemm_conv.py:5735`, `:5767`), and `kernel_sweep` labels cluster jobs `cluster`. The other `two_stage` readers are ABI, workspace and grid sites that both forms need (read). `test_only_the_two_stage_form_needs_stage2` passes (ran). `Runtime._launch_cluster` holds the only `hipDrvLaunchKernelEx` call (`hip_module.py:707`), and `test_launch_and_prepared_launch_share_the_cluster_launch` passes on all four CI shards (read)
bartekxk (10-26): is forward split-K auto-deduced, and what does it buy? → (10-29) "Please investigate auto deduce for this feature"	🟢 Done `77f43d41`. The 10-27 reply covers the benefit (small-output, deep-reduction layers). `select_split_k_fwd` (`helpers/split_k.py:409`) aims at 8 workgroups per CU. It splits only a grid under half of that, over a reduction of at least 512, and picks a power of two capped at 32 and at 2 K tiles per slice. `resolve_conv_split_k("fwd")` applies it on gfx1250 only and holds it to gridDim.z and the i32 workspace (`conv_split_k.py:337`, `:350`). The builder and `lower_conv_implicit_gemm` resolve `-1` before validation and engine choice, and the AOT cache shares the arch tuple. The new policy, resolver, auto-build, backend and benchmark-mode tests pass on the merge (ran) and on all four CI shards (read). `77f43d41` touches no golden and both goldens still pass, so the IR is unchanged. Threads resolved. See N1
Ours (10-27) N1 🟡: `test_grid_has_one_binary_per_degree` failed on llvm22 hosts; no full CI at `57ce3193` (golden conflict)	🟢 FIXED `fe16a31a`. The test pins `ROCKE_LLVM_FLAVOR=llvm23` (`test_conv_fwd_split_k.py:284`) and passes at llvm22 here (ran) and on all four CI shards (read). 🟢 ADDRESSED: five develop merges resolved the golden conflicts; the head merges cleanly and has a full CI run
CongMa13 (10-28): no new issues, endorsed the existing items; APPROVED at `dd07f5a3`	🟡 Stale, still counts. It predates `77f43d41` (the auto-deduced degree: a new policy, the forward resolver branch, builder and backend resolution, new `--split-k` semantics) and two develop merges. Develop does not dismiss stale approvals. The review decision is still REVIEW_REQUIRED with no reviewer requested, so a required code-owner approval is missing (`platform/`: rocke-platform, `library/`: rocke-lib) (ran: gh)
Bots / CI	🟢 At `77f43d41`: 44 succeeded, 51 skipped, 1 cancelled (a superseded `therock-pr-bot` run; the later one passed). Pre-commit, clang-tidy, ASAN, Component CI and the hipDNN superbuild pass. Multi-Arch CI tested gfx94X (Linux) and gfx110X (Windows), and all four `Test hipkernelprovider` shards pass. In `rocke_library_conv_pytest`, `test_conv_fwd_split_k.py` gives 42 passed and 7 skipped on every shard: the three GPU numerics classes (torch is not importable) and four C++-engine tests (CI builds no `rocke_engine`). The gfx125X build stage passed, but the gfx125X and gfx950 test jobs were skipped, so no gfx1250 kernel ran in CI. The only hardware evidence is the author's report that the numerics, `TestSplitKAutoNumerics` included, pass on gfx1250 (read). MERGEABLE, blocked on the review. Auto-merge (squash) is armed with a commit message saved before `77f43d41`, which does not mention the auto-deduced degree (ran: gh)
Our checks	🟢 • Merge: clean with develop `f22bed43` (7 behind, 17 ahead) (ran). • Interdiff since `57ce3193`: two PR commits, `fe16a31a` (above) and `77f43d41` (10 files). The five develop merges renumber the PR's parity configs 36–49 to 40–53, otherwise unchanged, and re-bless both goldens (ran: git). • Build: native build OK; 33 ctests pass (1 skipped) (ran). • Byte identity: GREEN at llvm20 / 22 / 23 for `conv_abi` (19 configs), `conv_fwd_split_k_reduce` (11) and `conv_implicit_gemm` (54), 0 bad. The cluster configs 44–48 need llvm23, so only that run compares them. Every family is GREEN at llvm22 (ran). • Goldens: the representative gate passes at all three flavors. Its only delta vs develop is the PR's split-K cases (4 at llvm20 / 22, where the cluster case is an expected failure; 5 at llvm23) plus counters. The differential golden (ll at llvm20–24, and canonical) has no unblessed change and nothing missing, and every IR config the PR emits is blessed (ran). • Tests: `test_conv_fwd_split_k.py` on the merge, with its C++ engine binding: 46 passed, 3 skipped (the torch-gated GPU numerics) (ran). The checks' seven-file run stopped at its first failure, `test_conv_dgrad_patch.py::TestPatchIdentity::test_off_keeps_the_historical_hash`, which fails on develop `f22bed43` too (ran). Its hand-kept list of hash-neutral identity fields lacks develop's `uniform_k_decode`, `tap_outer_k` and `dy_halo`, and no CI lane runs the file. This PR's `split_k_cluster` is one more such field: the PR added it to the lists in `test_conv_wgrad_prefetch_addr.py:326` and `test_conv_dgrad_spec_policy.py:354`, but not here. • Compile* (ROCm clang 24, -O3, on the merge): the 9 gfx1250 split-K parity configs of `conv_implicit_gemm` (40–48) compile, the 5 cluster ones at llvm23, and so do the 9 configs of `conv_fwd_split_k_reduce`. 0 failures, 0 newly rejected, no spills. The rejections are the PR's must-reject configs (49–53, and the reduce family's 9–10) and 5 that develop rejects too. `conv_abi`'s "0 compiled" is expected: it emits ABI lists, not IR (ran). Our first round's compile of all 1,860 AOT jobs was not repeated; the split-K grid is unchanged (read). • With #13936 (split-K wgrad, same base): the two PRs conflict only in one docstring bullet of `conv_split_k.py` and in the representative golden, and the code merges cleanly and agrees. #13936's bf16 cap of 12, from bf16's unit roundoff 2⁻⁸, bounds wgrad's 16-bit atomics only; this PR's forward partials meet in fp32 slabs and take no cap. #13936 changes only `_resolve_wgrad`, so `resolve_conv_split_k` composes (ran: git merge-tree). This PR does not touch `check_byte_identity.py`, so #14387 / #14392's rename problem (`LLVM23_ONLY_FAMILIES` → `TARGET_OP_PROFILE_FAMILIES`) does not apply (ran: git). • Rules 1–3 clean; the two new commits add no links or tracker IDs (ran)

Net-new:

#	Sev	file:line	Finding
N1	🟡	`kernels/common/conv_implicit_gemm.py:705`, `conv_split_k.py:321`	Auto (`split_k=-1`) checks only the arch, pointwise and 3-D before it picks a degree. It does not check the split form's own gates (`fwd_split_k_reason`, `conv_implicit_gemm.py:653`). On gfx1250, an auto spec with a fused ReLU, bias or clamp `acc_epilogue` resolves to a split degree: 4 on the PR's test problem, 16 on a batch-1 7×7 512→512 layer. The build then raises "split-K forward supports only the identity acc_epilogue", although the plain kernel builds (ran: probe on the merge). No shipped caller passes `-1` for forward yet, but the dispatcher wiring the author plans would hit it on fused-epilogue layers → resolve to 1 where `fwd_split_k_reason` rejects the split form, as the wgrad resolver does through `epilogue`, and add a test

Utility: 🟢 Positive — deterministic split-K for small-output gfx1250 forward convolutions, now with an auto-deduced degree.

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

Follow-up review: a local qwen3.6 model made the independent first pass. The orchestrator (Claude Opus 5.5) verified it. The checks were an interdiff of the PR's own changes since the head we reviewed (57ce3193). On a clean merge with develop f22bed43 they covered the native build and ctests, byte identity at llvm20 / 22 / 23, the representative and differential goldens, a ROCm clang 24 compile of the touched families' parity configs, the PR's split-K suite with that tree's C++ engine binding, and a probe of auto with fused epilogues. They also covered the failing patch test on develop, a simulated merge with #13936, and the threads, approvals and CI logs at head. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

Conflicts: - library/kernels/common/conv_split_k.py: module docstring. Kept develop's WMMA note on the wgrad reductions (#13936) and this branch's forward split-K bullet after the gridDim.z one. - platform/tests/golden/rocke_representative_ir_sha256.json: re-blessed with rocke_ir_parity_harness.py --write. Every case hash equals develop's or this branch's; the aliases block keeps develop's format. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…nnot take resolve_fwd_split_k_spec checked only the arch and the problem (pointwise, 3-D) before taking the policy's degree. On gfx1250 an auto spec with a fused accumulator epilogue (bias, scale, ReLU, clamp) resolved to a split degree, and the build then raised "split-K forward supports only the identity acc_epilogue" although the plain kernel builds. A resolved degree > 1 is now checked against fwd_split_k_reason (with the cluster settings left out, so their own errors still come from validation); a spec the split form rejects resolves to 1. The identity epilogue resolves as before, and no emitted IR changes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-10-29T10:23:37Z

@cgmillette, thanks. Replies to the 10-29 follow-up:

N1 (auto with a fused epilogue raises): fixed in 159563d.
- resolve_fwd_split_k_spec now checks a resolved degree > 1 against fwd_split_k_reason. A spec the split form rejects, such as one with a fused bias, scale, ReLU or clamp acc_epilogue, resolves to 1 and builds the plain kernel with its epilogue.
- The cluster settings are left out of that check, so their own errors still come from validation.
- New test: TestSplitKAutoBuild.test_auto_with_a_spec_the_split_cannot_take_builds_the_plain_kernel covers the four epilogues and checks that the identity epilogue still resolves to 4. No emitted IR changes.
test_conv_dgrad_patch.py::test_off_keeps_the_historical_hash: left alone. It fails the same way on develop (checked again on the new merge base): its hand-kept list lacks develop's own hash-neutral fields (uniform_k_decode, tap_outer_k, dy_halo*). Adding only split_k_cluster would not make it pass, so it belongs in a develop fix.

jakpiase · 2026-10-29T10:23:38Z

Merged develop (9d729d3280) in 741c104.

Conflicts:
- The conv_split_k.py module docstring: kept develop's WMMA wgrad note (feat(rocke): split-K wgrad on gfx1250 WMMA #13936) and this PR's forward bullet.
- The representative IR golden: re-blessed with rocke_ir_parity_harness.py --write. Every case hash equals develop's or this branch's, and no case is missing (309 cases at llvm20 / 22, 318 at llvm23).
Checks (CPU only, on 159563d):
- check_byte_identity.py --only conv,gfx1250: GREEN at llvm20 / 22 / 23 (21 / 21 / 19 families, 0 bad).
- Representative golden --check: OK at llvm20 / 22 / 23 and the llvm24 alias.
- run_diff.py --python-only --check-golden (ll at llvm20 / 22 / 23, and ir_canonical): GOLDEN OK, nothing missing.
- test_conv_fwd_split_k.py with this tree's C++ engine binding: 47 passed, 3 skipped (the gfx1250 GPU numerics).
- test_conv_split_k.py 206 passed, test_conv_dgrad_spec_policy.py 24 passed, test_conv_wgrad_prefetch_addr.py 14 passed, test_cluster_launch.py + test_llvm24_profile.py 29 passed.
- test_conv_dgrad_patch.py: 1 failure, which also fails on develop.
GPU tests were not run this round.

Brings in #14392 (wave32 WMMA direct and small-channel forward conv on gfx1250). Conflicts: - platform/CMakeLists.txt: keep both installed tests (test_direct_fwd_wmma.py and test_conv_fwd_split_k.py). - rocke_engine.cpp: keep both instance headers (instance_conv_fwd_small_wmma.h, instance_conv_fwd_split_k_reduce.h). - Representative IR golden: re-blessed with rocke_ir_parity_harness.py --write; every case hash equals develop's or this branch's. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-10-29T12:51:21Z

Merged develop (050acb2a69, #14392) in 52044d6.

Conflicts:
- platform/CMakeLists.txt: kept both installed tests (test_direct_fwd_wmma.py, test_conv_fwd_split_k.py).
- rocke_engine.cpp: kept both instance headers (instance_conv_fwd_small_wmma.h, instance_conv_fwd_split_k_reduce.h).
- Representative IR golden: re-blessed with rocke_ir_parity_harness.py --write. Every case hash equals develop's or this branch's, and none is missing (311 ok + 19 expected failures at llvm20 / 22, 330 ok at llvm23).
- check_byte_identity.py, SUPPORT_MATRIX.md and the parity harness merged cleanly. This PR registers no llvm23-only family (its cluster configs are in conv_implicit_gemm), so TARGET_OP_PROFILE_FAMILIES is unchanged. No parity config renumbering was needed.
Checks (CPU only, on 52044d6):
- check_byte_identity.py --only conv,gfx1250: GREEN at llvm20 / 22 / 23 (23 / 23 / 21 families, 0 bad).
- Representative golden --check: OK at llvm20 / 22 / 23 and the llvm24 alias.
- run_diff.py --python-only --check-golden (ll at llvm20 / 22 / 23, and ir_canonical): GOLDEN OK, nothing missing or changed.
- Native build OK; ctest 33 / 33 passed.
- test_conv_fwd_split_k.py with this tree's C++ engine binding: 47 passed, 3 skipped (gfx1250 GPU numerics). test_cluster_launch.py + test_llvm24_profile.py: 29 passed.
- Library CPU tests (split-K fwd, test_conv_split_k, dgrad spec policy, wgrad prefetch, builder signature contract, kernel label, feat(rocke): wave32 WMMA direct and small-channel forward convolution for gfx1250 #14392's test_fwd_small_wmma / test_direct_fwd_wmma, dgrad patch): 355 passed, 14 skipped, 1 failed. The failure is test_conv_dgrad_patch.py::test_off_keeps_the_historical_hash, which fails on develop 050acb2a69 too.
- black 25.12: clean.
GPU tests were not run this round.

Brings in #14385 (strided, dilated and packed-group WMMA direct dgrad). The representative IR golden was re-blessed with rocke_ir_parity_harness.py --write; every case hash equals develop's or this branch's, and cases and expected failures are the union of both. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-10-29T15:12:22Z

Merged develop (a0ff54234e, #14385) in ab72803.

Conflicts: only the representative IR golden (summary counters). Re-blessed with rocke_ir_parity_harness.py --write: cases and expected failures are the union of both sides, and every case hash equals develop's or this branch's (311 ok / 25 expected failures at llvm20 and llvm22, 336 ok at llvm23).
platform/CMakeLists.txt, rocke_engine.cpp, SUPPORT_MATRIX.md and the parity harness merged cleanly with both sides kept.
Checks (CPU only):
- check_byte_identity.py --only conv,gfx1250: GREEN at llvm20, llvm22 and llvm23.
- rocke_ir_parity_harness.py --check: OK at llvm20, llvm22, llvm23 (and the llvm24 alias).
- run_diff.py --python-only --check-golden (ll at all three flavors, and ir_canonical): GOLDEN OK, 0 missing.
- pytest (split-K, dgrad patch / spec policy, wgrad prefetch, builder signature, direct dgrad WMMA, cluster launch): 163 passed, 11 skipped, 1 failed. The failure, test_conv_dgrad_patch.py::TestPatchIdentity::test_off_keeps_the_historical_hash, also fails on develop.
- platform/tests/test_rocke.py: 330 passed, 110 skipped.
- black 25.12: clean.
No GPU runs on this merge.

Conflict: platform/tests/golden/rocke_representative_ir_sha256.json (summary counters only). Re-blessed with rocke_ir_parity_harness.py --write: every case hash equals develop's or this branch's; ok_cases 313/313/338 at llvm20/22/23. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-10-29T19:25:29Z

Merged develop (7e3d78a3d1) in e3352f1.

Conflicts: only the representative IR golden (summary counters). Re-blessed with rocke_ir_parity_harness.py --write: every case hash equals develop's or this branch's, nothing dropped. ok_cases is now 313/313/338 at llvm20/22/23 (develop's two new gfx942 f64 GEMM cases on top of this branch's).
Checks (CPU only, no GPU):
- black 25.12: clean.
- Library pytest (split-K fwd, builder signature contract, dgrad patch/spec policy, wgrad prefetch): 110 passed, 7 skipped, 1 failed. The failure, test_conv_dgrad_patch.py::TestPatchIdentity::test_off_keeps_the_historical_hash, also fails on pristine develop, so it is pre-existing.
- Platform test_rocke.py and runtime/test_cluster_launch.py: 357 passed, 110 skipped.
- Byte identity (--only conv,gfx1250) at llvm20/22/23: GREEN.
- Differential golden check: OK.
- Representative IR --check: OK.

cgmillette · 2026-10-29T22:23:52Z

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

PR #14393 — split-K forward implicit-GEMM convolution on gfx1250 (jakpiase, AICK-2778) — 🟡 Follow-up: no closer look needed — our N1 is fixed in `159563d9`, the four develop merges change only their conflict hunks, and bartekxk approved at head; the Windows tests need a re-run (a Chocolatey rate limit in setup), and the auto-merge message needs a refresh

Prior item	Status
Fixed at our 10-29 review: Copilot (10-26) ×9 (fused epilogues, channel-quad overrun, job scope, `--split-k-degrees`, `split_k <= 0`, NULL `dtype_d`, builder init, defaulted spec fields, cluster setup in the timed loop); bartekxk (10-26) ×2 (`two_stage` trap, duplicated cluster launch) and (10-29) auto-deduce; ours (10-27) N1 (llvm22 test failure)	🟢 Still FIXED. Since `77f43d41` the PR adds only `159563d9` (next row) and four develop merges. Each merge differs from git's own merge only in its conflict hunks: the representative golden, a `conv_split_k.py` docstring that keeps both sides' bullets, and both-sides keeps in `platform/CMakeLists.txt` and `rocke_engine.cpp` (ran: git). On the merge, `test_conv_fwd_split_k.py` with this tree's C++ engine binding gives 47 passed and 3 skipped (the torch-gated GPU numerics), the engine-agreement tests included (ran)
Ours (10-29) N1 🟡: auto (`split_k=-1`) with a fused epilogue raised instead of building the plain kernel	🟢 FIXED `159563d9`. `resolve_fwd_split_k_spec` (`conv_implicit_gemm.py:741`) checks a degree > 1 against `fwd_split_k_reason`, without the cluster settings, and falls back to 1. Probe on the merge at llvm22 and llvm23 (ran): on gfx1250, ReLU / bias / clamp / scale resolve to 1 and build the plain kernel under the explicit build's name. The identity epilogue still splits: 4 on the PR's test problem, 16 on a batch-1 7×7 512→512 layer, 8 and 4 with K = 10 and kpg = 6. On gfx1151, gfx1201, gfx950 and gfx942 every case resolves to 1 without raising. Through the backend (fp16), auto lowers to the explicit degree's IR in both engines, and C++ equals Python. The new test passes (ran). With `split_k_cluster` set, auto still reports the cluster's own errors (channels not a multiple of 4; llvm22, in both engines), as the reply says
Ours (10-29) note: `test_conv_dgrad_patch.py::test_off_keeps_the_historical_hash` fails on develop, and its list of hash-neutral fields also lacks this PR's `split_k_cluster`	🟡 DEFERRED: "belongs in a develop fix". Acceptable: it fails on develop too (now #14642), and no CI lane runs the file (ran). Whichever of this PR and the #14642 fix lands second must add `split_k_cluster` to that list
CongMa13 (10-28): APPROVED at `dd07f5a3`	🟡 Stale, still counts: predates `77f43d41`, `159563d9` and six develop merges
bartekxk (10-29): APPROVED at head `e3352f11`	🟢 At head. The review decision is now APPROVED, so the code-owner approval we found missing is in; no reviewer is still requested (ran: gh)
Ours (10-29) note: auto-merge is armed with a commit message saved before `77f43d41`	🟡 OPEN. The squash message saved on 10-28 still lacks the "Auto-deduced degree" section and the new `--split-k` semantics. With the approval in, the PR would land with that message once CI passes → re-arm auto-merge with the current PR body (ran: gh)
Bots / CI	🟡 At head `e3352f11` (22:11Z): pre-commit, clang-tidy, Component CI, the ASAN summary and both hipDNN superbuilds pass, and TheRock built math-libs for gfx110X, gfx125X and gfx950; gfx94X is still building. The one failure, Windows gfx110X `Test sanity`, failed in setup: `choco install ninja` got HTTP 429 (Too Many Requests) from the Chocolatey feed, the same as #14387's sanity job today. No test ran, and the Windows component tests, where `Test hipkernelprovider` runs, were skipped. So no rocKE test lane has run at head yet, and the PR is BLOCKED → re-run the failed job. The cancelled check is a superseded PR-bot run (ran: gh)
Our checks	🟢 • Merge: clean with develop `b0f07f86` (8 behind, 22 ahead), and still conflict-free with the newer develop `e8b95982` (ran). • Golden re-bless in the four merges (`741c1042`, `52044d6b`, `ab728037`, `e3352f11`): every case and expected failure equals the branch's or develop's entry, none is dropped, and the counters match a recount (313 / 313 / 338 ok cases). In `e3352f11`, where the sides differ (seven gfx950 attention cases develop re-blessed), it takes develop's. Against develop, the merged golden adds only the PR's split-K cases: 4 at llvm20 / 22, plus the cluster case as an expected failure, and 5 at llvm23 (ran). • Build: native build OK; 33 ctests pass (1 skipped) (ran). • Byte identity: GREEN at llvm20 / 22 / 23 for the 33 touched families, `conv_implicit_gemm` (54 configs), `conv_fwd_split_k_reduce` (11) and `conv_abi` (19) included, and for every family at llvm22. The library C++ parity tests pass at all three (ran). • Goldens: the representative gate passes at all three flavors. The differential golden (ll at llvm20–24, and canonical) has no unblessed change and nothing missing. Its 176 NEW entries (111 at llvm20 / 22) are develop's own: every config the PR's three families emit is blessed, and the count is 4 below develop's because the PR also blesses develop's `conv_implicit_gemm` 36–39 (ran). • Tests: the PR's seven test files give 494 passed, 19 skipped and 1 failed, `test_off_keeps_the_historical_hash` (above) (ran). • Compile (ROCm clang 24, -O3) of the touched conv families: 0 failures, 0 newly rejected, no spills (ran). • `159563d9` adds no links or tracker IDs, and the rule scan is clean (ran)

Net-new: none.

Utility: 🟢 Positive — deterministic split-K for small-output gfx1250 forward convs; auto degree falls back cleanly.

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

Follow-up review: a local qwen3.6 model made the independent first pass. The orchestrator (Claude Opus 5.5) verified it. The checks covered the PR's one new commit since the head we reviewed (77f43d41) and its four develop merges, each compared with git's own merge, with a case-by-case audit of the golden re-bless. On a clean merge with develop b0f07f86 they covered the native build and ctests, byte identity at llvm20 / 22 / 23, the representative and differential goldens, a ROCm clang 24 compile of the touched conv families, the PR's tests with that tree's C++ engine binding, and a probe of auto split-K across epilogues, channel counts, targets, flavors and both engines. They also covered the threads, approvals, auto-merge state and CI logs at head. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

cgmillette

lgtm

…250-fwd-splitk Conflicts: only the representative IR golden (expected failures and summary counters). Re-blessed with rocke_ir_parity_harness.py --write; every case hash equals develop's or this branch's and nothing is dropped. ok_cases is now 313/313/342 at llvm20/22/23. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Brings in #14393 (gfx1250 split-K forward conv) and #14379 (conv sweep benchmarks). Only the representative IR golden conflicted, on its summary counters: develop's split-K cases and this branch's five depthwise dgrad cases per flavor do not overlap. Merged case by case against the merge base, with the counters recomputed from the merged cases. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

## Motivation On wave32 targets, the row-streaming depthwise dgrad (`DirectDepthwiseDgradStreamSpec`, AOT variant `direct_depthwise_dgrad`) runs one channel per lane. It uses 16-bit loads and stores and scalar FMAs. Depthwise dgrad also has no direct kernel for filters past 11x11 or for non-square decomposed large kernels (51x5 / 5x51), so those layers fall back to implicit GEMM. This PR does three things: - it adds channel packing to the row-streaming depthwise dgrad; - it adds a rotated-filter dgrad mode to the output-stationary depthwise kernel, with per-axis padding for non-square filters; - it mirrors both in the C++ engine with byte-identical IR. The C++ engine had no port of the row-streaming depthwise dgrad, so this PR also adds its scalar form. Like `conv_direct_dgrad_wmma`, the new binaries are served by the direct-conv AOT cache. There is no dispatcher wiring. JIRA ID: AICK-2337 ## Technical Details **Channel-packed row-streaming dgrad (`DirectDepthwiseDgradStreamSpec.cpl`, `library/kernels/common/conv_direct_grouped.py`)** - With `cpl` 2, 4 or 8, every lane owns `cpl` adjacent channels: - dY rows load and dX rows store as one `cpl`-wide vector per lane; - every tap runs as `cpl / 2` packed `<2 x f32>` FMAs over channel pairs. - The schedule does not change: dY rows stream into `stride * KH` circular accumulator slots with a few rows prefetched, and a tail loop flushes the rows past the last dY row. - `cpl == 1` emits the unchanged kernel. - Validator rules: - `cpl` in {1, 2, 4, 8} and `groups % cpl == 0`; - `cpl * KH * KW <= 200` for the packed weights; - `cpl * block_w * KH <= 192` for the accumulator slots. - `block_ch` grows by `cpl`, and the kernel name gains `cpl{N}`. **Rotated-filter dgrad (`DirectDepthwiseTiledSpec.dgrad`)** - A stride-1 depthwise dgrad is the forward correlation of dY with the filter rotated by 180 degrees. The padding is `KH-1-PAD` rows and `KW-1-PAD_W` columns. - The output-stationary forward kernel computes it through the dgrad kernarg block: - the extents and the A-side strides swap roles; - weight row `r` is read from row `KH-1-r`, reversed; - the epilogue addresses dX through its kernarg strides; - the grid tiles `H x W`. - Only one filter row is register-resident, so this mode takes filters up to 64 on a side, including non-square ones. The forward mode keeps its 32 limit. - The kernel name gains the `dg` flag. `dgrad=False` emits the unchanged kernel. - The mode packs one channel per lane: both engines refuse `dgrad` with `cpl > 1` (the forward mode's channel-packed kernel, from #14392, is forward only). **Per-axis padding (`DirectConvProblem.pad_w`)** - `pad_w = -1` means `PAD`. `PAD_W` is the effective column padding, and `Wo` uses it. - Only the rotated-filter dgrad takes a problem with `PAD_W != PAD`: - the forward tiled mode refuses one; - the WMMA convolutions (`direct_fwd_wmma`, `fwd_smallc_wmma`, `fwd_smallk_wmma` through `wmma_fwd_gate`, `direct_dgrad_wmma`, and #14387's stem dgrad kernels `dgrad_patch_wmma` and `dgrad_smallc_wmma`) refuse one in their validators, in both engines. They offset columns by `PAD`. The check is `one_padding_reason` / `rocke_direct_conv_one_padding_reason`, shared with the row-streaming dgrad; - `direct_spec_for_problem` refuses one for every other variant; - `direct_job_for_identity` builds non-square / per-axis identities only for the new variant. - `kernel_identity` matches `pad_w` against `PAD_W`, and `conv_args` passes `pW = PAD_W`. - The direct problem's capability key (`DIRECT_PROBLEM_CAPABILITY_ATTRS`, from #14379) includes `pad_w`, so problems that differ only in their column padding never share a cached support verdict. - The benchmark accepts `-p != -q` only for a stride-1 depthwise dgrad. **AOT (`library/kernels/conv/direct_variants.py`, `direct_aot_jobs.py`)** - `direct_depthwise_dgrad` gains the `DW_DGRAD_CPL_GRID` knobs. The `cpl == 1` knobs keep their identities. - New variant `direct_depthwise_tiled_dgrad`, in `DIRECT_WAVE_PARAMETRIC`: - filters 13x13 to 31x31, plus 51x5 and 5x51, all at "same" padding; - tiles `DW_TILED_DGRAD_BLOCK_W` x `DW_TILED_DGRAD_BLOCK_H`. - `DirectCaps` gains `kw` / `pad_w` for those rows. **C++ engine** - New `rocke/instance_conv_direct_grouped_dw_dgrad.h` and `conv_direct_grouped_build_depthwise_dgrad_stream.cpp` port `DirectDepthwiseDgradStreamSpec`. This covers the defaults, `block_ch`, the kernel name, `validate` / `is_valid` (same order and reason text as Python), and the scalar and `cpl > 1` builders. They also add `_new` and `lower_to_llvm` entries. - `rocke_direct_depthwise_tiled_spec_t.dgrad` covers the builder mode, the `dg` name flag, the dgrad filter limit and the padding checks of `validate()`. - `rocke_direct_conv_problem_t.pad_w` (default -1), `rocke_direct_conv_problem_pad_w` and `rocke_direct_conv_one_padding_reason`. Every C++ mirror of `DirectConvProblem.Wo` in the direct-conv family uses it, as the Python property does. - The depthwise kinds have no `rocke_engine` binding, here or on `develop`. The parity pair is the cross-engine gate. **Docs:** `dsl_docs/instances/convolution.md` (the dgrad mode in the tiled section, and a new row-streaming depthwise dgrad section), `dsl_docs/instances/index.md`, `SUPPORT_MATRIX.md`. ## Test Plan - `check_byte_identity.py`, all families, at llvm20, llvm22 and llvm23. - `conv_direct_grouped_emit.{py,c}` parity configs 179-200 (after develop's wave32 forward configs 166-178): - row-streaming dgrad, scalar and `cpl` 2 / 4 / 8, in fp16 and bf16, at stride 1 and 2, on wave32 (gfx1250) and wave64 (gfx950); - rotated-filter dgrad on 13x13, 31x31, 51x5 (`pad_w`), 5x51 (`pad_w`, unrolled rows), a 7x7 with a padding that differs per axis and is not "same", and a 5x5 with `pad_w = 0` (the kernel name carries the per-axis padding); - nine specs both engines reject: - per-axis padding on the forward tiled mode; - the dgrad mode at stride 2; - a 65-row filter; - `groups % cpl != 0`; - `cpl = 3`; - the packed-accumulator bound; - `pad_w < -1` on the rotated-filter dgrad; - per-axis padding on the row-streaming kernel; - `pad_w < -1` on the row-streaming kernel. - Differential golden (`tests/instances/differential/golden/llvm_gfx_all.json`): the new `conv_direct_grouped` configs that both engines emit are blessed for `ll` (llvm20 / llvm22 / llvm23) and `ir_canonical`. - Golden IR: five new gfx1250 cases (row-streaming scalar, `cpl` 2, `cpl` 8 stride 2 bf16; rotated-filter 31x31 and 51x5) at every flavor. - `library/tests/test_direct_depthwise_dgrad_v2.py`, now in CTest: validators, AOT grid entries (including the packed-entry spill gate), the padding gate, and numerics against a numpy reference on any arch rocke models (each case built at the arch's wave size; cases the arch's validator refuses are skipped). - `rocke_conv_direct_dw_dgrad_host` ctest: the stream name-overflow cleanup, `pad_w < -1`, the per-axis padding name, the `dgrad` / `cpl > 1` refusal, and the per-axis padding refusal of the six WMMA kernels (including the two stem dgrad kernels). It is listed in the `quick` and `ffm-quick` categories of `ROCKE_ENGINE_test_categories_external.yaml`, so CI runs it. - Targeted CPU pytest, with and without the `rocke_engine` binding on the path. - black 25.12.0 / clang-format 18.1.4 on the changed files. ## Test Result - Byte identity GREEN at llvm20, llvm22 and llvm23. `conv_direct_grouped` has 201 configs: the 13 new compiled configs are IDENTICAL, and the 9 new reject configs are BOTH_REJECTED. - Golden `--check` passes at every flavor. Every existing hash is reproduced: `cpl == 1` and `dgrad=False` emit the old IR. The golden diff adds case keys and updates the summary counters. - Differential golden `run_diff.py --python-only --check-golden` is GOLDEN OK for `ll` at every flavor and for `ir_canonical`, with 0 unblessed changes and 0 missing. The bless adds 13 `conv_direct_grouped` configs (179-190, 197) to each section; no existing hash changes or drops. - `test_direct_depthwise_dgrad_v2.py` on gfx1250 (before the latest develop merge): 11 passed, 10 subtests (the GPU cases). - After merging develop at `311cb32338e` (#14387, wave32 WMMA stem dgrad), CPU only: - byte identity GREEN at llvm20, llvm22 and llvm23 for `--only conv,gfx1250`, including `conv_direct_grouped` (201 configs), `conv_dgrad_patch_wmma` (9) and `conv_dgrad_smallc_wmma` (9); - golden `--check` passes (llvm20 / 22 / 23 / 24); the differential golden is GOLDEN OK for `ll` at every flavor and for `ir_canonical`, with 0 unblessed changes and 0 missing; - native build and ctest: 34/34 passed, including `rocke_conv_direct_dw_dgrad_host`; - `test_direct_depthwise_dgrad_v2.py`, `test_direct_kernel_cache.py`, `test_direct_dgrad_wmma.py`, `test_direct_fwd_wmma.py`, `test_fwd_small_wmma.py`, `test_dgrad_stem_wmma.py`, `test_kernel_label.py`: 213 passed, 8 skipped with the `rocke_engine` binding; 203 passed, 18 skipped without it. The skips need a GPU or the binding; - `test_direct_depthwise_fwd_cpl.py`, `test_conv_direct_grouped_backend_parity.py`, `test_direct_depthwise_col_validate.py`, `test_conv_dgrad_depthwise_windowed.py`: 76 passed, 8 skipped with the binding; 70 passed, 14 skipped without it. - The merge's only conflict was the llvm23 summary counters of the representative golden: develop adds 4 gfx1250 stem-dgrad cases and this PR 5 gfx1250 depthwise-dgrad cases. Cases and expected failures are the union of both sides, and the counters are recomputed from them: `ok_cases` is 314 / 314 / 342 (llvm20 / 22 / 23). Every case hash equals develop's or this branch's. - Develop adds no `conv_direct_grouped` parity configs and changes no `conv_direct_grouped` entry of the differential golden (#14387 adds its own `conv_dgrad_patch_wmma` / `conv_dgrad_smallc_wmma` sections), so this PR's configs keep 179-200. - After that merge, #14387's two stem dgrad kernels also refuse a per-axis padding (they offset columns by `PAD`); their emitted IR is unchanged. - After merging develop at `32aed8707aa` (#14393 split-K forward conv, #14379 conv sweep benchmarks), CPU only: - the only conflict was again the representative golden's summary counters; merged case by case against the merge base, `ok_cases` is 318 / 318 / 347. The golden differs from develop only by this PR's 5 cases per flavor, and every hash equals develop's or this branch's; - #14379's `test_capability_keys.py` caught that its direct capability key left out `pad_w`; fixed in `23b898c8d2c` with a regression test; - byte identity GREEN at llvm20, llvm22 and llvm23 for `--only conv,gfx1250` (`conv_direct_grouped` 201 configs); golden `--check` passes (llvm20 / 22 / 23 / 24); the differential golden is GOLDEN OK for `ll` and `ir_canonical`; - native build and ctest: 34/34 passed; - the seven test files above: 213 passed, 8 skipped with the binding; the four above: 76 passed, 8 skipped with the binding; - develop's new sweep tests (`test_capability_keys`, `test_run_conv_shards`, `test_module_cache`, `test_aot_launch_zeroing`, `test_batched_timing`, `test_early_stop`, `test_streamk_sweep_launch`) plus the grouped-conv dispatch, split-K forward and this PR's tests: 445 passed, 10 skipped, 1 failed. The failure, `test_early_stop.py::test_staged_timing_reports_the_median_of_stalled_samples`, fails identically on `develop`. - After merging develop at `b33ff848a7b` (#14494 split the rocKE CMake test registration into `tests/cmake/*.cmake`), the `rocke_conv_direct_dw_dgrad_host` registration moves verbatim into `platform/tests/cmake/InstanceTests.cmake`; `tests/CMakeLists.txt` is develop's unchanged. CPU only: byte identity GREEN at llvm20 / 22 / 23 (`conv_direct_grouped` 201 configs), golden `--check` passes (llvm20 / 22 / 23 / 24), the differential golden is GOLDEN OK, native ctest 34/34 (`rocke_conv_direct_dw_dgrad_host` passes), `test_cmake_test_registration.py` 5 passed, and the test counts above are unchanged. - After merging develop at `64c2954a8e` (#14679 split the representative-IR golden into per-family shards), this PR's 5 cases per flavor move unchanged into the `conv_direct` and `conv_direct_depthwise_tiled` shards under `platform/tests/golden/rocke_representative_ir_sha256/`; the shards differ from develop's by those additions only. CPU only: golden `--check` against the shard directory passes (llvm20 / 22 / 23 / 24), byte identity GREEN at llvm20 / 22 / 23, the differential golden is GOLDEN OK, native ctest 34/34, and `test_ir_golden.py` / `test_rocke_ci_static.py` / `test_cmake_test_registration.py` pass; the test counts above are unchanged. - The full CPU pytest suite, run before the latest develop merges, showed no new failures against `develop`. The remaining failures are pre-existing and identical on `develop`: `test_conv_abi.py::test_cache_applies_runtime_launch_limits`, six cases in `test_conv_batch_and_miopen_parse.py`, `test_conv_dgrad_patch.py::TestPatchIdentity::test_off_keeps_the_historical_hash`, and `test_engine_linkage.py::test_installed_pytest_gets_distinct_artifact_environment_variables`. - No GPU run on gfx950 / gfx942 (no hardware here). Those targets are pinned by the parity and golden gates. - No performance results are included in this PR. ## Submission Checklist - [x] Both engines emit byte-identical IR for every new code path; the parity pair and golden cases are added. - [x] Defaults keep every existing configuration and its IR unchanged. - [x] New C++ code is in new files where the change allowed it. - [x] CTest entry, docs and support-matrix rows. - [ ] Dispatcher wiring: not part of this PR; the new binaries are served by the direct-conv AOT cache. --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

assistant-librarian · 2026-11-03T11:07:14Z

TheRock Submodule Bump Activity

Newest first

2026-11-03 — Included in TheRock via ROCm/TheRock#9698.

jakpiase and others added 7 commits October 26, 2026 14:08

test(rocke): pin the split-K forward test copies

9717251

The GPU numerics of test_conv_fwd_split_k copy through pinned host buffers: a pageable copy registers the host range with the driver, which must not overlap a kernel in flight. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase requested review from a team as code owners October 26, 2026 15:24

github-actions Bot added documentation project: hip-kernel-provider rocke: platform Touches rocKE platform labels Oct 26, 2026

jakpiase requested a balanced review from Copilot October 26, 2026 15:25

Copilot started reviewing on behalf of jakpiase October 26, 2026 15:26

View session

Copilot AI reviewed Oct 26, 2026

View reviewed changes

assistant-librarian Bot added the organization: ROCm label Oct 26, 2026

jakpiase and others added 3 commits October 26, 2026 16:08

bartekxk reviewed Oct 26, 2026

View reviewed changes

Comment thread ...roviders/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_implicit_gemm_conv.py Resolved

cgmillette mentioned this pull request Oct 27, 2026

fix(rocke): repair the differential golden and run its check in CI #13974

Merged

1 task

jakpiase and others added 2 commits October 27, 2026 10:51

jakpiase and others added 2 commits October 27, 2026 14:59

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-fwd-splitk

fcf7ad5

cgmillette mentioned this pull request Oct 29, 2026

[rocke] test_conv_dgrad_patch.py::TestPatchIdentity::test_off_keeps_the_historical_hash fails on develop since #13075 (the test's copy of the hash rule misses the new dgrad fields) #14642

Open

jakpiase and others added 2 commits October 29, 2026 10:11

bartekxk approved these changes Oct 29, 2026

View reviewed changes

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-fwd-splitk

26a31e3

cgmillette approved these changes Oct 30, 2026

View reviewed changes

cgmillette and others added 2 commits October 29, 2026 21:18

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-fwd-splitk

60281c9

jakpiase merged commit 4124e7b into develop Oct 30, 2026
97 of 98 checks passed

jakpiase deleted the users/jakpiase/rocke/gfx1250-fwd-splitk branch October 30, 2026 06:41

jakpiase mentioned this pull request Oct 30, 2026

feat(rocke): channel-packed and large-filter depthwise dgrad #14386

Merged

4 of 5 tasks

cgmillette mentioned this pull request Oct 30, 2026

test(rocke): qualify rocKE LLVM 24 lane #14692

Merged

8 tasks done

aledudek mentioned this pull request Nov 4, 2026

feat(rocke): TF32 compute for the conv kernels (native XF32 on gfx942, BF16x3 on gfx950) #14114

Open

1 task

Conversation

jakpiase commented Oct 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Changes

Testing

Uh oh!

therock-pr-bot Bot commented Oct 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

therock-pr-bot Bot commented Oct 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

bartekxk commented Oct 26, 2026

Review: PR #14393 — split-K forward implicit-GEMM convolution on gfx1250

Summary

Findings

1. Cluster split-K forward's KernelIdentity.two_stage=True is a landmine gated only by an incidental not split_k_cluster check

2. Minor: prepare_launch's new cluster branch duplicates _enqueue's cluster logic instead of reusing it

Overall assessment

Uh oh!

Uh oh!

cgmillette commented Oct 27, 2026

Uh oh!

jakpiase commented Oct 27, 2026

Uh oh!

cgmillette commented Oct 29, 2026

Uh oh!

jakpiase commented Oct 29, 2026

Uh oh!

jakpiase commented Oct 29, 2026

Uh oh!

jakpiase commented Oct 29, 2026

Uh oh!

jakpiase commented Oct 29, 2026

Uh oh!

jakpiase commented Oct 29, 2026

Uh oh!

cgmillette commented Oct 29, 2026

Uh oh!

cgmillette left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Nov 3, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

jakpiase commented Oct 26, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 26, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 26, 2026 •

edited

Loading

1. Cluster split-K forward's `KernelIdentity.two_stage=True` is a landmine gated only by an incidental `not split_k_cluster` check

2. Minor: `prepare_launch`'s new cluster branch duplicates `_enqueue`'s cluster logic instead of reusing it