Skip to content

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

Merged
jakpiase merged 23 commits into
developfrom
users/jakpiase/rocke/gfx1250-dgrad-depthwise
Oct 30, 2026
Merged

feat(rocke): channel-packed and large-filter depthwise dgrad#14386
jakpiase merged 23 commits into
developfrom
users/jakpiase/rocke/gfx1250-dgrad-depthwise

Conversation

@jakpiase

@jakpiase jakpiase commented Oct 26, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

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 feat(rocke): wave32 WMMA direct and small-channel forward convolution for gfx1250 #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 feat(rocke): wave32 WMMA stem dgrad for gfx1250 (patch embeddings, small-C stems) #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 feat(rocke): improve conv sweep benchmarks #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 (feat(rocke): wave32 WMMA stem dgrad for gfx1250 (patch embeddings, small-C stems) #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 (feat(rocke): wave32 WMMA stem dgrad for gfx1250 (patch embeddings, small-C stems) #14387 adds its own conv_dgrad_patch_wmma / conv_dgrad_smallc_wmma sections), so this PR's configs keep 179-200.
  • After that merge, feat(rocke): wave32 WMMA stem dgrad for gfx1250 (patch embeddings, small-C stems) #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 (feat(rocke): split-K forward implicit-GEMM convolution on gfx1250 #14393 split-K forward conv, feat(rocke): improve conv sweep benchmarks #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;
    • feat(rocke): improve conv sweep benchmarks #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 (refactor(hip-kernel-provider): split rocKE CMake test registration #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 (refactor(rocke): split the representative-IR golden into per-family shards #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

  • Both engines emit byte-identical IR for every new code path; the parity pair and golden cases are added.
  • Defaults keep every existing configuration and its IR unchanged.
  • New C++ code is in new files where the change allowed it.
  • 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.

jakpiase and others added 2 commits October 26, 2026 14:11
@jakpiase @claude
feat(rocke): channel-packed and large-filter depthwise dgrad
33eb59a 
Depthwise dgrad on wave32 ran one channel per lane with 16-bit loads and
stores and scalar FMAs, and had no kernel at all for filters past 11x11
or for the non-square SLaK filters.

- DirectDepthwiseDgradStreamSpec.cpl (2, 4, 8): every lane owns cpl
  adjacent channels; dY and dX move as one vector per lane and the taps
  run as packed <2 x f32> FMAs over channel pairs. cpl == 1 emits the
  unchanged kernel. A validator bound keeps the packed weights and
  accumulators out of scratch.
- DirectDepthwiseTiledSpec.dgrad: a stride-1 depthwise dgrad is the
  forward correlation of dY with the filter rotated by 180 degrees at
  padding K-1-P, so the output-stationary forward kernel serves it
  through the dgrad kernarg block. New AOT variant
  direct_depthwise_tiled_dgrad for 13x13 .. 31x31 and 51x5 / 5x51.
- DirectConvProblem.pad_w / DirectCaps.kw, pad_w: per-axis padding and
  non-square filters for that variant only; every other variant is
  refused such a problem, and the benchmark accepts -p != -q only for a
  stride-1 depthwise dgrad.
- AOT: cpl knobs for direct_depthwise_dgrad, the new variant's caps and
  knobs.

Tests: library/tests/test_direct_depthwise_dgrad_v2.py (validators, AOT
grid entries, the padding gate, gfx1250 numerics against torch).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
feat(rocke): mirror the depthwise dgrad additions in the C++ engine
3ab81cc 
The C++ engine builds the channel-packed and large-filter depthwise dgrad
kernels byte for byte like the Python builders:

- rocke_direct_depthwise_dgrad_stream_spec_t ports
  DirectDepthwiseDgradStreamSpec, which the engine did not cover yet: spec
  defaults, block_ch, kernel name, validate / is_valid (including the cpl
  rules: 1, 2, 4 or 8, groups % cpl == 0, the packed weight and accumulator
  bounds), the scalar ho-streaming builder and its cpl > 1 form with
  cpl-wide vector I/O and packed <2 x f32> FMAs. New files
  conv_direct_grouped_build_depthwise_dgrad_stream.cpp and
  rocke/instance_conv_direct_grouped_dw_dgrad.h.
- rocke_direct_depthwise_tiled_spec_t gains dgrad: the dgrad kernarg block
  with the extents swapped, the rotated weight rows, padding KH-1-PAD /
  KW-1-PAD_W, dX addressed through its strides; the "dg" name flag, filters
  up to 64 on a side in that mode, and the padding checks of validate().
- rocke_direct_conv_problem_t gains pad_w (-1: PAD) and
  rocke_direct_conv_problem_pad_w; every C++ mirror of DirectConvProblem.Wo
  uses it, as the Python property does.
- Parity: conv_direct_grouped_emit.{py,c} configs 171-188 (scalar and cpl
  2 / 4 / 8 streaming dgrad, fp16 and bf16, stride 1 and 2, wave32 and
  wave64; tiled dgrad 13x13, 31x31, 51x5, 5x51 and a per-axis padding; six
  specs both engines reject).
- Golden: five gfx1250 cases in the representative IR harness (new keys
  only; existing hashes unchanged).
- CTest runs test_direct_depthwise_dgrad_v2.py; convolution.md, index.md and
  SUPPORT_MATRIX.md describe the streaming dgrad, cpl and the dgrad mode.

The depthwise kinds have no pybind route, so the parity pair is the
cross-engine gate.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase
jakpiase requested review from a team as code owners October 26, 2026 15:07
@jakpiase
jakpiase requested a balanced review from Copilot October 26, 2026 15:10
@therock-pr-bot

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

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

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

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

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

🙋 Wish to Override Policy?

@therock-pr-bot

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

Copy link
Copy Markdown

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

Copilot AI reviewed Oct 26, 2026
View reviewed changes

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Unresolved error-path safety and asymmetric-padding correctness issues prevent approval.

4 open issues
What changed in this PR

Extends rocKE’s direct-convolution kernels with channel-packed depthwise backward-data computation and support for large, non-square filters. The new configurations use the ahead-of-time (AOT) cache; dispatcher wiring remains out of scope.

Changes:

  • Adds packed row-streaming and rotated-filter tiled dgrad kernels, with matching C++ implementations.
  • Extends per-axis padding, launch geometry, and AOT identities and configurations.
  • Adds numerical tests, cross-engine parity cases, golden IR entries, and documentation.

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

File Description
platform/​tests/​instances/​rocke_ir_parity_harness.py Registers five representative dgrad cases.
platform/​tests/​golden/​rocke_representative_ir_sha256.json Adds hashes across three LLVM flavors.
platform/​python/​rocke/​instances/​SUPPORT_MATRIX.md Documents both dgrad variants and constraints.
platform/​dsl_docs/​instances/​index.md Lists the extended depthwise specs.
platform/​dsl_docs/​instances/​convolution.md Describes packing, rotated filters, and padding.
platform/​cpp/​instances/​common/​conv_direct_nongrouped.cpp Uses effective column padding for output width.
platform/​cpp/​instances/​common/​conv_direct_grouped_spec_and_validate.cpp Extends defaults, naming, and tiled validation.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_depthwise.cpp Updates depthwise output-width calculation.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_depthwise_tiled.cpp Implements rotated-filter dgrad emission.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_depthwise_spatial.cpp Updates spatial output-width calculation.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_depthwise_dgrad_win.cpp Uses column padding in geometry validation.
platform/​cpp/​instances/​common/​conv_direct_grouped_build_depthwise_dgrad_stream.cpp Adds scalar and packed streaming builders and helpers.
platform/​cpp/​include/​rocke/​instance_conv_direct_wgrad_internal.hpp Updates shared wgrad output-width calculation.
platform/​cpp/​include/​rocke/​instance_conv_direct_grouped.h Exposes column padding and tiled dgrad fields.
platform/​cpp/​include/​rocke/​instance_conv_direct_grouped_dw_dgrad.h Declares streaming dgrad APIs and bounds.
platform/​CMakeLists.txt Registers the new numerical test suite.
library/​tests/​test_direct_kernel_cache.py Restricts forward expectations to forward variants.
library/​tests/​test_direct_depthwise_dgrad_v2.py Tests validation, AOT entries, and GPU numerics.
library/​tests/​parity/​conv_direct_grouped_emit.py Adds Python parity and rejection cases.
library/​tests/​parity/​conv_direct_grouped_emit.c Adds matching C++ parity cases.
library/​kernels/​conv/​kernel_identity.py Matches effective column padding.
library/​kernels/​conv/​direct_variants.py Registers tiled dgrad and padding gates.
library/​kernels/​conv/​direct_aot_jobs.py Adds packed configurations and large-filter capabilities.
library/​kernels/​common/​conv_direct_grouped.py Implements packing, tiled dgrad, and launch geometry.
library/​kernels/​common/​conv_args.py Passes effective column padding into launch arguments.
library/​benchmarks/​common/​benchmark_direct_conv.py Accepts per-axis padding for depthwise dgrad requests.

Review effort: Balanced

Comment thread dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped.py Resolved
Comment thread ...der/rocke/platform/cpp/instances/common/conv_direct_grouped_build_depthwise_dgrad_stream.cpp Resolved
Comment thread dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/benchmark_direct_conv.py Resolved
Comment thread dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped.py Resolved
jakpiase and others added 2 commits October 26, 2026 16:13
@jakpiase @claude
fix(rocke): harden the depthwise dgrad padding and lowering error path
4aa1ba0 
- Name a per-axis padding in DirectDepthwiseTiledSpec kernel names: specs
  that differ only in pad_w bake different column offsets, so the name
  gains a p{PAD}x{PAD_W} part when PAD_W != PAD. Names with one padding
  for both axes are unchanged. Re-bless the one golden case it renames.
- Only -1 is the pad_w sentinel: the tiled spec (both modes), the
  row-streaming dgrad spec and direct_spec_for_problem refuse values
  below -1 instead of reading them as PAD.
- The row-streaming dgrad kernel offsets columns by PAD, so its
  validators refuse a problem whose PAD_W != PAD.
- rocke_direct_depthwise_dgrad_stream_lower_to_llvm zero-initialises its
  builder: when the kernel name overflows, _new() fails before
  rocke_ir_builder_init and the failure path reads and frees the builder.

Both engines change alike. Parity configs 189-192 add a per-axis-padded
tiled dgrad spec and three specs both engines reject. A host-only C++ test
covers the name overflow path, the refusals and the kernel name.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
fix(rocke): keep per-axis padding off the uncached direct conv sweep
864d8a3 
parse_miopen_cmd_direct admits a per-axis padding for a stride-1
depthwise dgrad, which only the rotated-filter tiled kernel computes and
only --run-from-cache serves. The uncached sweep builds the row-streaming
kernel, which offsets columns by PAD, so it now skips such a case with a
pointer to --run-from-cache (and returns 2 when no case is left).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@cgmillette

cgmillette commented Oct 27, 2026

Copy link
Copy Markdown
Contributor

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

PR #14386 — channel-packed and large-filter depthwise dgrad — 🟠 First-round: closer look recommended — byte-identical in both engines, and all 10 numeric cases pass on a gfx1151 GPU, but CI checks none of the new kernels' numerics, which also ship to gfx942 / gfx950

Prior item Status
Copilot (10-26) ×4: per-axis padding missing from the kernel name; a builder freed before init on the long-name path; the uncached sweep admitting PAD_W != PAD; pad_w < -1 accepted 🟢 FIXED in 4aa1ba0b / 864d8a37 (read). The name carries p{PAD}x{PAD_W}, pad_w < -1 is refused, and the sweep refuses per-axis padding. The new host test rocke_conv_direct_dw_dgrad_host covers the long-name error path and passes (ran)
Bots / CI 🟢 Green at 864d8a37: 44 checks, including TheRock gfx94X / gfx110X and the hipDNN superbuilds (read)
Our checks 🟢
• Merge: clean with today's develop (ran).
• Build and byte identity: the native build and 33 ctests pass. Byte identity is GREEN for the conv families at llvm20 / 22 / 23, with conv_direct_grouped at 193 configs and 0 bad. The golden gate passes at all three flavors (ran).
• CPU suites: the direct-conv and kernel-cache suites fail only develop's pre-existing test_generic_rejects_match subtests, with 8 more passes (ran).
• gfx1151 GPU: the PR's 10 numeric cases pass against torch's conv2d_input (relative error < 1e-2), with the test's gfx1250-only gate widened to gfx1151 (ran). That covers the channel-packed stream (cpl 2 / 4 / 8, stride 2, 7x7) and the rotated-filter tiled dgrad up to 31x31 and 51x5 / 5x51.
• rule_scan is clean (ran)
# Sev file:line Finding
N1 🟠 library/tests/test_direct_depthwise_dgrad_v2.py (_gpu_skip) CI never checks these kernels' numerics. The only numeric test needs both torch and gfx1250, and CI has neither, so it skips in rocke_library_gpu_pytest everywhere. The AOT grid ships the new kernels to more than gfx1250: on gfx942 and gfx950 (wave64) it adds 536 channel-packed stream jobs and 167 rotated-filter tiled-dgrad jobs per arch, and gfx1151 / gfx1201 get the same (ran). We found no wave64 numeric run, ours or in the description. → Compute the reference with numpy (a depthwise dgrad is a short loop), and gate on the spec being valid for the running arch, with the wave size taken from the target. The gfx94X lane then checks the wave64 variants on every PR.

The worker's one item is dropped. It doubted that a zero-initialized builder can be freed, but = {} zero-initializes portably, and the host test exercises exactly that path.

Utility: 🟢 Positive — packs channels into the row-streaming depthwise dgrad (vector loads and packed FMAs), and serves depthwise dgrad filters up to 64 a side, including non-square ones, through the rotated-filter mode.

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. On a merge with today's develop the checks covered the native build and ctests, byte identity at three flavors, the golden gate and the direct-conv CPU suites against develop. They also included the PR's numeric cases on a gfx1151 GPU, a count of the AOT grid's new jobs per arch, a check of which arches the new specs validate on, and the CI state 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 cgmillette mentioned this pull request Oct 27, 2026
Merged
1 task
aledudek
aledudek approved these changes Oct 27, 2026
View reviewed changes

@aledudek aledudek left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Added 2 comments, overall LGTM

Comment thread dnn-providers/hip-kernel-provider/rocke/library/kernels/common/conv_direct_grouped.py Resolved
Comment thread dnn-providers/hip-kernel-provider/rocke/library/tests/test_direct_depthwise_dgrad_v2.py Resolved
jakpiase and others added 3 commits October 27, 2026 15:10
@jakpiase @claude
Merge branch 'develop' into users/jakpiase/rocke/gfx1250-dgrad-depthwise
044b384 
Resolve the representative IR golden conflict by keeping both sides' cases
and recomputing each flavor's summary counts.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
fix(rocke): gate spilling packed depthwise dgrad out of the AOT grid;…
8b07304 
… run its GPU test on every arch

The cpl > 1 validity bounds (cpl * KH * KW <= 200, cpl * block_w * KH <= 192)
let through binaries that spill: compiling the packed grid, a third of the
gfx950 entries spill to scratch (all of cpl 2 on 9x9, cpl 4 on 7x7, cpl 8
on 5x5) and about half of the gfx1151 / gfx1201 ones do. Add an arch-aware
policy gate to the AOT grid (_dw_dgrad_packed_spills), calibrated on the
compiled grid: per lane cpl * KH * (2 * KW + block_w / stride) <= 288, and on
gfx1250, which spills only in four-wave blocks, <= 448 there. The gfx942 /
gfx950 entries it drops all overflow the 256 architectural VGPRs; what it
keeps spills at most a few registers on a handful of entries. Validity is
unchanged, so the C++ engine and the IR are untouched; the bound comments in
both engines and the docs now say they are ceilings, not a spill model.

The GPU numerics test no longer needs gfx1250 and torch: it builds each case
at the running arch's wave size, skips a case the arch refuses, and checks
against a numpy reference, so wave64 CI lanes exercise the kernels too.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
Merge remote-tracking branch 'origin/develop' into users/jakpiase/roc…
27bb409 
…ke/gfx1250-dgrad-depthwise

Regenerate the IR golden after merging develop (#14349). This also
refreshes stale kernel_name/bytes metadata for
conv_direct_depthwise_tiled/gfx1250/dgrad_fp16_k51x5_w32 (missing the
_p25x2 padding suffix); its IR sha256 is unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@CongMa13

CongMa13 commented Oct 27, 2026

Copy link
Copy Markdown
Contributor

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

If you have already fixed these, please reply on each comment to say so and resolve the threads.

@jakpiase

jakpiase commented Oct 27, 2026

Copy link
Copy Markdown
Contributor Author

@cgmillette re N1 (CI never checks the new kernels' numerics): addressed in 8b07304. test_direct_depthwise_dgrad_v2.py no longer gates on gfx1250 or torch: the GPU test builds each case at ArchTarget.from_gfx(GPU_ARCH).wave_size, skips (as a subtest) any case the running arch's validator refuses, and checks against a numpy reference (inputs via rocke.numeric.bf16). All 10 cases are valid at wave64 on gfx942 / gfx950, so the gfx94X CI lane now exercises the wave64 variants. The same commit also gates spilling channel-packed entries out of the AOT grid (_dw_dgrad_packed_spills).

@CongMa13 this was the only open item from your list; replied here, no further changes needed.

jakpiase and others added 4 commits October 28, 2026 14:29
@jakpiase @claude
Merge origin/develop into users/jakpiase/rocke/gfx1250-dgrad-depthwise
340c036 
No textual conflicts (platform/CMakeLists.txt auto-merged).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
test(rocke): bless depthwise dgrad conv_direct_grouped configs in dif…
83802de 
…ferential golden

Record the new conv_direct_grouped parity configs 171..192 (13 that both
engines emit; the reject-only specs carry no hash) in llvm_gfx_all.json
for ll[llvm20/llvm22/llvm23] and ir_canonical, via
run_diff.py --python-only --record-golden --only conv_direct_grouped.
Additions only; no existing hash changed or removed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
Merge remote-tracking branch 'origin/develop' into users/jakpiase/roc…
a25c234 
…ke/gfx1250-dgrad-depthwise

Golden summary counters recomputed from the merged cases.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase
Merge remote-tracking branch 'origin/users/jakpiase/rocke/gfx1250-dgr…
48bba2a 
…ad-depthwise' into users/jakpiase/rocke/gfx1250-dgrad-depthwise
@jakpiase

jakpiase commented Oct 29, 2026

Copy link
Copy Markdown
Contributor Author

@CongMa13 re your 10-27 note pointing at #14386 (comment) (N1, "CI never checks these kernels' numerics"): this was answered in #14386 (comment) and fixed in 8b07304. test_direct_depthwise_dgrad_v2.py no longer needs torch or gfx1250: it uses a numpy reference, builds each case at ArchTarget.from_gfx(GPU_ARCH).wave_size, and skips (as a subtest) only cases the running arch's validator refuses, so the gfx94X lane now runs the wave64 variants.

Re-verified after merging today's develop (48bba2a): the test passes on gfx1250 (11 passed, 10 GPU subtests); byte identity (--only conv,gfx1250) is GREEN at llvm20/22/23 with conv_direct_grouped at 193 configs, 0 bad; the representative golden --check and the differential golden --check-golden (ll at all three flavors, and ir_canonical) pass. All inline review threads are resolved; no other open items.

@cgmillette

cgmillette commented Oct 29, 2026

Copy link
Copy Markdown
Contributor

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

PR #14386 — channel-packed and large-filter depthwise dgrad (jakpiase, AICK-2337) — 🟢 Follow-up: all prior items fixed — the numeric cases now run and pass in CI on gfx942 and gfx110X, and the AOT grid drops the spilling channel-packed entries

Prior item Status
Copilot (10-26) ×4: per-axis padding missing from the kernel name; a builder freed before init on the long-name path; the uncached sweep admitting PAD_W != PAD; pad_w < -1 accepted 🟢 FIXED 4aa1ba0b / 864d8a37, unchanged at head. The name carries p{PAD}x{PAD_W} (conv_direct_grouped.py:9196, conv_direct_grouped_spec_and_validate.cpp:2951). The builder is zeroed (…_dgrad_stream.cpp:989), and the host test's name-overflow case passes on the merge (rocke_conv_direct_dw_dgrad_host) (ran). The sweep refuses per-axis padding (benchmark_direct_conv.py:2778), and both engines refuse pad_w < -1 (conv_direct_grouped.py:464, …_dgrad_stream.cpp:68). The name, pad_w and stream-padding tests pass in CI (read). The sweep's own test skips in every lane because the benchmark imports torch; the stream spec's refusal of per-axis padding (conv_direct_grouped.py:7110) is tested and guards the same path (read)
Ours 🟠 N1 (10-27) + aledudek (10-27): CI never checked the new kernels' numerics; the GPU test needed gfx1250 and torch 🟢 FIXED 8b073044. The test now runs on any arch rocke models, builds each case at that arch's wave size and checks it against a numpy reference (test_direct_depthwise_dgrad_v2.py:232, :285, :317). At head, TestDepthwiseDgradGpu::test_cases PASSED in rocke_library_gpu_pytest in both gfx94X shards (gfx942, wave64) and both gfx110X shards (wave32), with no case skipped (read: CI logs). The first pass's run at head on a gfx1151 GPU also passed all 10 (read: its log). All 10 cases are valid on gfx942, gfx950, gfx1151, gfx1201 and gfx1250, so none skips on an arch the AOT grid ships to (ran: probe)
aledudek (10-27): the packed accumulator bound has no stride factor, and many packed AOT binaries spill 🟢 FIXED 8b073044. Validity is unchanged and is now documented as a ceiling in both engines and the docs. An arch-aware gate drops spilling packed entries from the AOT grid (direct_aot_jobs.py:1232, applied at :1638). It drops 258 of the 536 packed bf16 entries on gfx942, gfx950, gfx1151 and gfx1201, including every cpl 2 9x9, cpl 4 7x7 and cpl 8 5x5 entry they named, and only the 28 four-wave entries on gfx1250 (ran: probe). Its test passes in CI (read) and on the merge (ran). The remaining-spill counts are the author's compiles; our compile check built none of this family's configs (ran)
aledudek (10-27) APPROVED at 864d8a37 🟡 Stale — predates 8b073044 (their two items) and 83802de0 (the differential-golden bless); the decision is still REVIEW_REQUIRED (read)
CongMa13 (10-27): our 10-27 comment had no author reply 🟢 ADDRESSED — jakpiase replied on 10-27 and 10-29, pointing to 8b073044; that fix is verified above (read)
Bots / CI 🟢 At 48bba2a0, TheRock Multi-Arch CI is green. gfx94X (Linux) and gfx110X (Windows) ran the hipkernelprovider shards, the sanity check and the hipDNN samples and integration tests. gfx950 and gfx125X (gfx1250) built, but every test job on both was skipped, so neither ran these numerics. The hipDNN superbuilds, pre-commit and the policy bot pass; an earlier, superseded policy-bot run was cancelled. REVIEW_REQUIRED (read)
Our checks 🟢
• Interdiff since 864d8a37: 8b073044 (the spill gate, the test rewrite, and reworded comments and docs) and 83802de0 (the differential-golden bless). The kernel-builder edits in both engines are comments only. Three of the four develop merges conflicted, only in the representative golden's counters; one also refreshed the recorded name and size of the PR's own 51x5 case to carry p25x2 (ran: range and remerge diffs).
• Merge: clean with develop 218c5ec6 (5 behind, 11 ahead), including #14415's llvm24 lane (ran).
• Build and IR: the native build and its 34 ctests pass (one skipped). Byte identity is GREEN at llvm20 / 22 / 23, with conv_direct_grouped at 193 configs and 0 bad, and the library C++/Python parity passes (ran).
• Goldens: the representative gate passes at every flavor. The differential golden is GOLDEN OK at llvm20 / 22 / 23 / 24 and ir_canonical, with 0 unblessed changes and 0 missing (ran). Against develop, the bless only adds conv_direct_grouped configs 171–182 and 189, the 13 non-reject parity specs, to ll llvm20 / 22 / 23 and ir_canonical. No other family changes, and develop's llvm24 → llvm23 alias is untouched, so llvm24 is covered. The representative golden adds 5 gfx1250 cases per flavor plus counters, and no existing hash changes (ran: JSON diff).
• Tests: the PR's two test files give 123 passed, 5 skipped on the merge. Rules 1–3 are clean, and no links or tracker IDs were added (ran)

Net-new: none.

Utility: 🟢 Positive — channel-packed and large-filter depthwise dgrad, with numerics now checked in CI.

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 change since the head we reviewed (864d8a37), including the develop merges' conflict resolutions. On a merge with develop 218c5ec6 they covered the native build and ctests, byte identity and the library parity at three flavors, the representative and differential goldens with a JSON diff of both against develop, and the PR's two test files. They also included a CPU probe of the AOT spill gate and of each GPU case's validity per arch, and the 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 cgmillette mentioned this pull request Oct 29, 2026
Open
@jakpiase @claude
Merge remote-tracking branch 'origin/develop' into users/jakpiase/roc…
5815aa9 
…ke/gfx1250-dgrad-depthwise

Conflict only in the representative golden's summary counters
(develop added conv_wgrad/gfx1250/n1h8c16k32r3_spk4, this branch adds
5 gfx1250 depthwise-dgrad cases); counters recomputed from the merged
cases: ok_cases 310/310/318, gfx1250 32/32/40.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase

jakpiase commented Oct 29, 2026

Copy link
Copy Markdown
Contributor Author

Merged current develop (9d729d3280b) in 5815aa9.

  • Conflict: only the summary counters of platform/tests/golden/rocke_representative_ir_sha256.json. Develop added conv_wgrad/gfx1250/n1h8c16k32r3_spk4 (feat(rocke): split-K wgrad on gfx1250 WMMA #13936) and this PR adds 5 gfx1250 depthwise-dgrad cases. I recomputed the counters from the merged cases: ok_cases 310 / 310 / 318 and gfx1250 32 / 32 / 40 (llvm20 / 22 / 23). Every case hash equals develop's or this branch's, and no case differs between the two.
  • Develop adds no conv_direct_grouped parity configs, so this PR's configs keep their numbers (171-192).

Checks on the merge (CPU only):

  • check_byte_identity.py --only conv,gfx1250: GREEN at llvm20 / 22 / 23. conv_direct_grouped has 193 configs, 0 bad.
  • Representative golden --check: IR parity OK for llvm20, 22, 23 and 24.
  • Differential golden run_diff.py --python-only --check-golden: GOLDEN OK for ll (llvm20 / 22 / 23 / 24) and ir_canonical, with 0 unblessed changes and 0 missing.
  • test_direct_depthwise_dgrad_v2.py + test_direct_kernel_cache.py: 127 passed, 1 skipped (the GPU case, which needs a GPU).
  • black 25.12 --check on the 10 Python files changed: clean.

No GPU tests were run for this merge.

@jakpiase @claude
Merge origin/develop into users/jakpiase/rocke/gfx1250-dgrad-depthwise
68ffa19 
Develop brought #14392 (channel-packed depthwise forward), which also
extends DirectDepthwiseTiledSpec and the wave32 parity block.

- DirectDepthwiseTiledSpec (Python and C++): keep both new fields, dgrad
  and cpl. The kernel name carries the per-axis padding part, the cpl part
  and the ur/dg/bf16 flags. The rotated-filter dgrad mode packs one channel
  per lane, so both engines now refuse dgrad with cpl > 1 (Python and host
  tests added).
- conv_direct_grouped parity: develop's wave32 configs end at 178, so this
  branch's depthwise dgrad configs move from 171..192 to 179..200. The
  differential golden's entries move with them, hashes unchanged; docs and
  the harness comment follow.
- direct_aot_jobs imports, test_direct_kernel_cache wave32 test: union.
- Representative golden re-blessed; every case hash equals develop's or
  this branch's.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase

jakpiase commented Oct 29, 2026

Copy link
Copy Markdown
Contributor Author

Merged current develop (050acb2a694, which brings #14392) in 68ffa19.

Conflicts and resolution:

  • DirectDepthwiseTiledSpec (Python, C++ header, spec/name code): both branches added a field (dgrad here, cpl in feat(rocke): wave32 WMMA direct and small-channel forward convolution for gfx1250 #14392). Kept both. The name keeps this branch's order (per-axis padding part, then develop's cpl{N} part, then the ur/dg/bf16 flags), so cpl == 1 names are unchanged on either side. The rotated-filter dgrad mode has no packed path, so both engines now refuse dgrad with cpl > 1; covered by a new check in test_direct_depthwise_dgrad_v2.py and in the rocke_conv_direct_dw_dgrad_host ctest.
  • conv_direct_grouped_emit.{py,c}: develop's wave32 configs now end at 178, so this PR's depthwise dgrad configs move from 171-192 to 179-200. The differential golden entries moved with them, hashes unchanged. Docs and the harness comment updated to match.
  • direct_aot_jobs.py imports and the wave32 test in test_direct_kernel_cache.py: union of both sides.
  • Representative golden: re-blessed with --write; every one of the 954 case hashes equals develop's or this branch's, none dropped.

Checks on the merge (CPU only):

  • Byte identity --only conv,gfx1250: GREEN at llvm20 / 22 / 23, conv_direct_grouped 201 configs, 0 bad.
  • Representative golden --check: OK at llvm20 / 22 / 23 / 24.
  • Differential golden: GOLDEN OK for ll (llvm20 / 22 / 23 / 24) and ir_canonical, 0 unblessed changes, 0 missing.
  • Native build + ctest: 34/34 passed.
  • pytest: test_direct_depthwise_dgrad_v2.py + test_direct_kernel_cache.py 130 passed, 1 skipped; 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 70 passed, 14 skipped.
  • black 25.12.0 and clang-format 18.1.4 clean on the changed files.
  • No GPU run this round.

@cgmillette

cgmillette commented Oct 29, 2026

Copy link
Copy Markdown
Contributor

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

PR #14386 — channel-packed and large-filter depthwise dgrad (jakpiase, AICK-2337) — 🟡 Follow-up: merge verified, minor padding gap

Prior item Status
Copilot (10-26) ×4: per-axis padding missing from the kernel name; a builder freed before init on the long-name path; the uncached sweep admitting PAD_W != PAD; pad_w < -1 accepted 🟢 FIXED 4aa1ba0b / 864d8a37, unchanged since our last review (ran: interdiff). A per-axis padding still reaches the name (…_f51x5_p25x2_…), and an equal one keeps the old name (ran: probe). The builder is still zeroed (…_dgrad_stream.cpp:989), the sweep still refuses per-axis padding (benchmark_direct_conv.py:2860), and rocke_conv_direct_dw_dgrad_host passes at head (ran)
Ours 🟠 N1 (10-27) + aledudek (10-27): CI never checked the new kernels' numerics 🟢 FIXED 8b073044. The merge only adds the cpl check to the test file (ran: interdiff). At head, TestDepthwiseDgradGpu::test_cases PASSED in rocke_library_gpu_pytest in both gfx94X shards (read: CI logs)
aledudek (10-27): the packed accumulator bound has no stride factor, and many packed AOT binaries spill 🟢 FIXED 8b073044; the gate is unchanged (direct_aot_jobs.py:1619, applied at :2041). It still drops 258 of the 536 packed bf16 entries on gfx942, gfx950, gfx1151 and gfx1201, and 28 on gfx1250 (ran: probe)
aledudek (10-27) APPROVED at 864d8a37 🟢 Still counts: develop keeps approvals across pushes. Since then: 8b073044 (their two items, fixed above), 83802de0 (the differential-golden bless) and the develop merges (read)
CongMa13 (10-27): our 10-27 comment had no author reply 🟢 ADDRESSED — jakpiase replied on 10-27 and 10-29 (read)
jakpiase (10-29 10:11): develop merge 5815aa93; only the representative golden's counters conflicted 🟢 Verified — the PR's own patch changed only in those counters, so 5815aa93 is equivalent to the 48bba2a0 we reviewed (ran: interdiff, remerge-diff)
jakpiase (10-29 12:53): develop merge 68ffa198 (#14392); kept both dgrad and cpl; refuses dgrad with cpl > 1 🟢 Verified.
• Fields: dgrad, then cpl, defaulting to False and 1, in Python (conv_direct_grouped.py:9407), the C struct (instance_conv_direct_grouped.h:525) and its defaults (…_spec_and_validate.cpp:2938) (read).
• Name: both engines append the per-axis padding part, then cpl{N}, then the ur / dg / bf16 flags (conv_direct_grouped.py:9453, …_spec_and_validate.cpp:3002). The forward names at cpl 1, 2 and 4 equal develop's (ran: probe), and the dgrad names and IR equal this PR's before the merge (ran: both goldens). No valid spec carries both parts, because each mode refuses the other's.
• Refusal: the same condition (dgrad && cpl != 1), message and place (right after the stride and padding check) in conv_direct_grouped.py:9486 and …_dgrad_stream.cpp:1070 (read). The new pytest case passes here (ran) and in CI (read), and the new host-ctest case passes here (ran).
• Cache keys: every develop AOT identity of the six depthwise variants is still produced, unchanged, on gfx942, gfx950, gfx1151, gfx1201 and gfx1250 in both dtypes, #14392's packed tiled forward entries included. The PR's 167 tiled dgrad entries per arch are all cpl 1, so the refusal drops none (ran: probe).
• #14392 kept: its six packed configs (171–176) emit the same IR as on develop (ran), and no develop hash changes in either golden (ran: JSON diff)
Bots / CI 🟢 At 68ffa198: the policy bot, pre-commit, clang-tidy, base freshness, Component CI and the hipDNN superbuilds (Linux, Windows) pass; a superseded policy-bot run was cancelled. TheRock Multi-Arch CI built gfx94X, gfx950 and gfx125X. On gfx94X, the sanity check, the hipDNN samples and integration tests, and both hipkernelprovider shards pass (49 ctests each). The gfx950 and gfx125X test jobs are skipped, and Windows gfx110X is still building (14:21Z). Neither shard runs rocke_conv_direct_dw_dgrad_host, which is missing from ROCKE_ENGINE_test_categories_external.yaml, so the C++ engine's checks for this PR, the new cpl refusal among them, run only locally (read). Mergeable; BLOCKED, with a rocke-platform team review requested (read)
Our checks 🟢
• Interdiff since 48bba2a0: 5815aa93 changes only the golden counters. 68ffa198 changes the PR's own patch in the 13 files of the author's list: the spec merge above, the parity configs renumbered from 171–192 to 179–200 with their golden entries, the docs, and the union in the kernel-cache test (ran: interdiff, remerge-diff).
• Merge: the head contains develop 050acb2a (0 behind); develop's newer bbd0df28 touches only rocBLAS (ran).
• Build and IR: the native build and its 34 ctests pass (one skipped). Byte identity is GREEN at llvm20 / 22 / 23 for conv_direct_grouped (201 configs, 0 bad), and the library C++/Python parity passes (ran). 15 of those configs build DirectDepthwiseTiledSpec: develop's forward 57–61 and 168, #14392's packed 173–175, and this PR's rotated-filter dgrad 186–190 and 197 (per-axis padding in 188–190 and 197). Six more are specs both engines reject (62, 178, 191–193, 198).
• Goldens: the representative gate passes at llvm20–24, and the differential golden is GOLDEN OK at every flavor and ir_canonical (0 unblessed changes, 0 missing). Against develop, the PR adds only its 13 configs to each section (179–190 and 197, the old 171–182 and 189 with the same hashes) and 5 representative cases per flavor. The 115 / 163 new unblessed configs are develop's, #14392's 171–176 among them (ran: JSON diff).
• Compile (ROCm clang 24): 160 compiled, 0 failures, 0 new rejections. The 8 rejections without a develop baseline are this PR's reject specs 191–196, 198 and 199; 200, the last one, lies past the check's index range (ran). Of the 8 configs that use scratch, 7 are develop's, with unchanged IR. The 8th is this PR's parity config 185 (177 before the renumbering: row-streaming cpl 4, 7x7, gfx950), and the spill gate keeps that entry out of the AOT grid on gfx942, gfx950, gfx1151 and gfx1201 (ran: single-config compiles, probe).
• Tests: the PR's two test files give 124 passed and 7 skipped here; the skips need torch or a GPU (ran). Rules 1–3 are clean (ran)
# Sev file:line Finding
N1 🟡 kernels/common/conv_fwd_wmma_gate.py:35, conv_direct_dgrad_wmma.py:241; C++ conv_direct_fwd_wmma.cpp:89, conv_fwd_small_wmma.cpp:90, conv_direct_dgrad_wmma.cpp:373 Develop's WMMA conv kernels accept a per-axis padding. With PAD = 1 and pad_w = 0, the validators of #14392's direct_fwd_wmma and fwd_smallc_wmma and of the older direct_dgrad_wmma pass, and the IR is the one for pad_w = -1 (columns offset by PAD). Python's p.Wo now follows PAD_W (22 columns on the fwd config) while their C++ mirrors use PAD (24), so a direct caller silently gets another convolution, the engines disagree, and the description's "every C++ mirror of Wo" no longer holds. Only direct_spec_for_problem refuses it, so no AOT, dispatch or benchmark path reaches it (ran: probe); the dgrad-WMMA site predates this merge → refuse PAD_W != PAD in wmma_fwd_gate and in the dgrad-WMMA shape check, in both engines, as the forward tiled kernel does

Utility: 🟢 Positive — channel-packed and large-filter depthwise dgrad, both engines, merged cleanly with #14392.

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 and a remerge-diff of the PR's own change since the head we reviewed (48bba2a0). On the head, which contains develop 050acb2a, they covered the native build and ctests, byte identity and the library parity at three flavors, both goldens with a JSON diff against develop and the previous head, the PR's two test files, and a ROCm clang 24 compile of conv_direct_grouped. They also included single-config compiles of the PR's and #14392's configs, CPU probes of the tiled spec's names and refusals, of the depthwise AOT cache keys against develop, of the spill gate and of the WMMA kernels' padding checks, and the 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
cgmillette approved these changes Oct 29, 2026
View reviewed changes

@cgmillette cgmillette left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@cgmillette cgmillette mentioned this pull request Oct 29, 2026
Merged
5 of 7 tasks
jakpiase and others added 2 commits October 29, 2026 15:05
@jakpiase @claude
Merge origin/develop into users/jakpiase/rocke/gfx1250-dgrad-depthwise
5a2270e 
Brings in #14385 (strided, dilated and packed-group WMMA direct dgrad).
The only conflict was the summary counters of the representative golden;
re-blessed with --write, and every case hash equals develop's or this
branch's. Develop adds no conv_direct_grouped parity configs and does not
touch the differential golden, so this PR's configs keep 179-200.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
fix(rocke): refuse a per-axis padding in the WMMA convolutions
e6d8a84 
DirectConvProblem.pad_w, added by this PR, makes p.Wo follow PAD_W, but
the WMMA forward (direct, small C, small K) and dgrad kernels offset
columns by PAD and their C++ mirrors size Wo from PAD. With PAD_W != PAD
their validators passed and built another convolution, and the engines
disagreed on Wo. Only direct_spec_for_problem refused it.

Move the row-streaming dgrad's padding check into one_padding_reason
(Python) / rocke_direct_conv_one_padding_reason (C++) and apply it in
wmma_fwd_gate and in the dgrad WMMA validator, in both engines, with the
same message. Covered in test_direct_depthwise_dgrad_v2.py and in the
rocke_conv_direct_dw_dgrad_host ctest, which is now listed in the quick
and ffm-quick CTest categories so CI runs it.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase
jakpiase requested a review from a team as a code owner October 29, 2026 15:10
@jakpiase
jakpiase enabled auto-merge (squash) October 29, 2026 15:12
@cgmillette

cgmillette commented Oct 29, 2026

Copy link
Copy Markdown
Contributor

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

PR #14386 — channel-packed and large-filter depthwise dgrad (jakpiase, AICK-2337) — 🟢 Follow-up: WMMA padding fix verified

Prior item Status
Earlier items (10-26 to 10-29 12:53): Copilot ×4; our 🟠 N1 and aledudek's two items; CongMa13's reminder; the develop merges 5815aa93 and 68ffa198 (#14392) 🟢 FIXED / ADDRESSED / Verified as in our 14:26 post, and unchanged since (ran: interdiff). The only change is that the row-streaming dgrad's padding check is now a call to the shared helper below. It gives the same messages: reject configs 199 and 200 still fail with them (ran)
Our 10-29 🟡 N1: the WMMA convs accepted PAD_W != PAD 🟢 FIXED e6d8a84e. The row-streaming dgrad's check is now the shared one_padding_reason / rocke_direct_conv_one_padding_reason (conv_direct_grouped.py:472, …_dgrad_stream.cpp:1015). It ends wmma_fwd_gate (conv_fwd_wmma_gate.py:65, .cpp:94), which all three forward WMMA kernels call first. In the dgrad-WMMA shape check it follows the PAD check (conv_direct_dgrad_wmma.py:352, .cpp:407). Place, condition and message are the same in both engines, so the C++ Wo helpers, which still use PAD, now agree with Python on every spec they accept (read). With PAD 1, all four Python validators and builders refuse pad_w = 0 ("one padding for both axes (got PAD=1x0)") and pad_w = -2, and pad_w = -1 or 1 keeps develop's IR and name (ran: probe). The new pytest case (test_direct_depthwise_dgrad_v2.py:160, 4 subtests) and host-ctest case (conv_direct_dw_dgrad_host.cpp:229) check the same for each engine and pass (ran)
Our 10-29 Bots / CI note: rocke_conv_direct_dw_dgrad_host not in the CI tier YAML 🟢 FIXED e6d8a84e. It is now in quick and ffm-quick (ROCKE_ENGINE_test_categories_external.yaml:53, :95), next to rocke_wgrad_streamk_host, which the gfx94X shards ran at 68ffa198, and it is registered and installed the same way (read). No CI lane has run it yet (Bots / CI row)
Approvals: aledudek (10-27) at 864d8a37; cgmillette (10-29 14:47) at 68ffa198 🟢 Both still count: develop keeps approvals across pushes. Since the later one, the head gained 5a2270ec and e6d8a84e, both verified here (read)
jakpiase (10-29 15:05): develop merge 5a2270ec (#14385), no comment 🟢 Verified. The only conflict was the representative golden's llvm23 counters, resolved as develop's 331 cases (53 on gfx1250) plus this PR's 5 (ran: remerge-diff). Every case of both sides is kept with its hash, and every flavor's counters recompute exactly (ran: JSON diff)
Bots / CI 🟢 Nothing fails because of this PR, but no Linux test lane has run at e6d8a84e. At 16:34Z: 22 passed, 41 skipped, 1 running, 1 cancelled (a superseded policy-bot run) and 1 failed. The failure is the Linux compiler-runtime stage: its external checkout (git fetch) hung for 10 minutes until the runner's container hook failed, so nothing was built. Every Linux math-libs build and test job was then skipped → re-run that stage. Windows gfx110X math-libs has been building since 15:44Z. The policy bot, pre-commit, clang-tidy and both hipDNN superbuilds pass. The PR's kernels emit the same IR as at 68ffa198, where both gfx94X shards passed their GPU numerics. Auto-merge is on. The PR is BLOCKED: CODEOWNERS gives /dnn-providers/, and so the tier YAML, to hipdnn-reviewers, whose review is requested (read: CI logs, gh)
Our checks 🟢
• Interdiff since 68ffa198: 5a2270ec changes only the golden counters, and e6d8a84e is the 11-file fix above (ran).
• Merge: clean with develop bec4b98a (8 behind, 15 ahead). Its only rocKE commit, #13889 (attention), keeps all its golden hashes in the merge (ran).
• Build and IR: the native build and its 34 ctests pass (one skipped). Byte identity is GREEN at llvm20 / 22 / 23 for conv_direct_grouped (201 configs) and conv_direct_dgrad_wmma (32), and the library parity passes (ran). The forward WMMA families were not in that run; on the C++ side only their gate changed (read).
• Goldens: the representative gate passes (llvm20–24), and the differential golden is GOLDEN OK at every flavor and ir_canonical, with 0 unblessed changes and 0 missing. The PR still adds only its 13 configs per section, with the same hashes as before. The new unblessed configs at llvm23 / 24 and ir_canonical went from 163 to 180 because #14385 grew develop's conv_direct_dgrad_wmma from 7 to 24 emitted configs without blessing the new ones; llvm20 / 22 stay at 115 (ran: JSON diff).
• Compile (ROCm clang 24): conv_direct_dgrad_wmma: 24 compiled, its 8 negative configs rejected on develop too, no spills. conv_direct_grouped: 160 compiled, 0 failures, 0 new rejections. The 9 rejections with no develop baseline are the PR's reject specs 191–196 and 198–200; 200 now counts because the check no longer stops at 200 configs. Still 8 configs spill, as before, and the PR's one, config 185, has the same IR and scratch use (ran).
• Tests: the PR's two test files give 125 passed and 7 skipped (fu2: 124; the new WMMA test is the extra one). With this tree's engine, develop's test_direct_dgrad_wmma.py gives 30 passed and test_direct_fwd_wmma.py 10 passed; each skips only its gfx1250 GPU test (ran). rule_scan is clean (ran)

Net-new: none.

Utility: 🟢 Positive — channel-packed and large-filter depthwise dgrad in both engines; WMMA padding gap closed.

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 and a remerge-diff of the PR's own change since the head we reviewed (68ffa198). They also covered a CPU probe of the four WMMA validators against develop, the new host ctest, develop's WMMA test files, both goldens against develop and the previous head, and single-config emits and a compile of the reject and spilling configs. CI logs at head were read. On the merge with develop bec4b98a the checks also covered the native build and ctest, byte identity and the library parity at three flavors, both goldens, the PR's tests and a ROCm clang 24 compile. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

jakpiase and others added 4 commits October 29, 2026 19:14
@jakpiase @claude
Merge origin/develop into users/jakpiase/rocke/gfx1250-dgrad-depthwise
8f53114 
Conflict: summary counters of the representative IR golden. Develop
added 2 gfx942 cases and this branch 5 gfx1250 cases; the counters are
recomputed from the merged cases (ok_cases 314 / 314 / 338).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase
Merge branch 'develop' into users/jakpiase/rocke/gfx1250-dgrad-depthwise
8c2ec27
@jakpiase @claude
Merge origin/develop into users/jakpiase/rocke/gfx1250-dgrad-depthwise
b8bd509 
Brings #14387 (wave32 WMMA stem dgrad). The only conflict was the
llvm23 summary counters of the representative IR golden: develop added
4 stem-dgrad cases and this branch 5 depthwise-dgrad cases. Cases and
expected failures are the union of both sides with their hashes; the
summary counters are recomputed from them.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
fix(rocke): stem dgrad WMMA kernels refuse a per-axis padding
0b566b2 
#14387's dgrad_patch_wmma and dgrad_smallc_wmma take a DirectConvProblem
and offset columns by PAD, so like the other WMMA convolutions their
validators now refuse PAD_W != PAD and pad_w < -1, through the shared
one_padding_reason / rocke_direct_conv_one_padding_reason, in both
engines. The check follows each kernel's PAD check. Emitted IR is
unchanged. The pytest and rocke_conv_direct_dw_dgrad_host cases cover
both kernels.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@cgmillette

cgmillette commented Oct 30, 2026

Copy link
Copy Markdown
Contributor

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

PR #14386 — channel-packed and large-filter depthwise dgrad (jakpiase, AICK-2337) — 🟢 Follow-up: golden conflict resolved exactly

Prior item Status
Earlier items (10-26 to 10-29 16:40): Copilot ×4; aledudek's two items; our 🟠 N1 and CongMa13's reminder; our 🟡 N1 (WMMA per-axis padding) and tier-YAML note; the develop merges through 5a2270ec 🟢 FIXED / ADDRESSED / Verified as in our 16:40 post, and unchanged since: the seven earlier non-merge commits are identical (ran: range-diff)
Our 10-29 16:40 Bots / CI note: no CI lane had run rocke_conv_direct_dw_dgrad_host yet 🟢 Done in CI: at 0b566b2e both gfx94X hipkernelprovider shards (gfx942) ran it and test_direct_depthwise_dgrad_v2.py, including the GPU numerics case and the new stem case, and pass (read: CI logs)
Approvals: aledudek (10-27) at 864d8a37; cgmillette (10-29 14:47) at 68ffa198 🟢 Both still count: develop's rules keep approvals across pushes. Every later commit is verified here or in our 16:40 post (read: gh)
jakpiase (10-29 19:14 to 10-30 04:37): develop merges 8f53114b, 8c2ec278, b8bd5090 (#14387) 🟢 Verified. 8c2ec278 is clean. The other two conflict only on the representative golden's summary counters (after #14357's two gfx942 f64 GEMM cases and #14387's four llvm23 stem-dgrad cases); each is resolved as the sum of both sides, and no merge changes anything else (ran: remerge-diff). Against develop 311cb323, the golden at head keeps all 955 develop case entries and 56 expected failures unchanged, #14387's stem cases and #13889's 21 re-blessed attention hashes included. It adds only this PR's 5 cases per flavor, identical to e6d8a84e, and every flavor's counters recompute exactly (ok_cases 314 / 314 / 342, gfx1250 34 / 34 / 62). The differential golden is develop's plus this PR's unchanged 13 configs per section (ran: JSON diff)
jakpiase (10-30 04:49) 0b566b2e: #14387's stem dgrad WMMA kernels refuse a per-axis padding 🟢 Verified. dgrad_patch_wmma and dgrad_smallc_wmma offset columns by PAD, like the kernels of our 10-29 N1. The shared check now follows each kernel's patch or PAD check, at the same place and with the same message in both engines (conv_dgrad_patch_wmma.py:219 / .cpp:264, conv_dgrad_smallc_wmma.py:217 / .cpp:297), and both engines' builders validate first (read). They are the only DirectConvProblem users added since e6d8a84e (ran: git grep). The new cases (test_direct_depthwise_dgrad_v2.py:196, conv_direct_dw_dgrad_host.cpp:260) pass (ran). IR is unchanged: byte identity is GREEN for both families, their golden hashes are develop's, and no config is newly rejected (ran)
Bots / CI 🟢 Nothing fails. At 06:31Z: 38 check runs passed, 49 skipped, 1 running and 1 cancelled (a superseded policy-bot run). Both gfx94X hipkernelprovider shards pass all 50 ctests on gfx942, #14387's stem tests included; only test_benchmark_sweep_refuses_per_axis_padding skips there (no torch). The gfx950 and gfx125X test lanes are skipped; their math-libs builds pass. The policy bot, pre-commit, clang-tidy, Component CI and both hipDNN superbuilds pass. The Windows gfx110X math-libs build passed at 06:28Z; its test stage has just started, and the required Multi-Arch CI Summary waits for it. The PR is BLOCKED: that check is pending, and CODEOWNERS gives the tier YAML (/dnn-providers/) to hipdnn-reviewers, whose review is requested. Auto-merge (squash) is on (read: gh, CI logs)
Our checks 🟢 Merge with develop 2d1e830e is clean; develop's only later commit, #13471, touches composablekernel alone (head now 2 behind, 19 ahead). On the merge the native build and its 34 ctests pass (one skipped), and byte identity is GREEN at llvm20 / 22 / 23, conv_direct_grouped (201 configs) and both stem families (9 each, llvm23) included. The library parity and the representative golden gate (llvm20–24) pass, and the differential golden is GOLDEN OK at every flavor and ir_canonical, with 0 unblessed changes and 0 missing; its 115 / 180 new configs are develop's own. The PR's tests give 126 passed and 7 skipped (16:40: 125; the stem case is new). ROCm clang 24 compiles stem 6 + 6, dgrad-WMMA 24, forward-WMMA 20 and conv_direct_grouped 160 configs with 0 failures and 0 new rejections; still 8 conv_direct_grouped configs spill, as before. rule_scan is clean; no links added (ran)

Net-new: none.

Utility: 🟢 Positive — channel-packed and large-filter depthwise dgrad in both engines; per-axis padding refused elsewhere.

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

Follow-up review: a local qwen3.6 model made the independent first pass, and the orchestrator (Claude Opus 5.5) verified it. The checks were a range-diff and remerge-diffs since the head we reviewed (e6d8a84e), JSON diffs of both goldens against develop and that head, and the CI logs at head. On the merge with develop 2d1e830e they also covered the native build and ctest, byte identity and the library parity at three flavors, both goldens, the PR's tests and a ROCm clang 24 compile. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

jakpiase and others added 2 commits October 30, 2026 12:40
@jakpiase @claude
Merge origin/develop into users/jakpiase/rocke/gfx1250-dgrad-depthwise
f959e76 
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>
@jakpiase @claude
fix(rocke): key direct problems on their column padding
23b898c 
The direct predicate reads the column padding as PAD_W, which pad_w
overrides, but DIRECT_PROBLEM_CAPABILITY_ATTRS from #14379 left pad_w
out. Two problems differing only in pad_w then shared a capability key,
so a cache filter could answer a per-axis padding with the verdict for
an equal one. test_direct_problem_key_covers_every_extent_the_predicate_reads
caught it after the merge; a new test pins the separation directly.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@cgmillette

cgmillette commented Oct 30, 2026

Copy link
Copy Markdown
Contributor

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

PR #14386 — channel-packed and large-filter depthwise dgrad (jakpiase, AICK-2337) — 🟠 Follow-up: conflicts with #14494, re-merge needed

Prior item Status
Earlier items (10-26 to 10-30 06:33Z): Copilot ×4; aledudek's two items; our 🟠 N1 and CongMa13's reminder; our 🟡 N1 (WMMA per-axis padding) and tier-YAML note; the stem-dgrad padding refusal 0b566b2e; the develop merges through b8bd5090 🟢 FIXED / ADDRESSED / Verified as in our 06:33Z post, and unchanged since: the PR's own patch is identical in 37 of its 40 files, and the other three are the two rows below (ran: interdiff). rocke_conv_direct_dw_dgrad_host, which both gfx94X shards ran at 0b566b2e, passes here (ran)
Approvals: aledudek (10-27) at 864d8a37; cgmillette (10-29 14:47) at 68ffa198 🟢 Both still count: develop's rules keep approvals across pushes, and the decision is APPROVED. Both later commits are verified below (read: gh)
jakpiase (10-30 12:40): develop merge f959e76d (#14393, #14379) 🟢 Verified. Only the representative golden's counters conflicted, at llvm20 / 22. At llvm23 both sides had moved ok_cases from 337 to 342 (#14393's 5 split-K cases, this PR's 5), so git kept 342 without a conflict; the resolution sets 347 there and 318 at llvm20 / 22 (ran: remerge-diff). Against develop 32aed870 the golden keeps all 968 develop cases and 58 expected failures unchanged, #14393's 13 new cases included, and adds only this PR's 5 cases per flavor, identical to 0b566b2e. Every flavor's counters recompute exactly. The differential golden is develop's, unchanged, plus this PR's 13 configs per section, as before (ran: JSON diff)
jakpiase (10-30 12:40) 23b898c8: #14379's direct capability key left out pad_w 🟢 FIXED 23b898c8. KernelCache.compatible() memoises its answers under that key (kernel_cache.py:481), and the predicate reads the column padding as PAD_W (kernel_identity.py:581). Without pad_w in the key (:770), a --run-from-cache sweep could hand a per-axis-padding dgrad the kernels picked for the same shape at another column padding. #14379's guard test and the new one (test_capability_keys.py:262, :285) pass, 16 of 16 (ran). Develop is unaffected: its DirectConvProblem has one PAD, which the key holds. #14379's results and early-stop key already separates the paddings, through ConvGeometry.from_direct_problem's pW = PAD_W (read)
Bots / CI 🟢 Nothing fails. At 14:02Z, 78 check runs: 32 passed, 42 skipped, 3 running and 1 cancelled (a superseded policy-bot run). The policy bot, pre-commit, clang-tidy, Component CI, both hipDNN superbuilds, the gfx94X, gfx950 and gfx125X math-libs builds, and the gfx94X sanity check, hipDNN samples and integration tests pass. Both gfx94X hipkernelprovider shards are running, so no rocKE test lane has finished at this head yet (both passed at 0b566b2e); Windows gfx110X is still building. GitHub reports the PR CONFLICTING since #14494 landed (13:49Z), so auto-merge (squash, on) waits for a re-merge (N1); hipdnn-reviewers, the tier YAML's code owners, are still requested (read: gh)
Our checks 🟠 Merge with develop cd0d5ae0 (3 behind, 21 ahead) has one conflict, a single hunk in platform/tests/CMakeLists.txt (N1). The tier YAML and platform/CMakeLists.txt auto-merge, and both goldens merge as the head's: develop's are byte-identical to 32aed870's, so nothing needs re-blessing (ran: merge-tree). On the head itself, which contains develop 32aed870: the native build and its 34 ctests pass (one skipped), and byte identity is GREEN at llvm20 / 22 / 23 for 14 conv families, conv_direct_grouped (201 configs) and both stem families included. The library parity and the representative golden gate (llvm20–24) pass. The differential golden is GOLDEN OK at every flavor and ir_canonical, with 0 unblessed changes and 0 missing; its 111 / 176 new configs are develop's own (read: another PR's check on develop 32aed870 gives the same). The tests the PR adds or changes, now including test_capability_keys.py, give 142 passed and 7 skipped (06:33Z: 126 and 7). ROCm clang 24 compiles stem 6 + 6, dgrad-WMMA 24, forward-WMMA 20 and conv_direct_grouped 160 configs with 0 failures and 0 new rejections. The same 8 conv_direct_grouped configs spill as before: 7 of develop's and this PR's parity config 185 (row-streaming cpl 4, 7x7, gfx950), whose hashes are unchanged and which the AOT spill gate keeps out of the shipped grid. rule_scan is clean (ran)
# Sev file:line Finding
N1 🟠 platform/tests/CMakeLists.txt:518 Conflicts with develop since #14494 (13:49Z). #14494 moved every C++ test registration out of tests/CMakeLists.txt into tests/cmake/*.cmake, so this PR's rocke_conv_direct_dw_dgrad_host block is the merge's only conflict (ran: merge-tree). Taking develop's side would drop the host test from the build while the tier YAML still lists it, and keeping the block in the orchestrator fails #14494's test_cmake_test_registration.py → merge develop, take its include-only tests/CMakeLists.txt, and move the block unchanged into tests/cmake/InstanceTests.cmake after rocke_moe_fused_mega_spec_guards. Keep the auto-merged tier YAML lines, and leave both goldens as they are: neither conflicts

Utility: 🟢 Positive — channel-packed, large-filter depthwise dgrad in both engines; cache key covers column padding.

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

Follow-up review: a local qwen3.6 model made the independent first pass, and the orchestrator (Claude Opus 5.5) verified it. The checks were an interdiff and a remerge-diff since the head we reviewed (0b566b2e), JSON diffs of both goldens against develop and that head, a read of #14379's cache keys against the new column padding, the capability-key tests at head, a merge-tree against develop cd0d5ae0, and the CI state at head. On the head, which contains develop 32aed870, they also covered the native build and ctest, byte identity and the library parity at three flavors, both goldens, the PR's tests and a ROCm clang 24 compile. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

@cgmillette
cgmillette disabled auto-merge October 30, 2026 14:09
@cgmillette

cgmillette commented Oct 30, 2026

Copy link
Copy Markdown
Contributor

Automated merge-queue status (rocKE merge monitor)

Hi @jakpiase, I turned auto-merge off on this PR for now. Before a rocKE PR lands, the rocKE merge monitor checks it on top of today's develop branch, and this version would not land cleanly:

  • our latest review of this exact version found problems to fix first, and they are still open

The open items from our review, each with its fix:

What to do: address the open items in our latest review (#14386 (comment)). After that, turn auto-merge back on; the monitor re-checks every push.

Nothing else changed: the monitor does not approve, reject or merge PRs. It only turned auto-merge off to keep develop green. This comment is updated in place.

jakpiase and others added 2 commits October 30, 2026 15:31
@jakpiase @claude
Merge origin/develop into users/jakpiase/rocke/gfx1250-dgrad-depthwise
4bfa7ca 
#14494 moved every C++ test registration out of
platform/tests/CMakeLists.txt into tests/cmake/*.cmake, which made this
branch's rocke_conv_direct_dw_dgrad_host block the merge's only conflict.
CMakeLists.txt is now develop's orchestrator unchanged, and the block moves
verbatim (re-indented to file scope) into cmake/InstanceTests.cmake, between
the same two neighbours it sat between before. Its YAML tier entries are
unchanged and still apply, since the labels run after InstanceTests.cmake.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@jakpiase @claude
Merge origin/develop into users/jakpiase/rocke/gfx1250-dgrad-depthwise
ddf46f1 
#14679 split the representative-IR golden into per-family shards under
platform/tests/golden/rocke_representative_ir_sha256/ and deleted the
single file this branch had added cases to. The old file goes; this
branch's delta against the merge base (5 cases per flavor, no other
change) is applied to the conv_direct and conv_direct_depthwise_tiled
shards with their hashes unchanged, written the way the harness writes a
shard. The shards differ from develop's by those additions only.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@cgmillette

cgmillette commented Oct 30, 2026

Copy link
Copy Markdown
Contributor

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

PR #14386 — channel-packed and large-filter depthwise dgrad (jakpiase, AICK-2337) — 🟡 Follow-up: re-merge verified, CI running

Prior item Status
Earlier items (10-26 to 10-30 14:06Z): Copilot ×4; aledudek ×2; CongMa13; our 10-27/10-29 N1s; stem refusal; pad_w key; merges to f959e76d 🟢 FIXED / Verified, unchanged
• fixes 4aa1ba0b, 864d8a37, 8b073044, e6d8a84e, 0b566b2e, 23b898c8
• patch unchanged outside the two moves below (ran: interdiff)
• host ctest, PR tests pass on the merge (ran)
Approvals: aledudek (10-27) at 864d8a37; cgmillette (10-29) at 68ffa198 🟡 Stale
• aledudek's predates 8b073044, 83802de0, e6d8a84e, 0b566b2e, 23b898c8
• cgmillette's predates e6d8a84e, 0b566b2e, 23b898c8
• both still count: decision APPROVED (read: gh)
Our 10-30 N1: tests/CMakeLists.txt conflicts with #14494 🟢 FIXED 4bfa7ca0
• orchestrator = develop's; block moved verbatim to InstanceTests.cmake (ran: remerge-diff)
• #14494's registration test 5 ✓; host ctest ✓ on merge (ran)
• tier YAML still lists it: quick, ffm-quick (read)
jakpiase (10-30 16:27) ddf46f15: golden cases moved to #14679's shards 🟢 Verified
• old .json gone; cases in conv_direct (3), conv_direct_depthwise_tiled (2) (ran)
• hashes = 23b898c8's; develop's entries, schema, aliases unchanged (ran: JSON diff)
• touched shards: 37 cases × llvm20–24, 0 drift (ran)
Our 10-30 14:09Z note: auto-merge off for N1 🟢 Cleared
• nothing 🔴/🟠 open at ddf46f15
• jakpiase may turn auto-merge back on
Bots / CI 🟡 running
• 26 ✓ / 28 skip · 1 running · 2 queued · 0 ✗; 1 cancelled duplicate PR-bot run
• running: math-libs gfx950; queued: Windows gfx110X tests; gfx94X rocKE shards not started
• merge: mergeable; BLOCKED until Multi-Arch CI Summary posts
Our checks 🟢 Clean
• merge vs d3ffe88c: clean, 6 behind (ran)
• build + ctest 33 ✓ / 1 skip; byte identity GREEN ×3, 10 conv families (ran)
• parity ✓; golden gate ✓; differential: PR's 13 per section, NEW = develop's (ran)
• PR tests 142 ✓ / 7 skip; 216 configs compile, spills as before; rules 1–3 clean (ran)

Net-new: none.

Utility: 🟢 Positive — channel-packed, large-filter depthwise dgrad in both engines; follows develop's new layouts.

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

Follow-up review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: merge with develop, build + ctest, byte identity, library parity, representative and differential goldens, PR tests, compile, rules 1–3, #14494's registration test, interdiff and remerge-diffs since 23b898c8, CI at head. Flags whether a closer look is needed; not an approval or rejection.

@jakpiase
jakpiase merged commit 43d958b into develop Oct 30, 2026
95 of 96 checks passed
@jakpiase
jakpiase deleted the users/jakpiase/rocke/gfx1250-dgrad-depthwise branch October 30, 2026 19:32
@cgmillette cgmillette mentioned this pull request Oct 30, 2026
Merged
8 tasks done
@assistant-librarian

assistant-librarian Bot commented Nov 3, 2026

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

@aledudek aledudek mentioned this pull request Nov 4, 2026
Open
1 task
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

Copilot code review Copilot
Copilot review effort
Applies to this pull request for everyone.Learn more about Copilot code review.
Copilot left review comments
@cgmillette cgmillette cgmillette approved these changes
@aledudek aledudek aledudek approved these changes

Assignees

No one assigned

Labels

documentation organization: ROCm project: hip-kernel-provider rocke: platform Touches rocKE platform

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

5 participants

@jakpiase @cgmillette @CongMa13 @aledudek