feat(rocke): gfx1250 direct wgrad kernels with two-stage reduction by jakpiase · Pull Request #14833 · ROCm/rocm-libraries

jakpiase · 2026-11-01T19:23:45Z

Motivation

On gfx1250 (wave32), the direct-conv AOT cache has no weight-gradient kernel for dense / grouped 3x3 layers, for small-channel stem convolutions or for depthwise layers (the depthwise one is also missing on the other wave32 targets, gfx1151 and gfx1201). Those problems fall back to the implicit-GEMM wgrad. For the stems this means an odd GEMM-N (KH * KW * C, e.g. 27) that the split-K kernels cannot take.

This PR adds three direct wgrad kernels (two gfx1250 WMMA kernels and a VALU depthwise kernel) and their Stage 2 reduce helper. Both engines build them, and they are served from the direct-conv AOT cache. It also makes the run-from-cache path robust for problems whose activations exceed the kernels' 2^31-byte addressing, and for kernargs that do not fit.

Technical Details

Kernels (library/kernels/common/conv_direct_grouped.py, C++ twins in platform/cpp/instances/common/conv_direct_grouped_build_{wgrad_image,wgrad_small_c,depthwise_wgrad}.cpp):

Image-per-block direct wgrad (DirectConvWgradSpec(image_per_block=True), variant direct_wgrad on wave32):
- gfx1250 only; one workgroup per image and k x c tile on WMMA 16x16x32;
- stride 1; the validator's register estimate admits square filters up to 4x4 at one tile per wave, so 5x5 and 7x7 are refused; the cache carries 3x3;
- register blocking (wave_tiles_k / wave_tiles_c) and rows_per_sync rows per LDS sync;
- a row split over grid y (row_blocks);
- an optional TDM ring (tdm_stages, llvm23 or llvm24);
- an LDS-staged coalesced epilogue (lds_epilogue).
Small-C stems (DirectConvWgradSpec(small_c=True), variant direct_wgrad_small_c):
- gfx1250 only; a GEMM over the flattened (r, s, c) filter row for ungrouped stems (kpg a multiple of 8, any filter, stride and padding); the cache carries stems with 3 or 4 input channels;
- vectorised X loads (vec_x), sub_steps staging, an optional uniform wave index.
Depthwise VALU wgrad (DirectDepthwiseWgradSpec, variant direct_depthwise_wgrad):
- builds for every target at its wave size; the cache carries it for the wave32 targets (gfx11, gfx12, gfx1250): 3x3, 5x5 and 7x7, stride 1 and 2;
- lanes own channels and accumulate the taps in fp32 registers (r_split splits the filter rows over wave groups);
- an LDS reduction over waves;
- the channel count comes from the p_groups kernarg.
Two-stage reduction:
- each kernel adds into one of ws_replicas fp32 scratch slabs;
- the cached direct_wgrad_reduce helper (WgradReduceSpec, one binary per replica count) folds the slabs into dW;
- direct_kernel_sweep.plan_for builds the two-step plan;
- DirectPlan.zeroed names the scratch the caller zeroes.

Shared lane map: the three builders take the wave32 ds_load_tr16_b128 transpose-read lane map from ConvKOuterFragmentReader.wave32_lane_offsets (C++: rocke_conv_tr_lane_offsets32), which the implicit-GEMM fragment reader now calls too. The gfx1250 dgrad WMMA builders already on develop still compute the same map inline; this PR does not touch them.

Platform: a new IR op tile.buffer_atomic_fadd_f32 (raw_ptr_buffer_atomic_fadd with a cache-policy immediate) in both engines. It is appended to the opcode enum (value 231, after develop's tile.gpu_check_failed), and the intrinsic declaration is added to both lowerers and to the llvm20, llvm22, llvm23 and llvm24 arch-domain data, with a cell for every wired arch.

ABI: rocke_direct_conv_wgrad_spec_t gains fields at its end (image_per_block through uniform_wave), which changes its size, so ROCKE_ABI_VERSION (abi.h) and BINARY_ABI (portable_ir/src/abi.py) move to 11 (develop is at 10). A layout canary in platform/tests/instances/conv_direct_wgrad_refusals_host.cpp pins the struct size to the ABI version.

AOT (library/kernels/conv/direct_aot_jobs.py, direct_variants.py):

direct_wgrad_small_c and direct_depthwise_wgrad join the direct grid.
The wave32 image grid gains the two-stage knobs.
The reduce helper is cached per replica count.
The small-C rows cover:
- the C=3 3x3 / 7x7 / patchify stems;
- the stride-1 3x3 input layers;
- the unpadded 3x3/s2 stem;
- the C=4 3x3 latent-space input convolution.
The 16-wave vec_x=4 / sub_steps=4 small-C tile is cached with uniform_wave off.
Register-staged (tdm_stages=0) image tiles with more than one tile per wave are not cached at 16 waves, where the 256-VGPR cap of a block with more than 8 waves makes them spill; their TDM-staged twins are cached.
gfx950 grids are unchanged.

Run side (benchmark_direct_conv.py, benchmark_implicit_gemm_conv.py):

A plan whose kernargs do not fit (struct.error while binding) is skipped instead of ending the sweep.
wgrad problems past 2^31 bytes of activations run as batch slices that accumulate into one dW, zeroed once:
- implicit GEMM: only the packed-atomic split-K kernels are offered, and only the first slice clears dW;
- direct conv: only the one-step, scratch-free plans that add into the zeroed fp32 dW are offered.

AOT-cache-served; no dispatcher wiring. gfx1250 conv is served from the AOT cache, as it is for gfx950.

Test Plan

Byte identity (platform/tools/check_byte_identity.py, fresh --build-root per flavor) at llvm20, llvm22, llvm23 (full), and llvm24 (--only conv_direct_grouped).
Parity configs 201..241 in library/tests/parity/conv_direct_grouped_emit.{py,c}:
- wave32 wgrad: 201..231;
- depthwise wgrad: 232..238;
- cached small-C stems: 239..241;
- including specs both engines must reject.
IR golden:
- new harness cases conv_direct_wgrad/gfx1250/*;
- new keys in platform/tests/instances/differential/golden/llvm_gfx_all.json;
- rocke_ir_parity_harness.py --check and run_diff.py --check-golden.
C++ host tests: the buffer atomic checks run inside the existing rocke_buffer_rsrc_hip, rocke_opcode_values pins the new opcode, and the new rocke_conv_direct_wgrad_refusals_host runs every refusal case below through the C++ validators.
CPU tests:
- test_direct_wgrad_refusals.py: one reason-match case per refusal check of the image-per-block, small-C and depthwise validators and their arch gates, and a check that every refusal of a numeric knob or an arch names the value it got;
- the refusal text of each case lives in one table, library/tests/direct_wgrad_refusal_reasons.inc, that the C++ host test includes and the Python test parses, so both engines are held to the same text;
- test_direct_depthwise_wgrad_cachepolicy.py (per-arch cache policy of the depthwise atomics) runs in the installed conv pytest lane;
- test_conv_tr_lane_map.py: the shared wave32 lane map, lane by lane, and that each wave32 direct wgrad builder uses it;
- test_direct_kernel_cache.py: AOT rows and reduce helpers, the skipped 16-wave register-staged tiles, two-step plans, stem serving, kernarg skip, wgrad batch chunks, and grid_caches_variant against the enumerated jobs;
- test_direct_conv_correctness.py::TestDirectDepthwiseWgradGate: the depthwise numeric test's arch gate;
- platform/tests/tools/test_arch_domain_artifact.py and platform/tests/core/test_arch_domain_lane.py at ROCKE_LLVM_FLAVOR llvm20, llvm22, llvm23 and llvm24;
- test_kernel_label.py;
- test_builder_signature_contract.py;
- test_conv_batch_and_miopen_parse.py;
- platform/tests/core/test_buffer_atomic_fadd.py.
GPU numeric tests against torch in test_direct_conv_correctness.py: TestDirectConvWgradCorrectness (gfx1250: image-per-block, image row split with TDM, small-C) and TestDirectDepthwiseWgradCorrectness, which runs wherever the AOT grid caches the depthwise kernel (asked of the grid through grid_caches_variant: the wave32 targets, gfx1151 and gfx1201 included), and on gfx950. The wgrad test gate now admits gfx1250; TestDirectWgradAotKernelsCorrectness, whose cases are the gfx950 AOT wgrad kernels, stays gfx950-only.
Fresh cache: --compile-all --directions wgrad --arch gfx1250, then --run-from-cache --verify per family and dtype.

Test Result

Byte identity GREEN at llvm20, llvm22, llvm23 and llvm24.
IR golden and differential golden OK: new keys, plus this PR's own gfx1250 wgrad entries re-blessed after the lane-map refactor (the gfx1250 code objects of those configs are unchanged). Config 235 (gfx950 depthwise wgrad) is re-recorded after its atomics took the per-arch cache policy 0.
CPU tests: no new failures against the develop baseline.
Fresh gfx1250 direct wgrad cache: 1717 kernels per dtype (fp16, bf16), 0 compile failures.
gfx1250 GPU tests (test_direct_conv_correctness.py, test_direct_kernel_cache.py, test_direct_wgrad_specs.py, test_buffer_atomic_fadd.py): all pass.
--run-from-cache --verify smoke, fp16 and bf16: the C=3 stems (3x3 s1 p1 and unpadded 3x3 s2), the C=4 stem, 3x3 and 7x7 depthwise, and a dense 3x3 layer each run a new direct kernel, and every verified kernel passes. A bf16 3x3 layer with activations past 2^31 bytes runs as two batch slices accumulating into dW and verifies.

Submission Checklist

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

JIRA ID : AICK-2878

…T wiring Direct (non-GEMM) weight-gradient kernels for wave32 gfx1250, mirrored in the Python and C++ engines and served from the direct-conv AOT cache: - Image-per-block direct wgrad (direct_wgrad on wave32): one workgroup per image and k x c tile on WMMA 16x16x32, register blocking (wave_tiles_k/c), rows_per_sync, a row split over grid y (row_blocks), TDM staging (tdm_stages), a coalesced LDS epilogue (lds_epilogue). - Small-C stems (direct_wgrad_small_c): a GEMM over the flattened (r, s, c) filter row for 3 input channels, with vectorised X loads (vec_x), sub_steps staging and an optional uniform wave index. - Depthwise VALU wgrad (direct_depthwise_wgrad, DirectDepthwiseWgradSpec): lanes own channels and accumulate the taps in fp32 registers, with an LDS reduction over waves; the channel count comes from the p_groups kernarg. - Two-stage reduction: every kernel adds into one of ws_replicas fp32 scratch slabs; the cached direct_wgrad_reduce helper folds them into dW. The run-from-cache plan zeroes the scratch (DirectPlan.zeroed) and launches both steps; the helper identity goes through _helper_identity. Platform: the tile.buffer_atomic_fadd_f32 IR op (raw_ptr_buffer_atomic_fadd with a cache-policy immediate) in both engines, with a C++ host test. AOT: direct_wgrad_small_c and direct_depthwise_wgrad join the direct grid next to direct_wgrad, the wave32 image grid gains the two-stage knobs, and the reduce helper is cached per replica count. gfx950 grids are unchanged. No dispatcher wiring: gfx1250 conv is served from the AOT cache. Tests: parity configs 201..238 in conv_direct_grouped_emit.{py,c} (wave32 wgrad 201..231, depthwise wgrad 232..238), harness cases conv_direct_wgrad/gfx1250/* with their golden, CPU checks of the AOT grid and launch plans, and GPU correctness tests. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…channels - C=3: the stride-1 3x3 input layers (K 32 / 64) and the unpadded 3x3/s2 stem join the small-C rows; without a small-C binary they fell back to an implicit GEMM with an odd GEMM-N (27) that cannot split. - C=4: 3x3 s1 p1 stems with K 320 / 512 (the latent-space input convolution of the latent diffusion UNets). - The 16-wave vec_x=4 / sub_steps=4 small-C tile is cached with uniform_wave off: with it on, the binary spills SGPRs into VGPR lanes at the 128-VGPR cap and returns an all-zero dW. Grid and caps only; the kernels are unchanged. Parity configs 239..241 pin both engines on the cached stem tiles (C=4 with uniform_wave off, and the two new C=3 stems). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

A byte extent past 2^31 on an i32 kernarg raised struct.error while the run-from-cache sweep bound the plan's launches, and ended the whole direct sweep of the shard. Skip that plan instead, as an unloadable one already is. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

wgrad reduces over the batch, so a problem whose activations pass the kernels' i32 byte addressing was refused. It can still run as consecutive batch slices if every slice adds into the one dW, zeroed once: - implicit GEMM from the cache: _batch_chunks now slices wgrad; a sliced wgrad problem is offered only the packed-atomic split-K kernels (_wgrad_accumulates), dY / X are offset per slice, dW is shared, and only the first slice's launch clears dW (_cached_bind zero_dw). - direct conv from the cache: _batch_chunks slices wgrad too, the slices bind dY and X windows and the whole dW (_chunk_buffers), and only the one-step plans without scratch are kept (_accumulates_dw): their kernel adds into the zeroed fp32 dW, where a fold, a reduce or a scratch slab would overwrite it. Problems that fit keep the single-launch path unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

therock-pr-bot · 2026-11-01T19:27:51Z

✅ 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 · 2026-11-01T19:27:53Z

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

Copilot

🟡 Changes recommended

The public wgrad spec changes size without updating both ABI constants, allowing old callers to pass compatibility checks despite possible out-of-bounds reads or incompatible return layouts.

4 open issues

What changed in this PR

This PR extends rocKE’s direct-convolution AOT cache with gfx1250 weight-gradient support for dense/grouped layers, small-channel stems, and depthwise convolutions, without adding dispatcher wiring.

Changes:

Adds wave32 wgrad support, replicated scratch buffers, and cached second-stage reduction plans.
Introduces a buffer fp32 atomic-add IR operation with matching Python and C++ lowering.
Extends batch-sliced execution, validation tests, parity goldens, and support documentation.

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

File	Description
`rocke/platform/tests/instances/rocke_ir_parity_harness.py`	Adds representative gfx1250 wgrad cases.
`rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json`	Adds differential hashes for new configurations.
`rocke/platform/tests/golden/rocke_representative_ir_sha256/conv_direct_wgrad.json`	Records new IR hashes and expected flavor failures.
`rocke/platform/tests/core/test_opcode_values.cpp`	Pins the appended atomic opcode.
`rocke/platform/tests/core/test_buffer_atomic_fadd.py`	Tests Python atomic lowering and policy validation.
`rocke/platform/tests/core/test_buffer_atomic_fadd.cpp`	Adds corresponding C++ host tests.
`rocke/platform/tests/cmake/CoreTests.cmake`	Registers and installs the atomic host test.
`rocke/platform/python/rocke/instances/SUPPORT_MATRIX.md`	Documents gfx1250 AOT wgrad availability.
`rocke/platform/python/rocke/core/lower_llvm.py`	Lowers buffer atomic adds in Python.
`rocke/platform/python/rocke/core/ir.py`	Adds the atomic builder API and cache-policy constant.
`rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json`	Records atomic intrinsic architecture support.
`rocke/platform/dsl_docs/instances/index.md`	Indexes new wgrad specs and helpers.
`rocke/platform/dsl_docs/instances/convolution.md`	Describes kernels, scratch layout, and AOT serving.
`rocke/platform/cpp/instances/common/conv_direct_grouped_public_entry_glue.cpp`	Routes image and small-C specs to their builders.
`rocke/platform/cpp/instances/common/conv_direct_grouped_build_wgrad.cpp`	Extends wgrad staging and atomics for wave32.
`rocke/platform/cpp/include/rocke/ir.h`	Declares the atomic opcode and builder API.
`rocke/platform/cpp/include/rocke/instance_conv_direct_grouped.h`	Extends public wgrad specs and declarations.
`rocke/platform/cpp/include/rocke/instance_conv_direct_grouped_internal.h`	Adds wave32 builder state and entry points.
`rocke/platform/cpp/core/lower_llvm/mem.cpp`	Implements C++ atomic lowering.
`rocke/platform/cpp/core/lower_llvm/data.cpp`	Declares the LLVM atomic intrinsic.
`rocke/platform/cpp/core/ir/ir_flow.cpp`	Builds atomic operations with policy validation.
`rocke/platform/cpp/core/ir/core_types.cpp`	Registers the atomic name and effectful classification.
`rocke/library/tests/test_kernel_label.py`	Extends label checks to wave32-only variants.
`rocke/library/tests/test_direct_kernel_cache.py`	Tests AOT coverage, reduction plans, and batch slicing.
`rocke/library/tests/test_direct_conv_correctness.py`	Adds gfx1250 numerical wgrad coverage.
`rocke/library/tests/test_conv_batch_and_miopen_parse.py`	Tests wgrad slicing and accumulating-kernel eligibility.
`rocke/library/tests/test_builder_signature_contract.py`	Registers the depthwise wgrad spec contract.
`rocke/library/tests/parity/conv_direct_grouped_emit.py`	Adds Python parity and rejection configurations.
`rocke/library/tests/parity/conv_direct_grouped_emit.c`	Mirrors parity configurations in C++.
`rocke/library/kernels/conv/direct_variants.py`	Adds variants, scratch sizing, and reduction specs.
`rocke/library/kernels/conv/direct_aot_jobs.py`	Adds tuning grids, stem rows, and reduction jobs.
`rocke/library/benchmarks/common/kernel_cache.py`	Adds a cache namespace for wgrad helpers.
`rocke/library/benchmarks/common/direct_kernel_sweep.py`	Constructs and binds two-stage wgrad plans.
`rocke/library/benchmarks/common/benchmark_implicit_gemm_conv.py`	Accumulates oversized wgrad batches across slices.
`rocke/library/benchmarks/common/benchmark_direct_conv.py`	Handles scratch initialization, slicing, and kernarg skips.
`ROCKE_ENGINE_test_categories_external.yaml`	Categorizes the new atomic host test.

_{Review effort: Balanced}

rocke_direct_conv_wgrad_spec_t gained image_per_block through uniform_wave at its end, which changes its size. A caller built against the old header would pass the ABI check while allocating a smaller spec. Bump ROCKE_ABI_VERSION and the Python BINARY_ABI together, and pin the struct size next to the ABI level in conv_direct_dw_dgrad_host so the next change to the struct has to move both. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

A NaN rel_err compared false against `>= tol`, so non-finite depthwise wgrad output passed. Reject anything not below the tolerance, as the other wgrad helpers do. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The buffer_atomic_fadd_f32 docs said an out-of-range voffset is dropped. On gfx1250 a buffer atomic at the out-of-range sentinel faults, while a load there reads 0. Say that callers must guard invalid offsets (as the wgrad kernels do with scf_if), in both engines' docs and the convolution doc. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Link test_buffer_atomic_fadd.cpp into the already categorized rocke_buffer_rsrc_hip binary, whose self-check now calls it, instead of registering a separate rocke_buffer_atomic_fadd CTest. The coverage keeps running in every tier without a new test name, so the provider's test category list goes back to the develop version. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…fer atomics For raw or imported IR the two LLVM lowerers disagreed on the tile.buffer_atomic_fadd_f32 cachepolicy attr: Python truncated 20.9 to 20 while C++ read it as 0, a missing attr became 0, and values above 31 reached the intrinsic. Both lowerers now refuse a missing, non-integer (including bool) or out-of-range attr with a ValueError. Builder-made kernels always carry a valid policy, so emitted IR does not change. Tests: the Python test and the C++ checks (run by rocke_buffer_rsrc_hip) rewrite the attr to missing, 20.9, true, -1 and 32 and expect the refusal. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette · 2026-11-01T20:42:25Z

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

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase) — 🟠 First-round: arch-domain rows missing

Prior item	Status
therock-pr-bot (11-01) PR Description ❌	🟡 OPEN • body has no ISSUE/JIRA ID or closing keyword (read) • `Not ready to Review` stays until the author adds one
Copilot (11-01) wgrad spec grows, ABI not bumped	🟡 OPEN • +12 `int` fields; `spec_default` returns it by value (read) • no user outside the engine build in the provider (read) • develop is 6; takes develop + 1 when it lands, `abi.py` too
Copilot (11-01) ×3: NaN passes dw test; lowerer cachepolicy; OOB-drop doc	🟡 OPEN • `test_direct_conv_correctness.py:2937` `>=` lets NaN pass (read) • lowerers trust the attr; builders enforce 0..31 (read) • all 5 new atomics guarded; docstring only (read)
Bots / CI	🟡 running • fails: Multi-Arch CI Summary ×2 — superseded runs, cancelled; therock-pr-bot — row above • running: gfx94X hipkernelprovider ×2 (queued), hipdnn tests, gfx110X build; 46 ✓ • merge state: blocked; 4 Copilot threads open
Our checks	🟠 1 failing test • merge vs `fce4a134`, `61b3ac0e`: clean (ran) • ctest 35 ✓ / 1 skip; byte identity + goldens ×4 OK (ran) • tests 278 ✓ / 66 skip / 1 ✗ (N1); compile spills develop's only; rules 1–3 clean (ran)

#	Sev	file:line	Finding
N1	🟠	`core/lower_llvm.py:673`, `lower_llvm/data.cpp:291`, `intrinsic_arch_domain.llvm23.json:4583`	New decl key has arch-domain rows at llvm23 only • impact: `test_every_decl_key_is_covered_on_every_wired_arch` fails at llvm20/22/24 (ran); no CI lane runs it • fix: add its rows + `canonical` to those columns (`gen_arch_domain.py` per toolchain); llvm24 measured `ok` ×8 (ran)
N2	🟡	`kernels/conv/direct_aot_jobs.py:2505`, `conv_direct_grouped.py:5200`	LDS-epilogue fallback caches 32 spilling 16-wave image binaries • impact: 256 VGPRs + scratch (ran: ROCm clang 24); `direct_aot_jobs.py:664` says the estimate drops spills • fix: count the 16-wave cap in the estimate, or skip spilling fallbacks

Utility: 🟢 Positive — gfx1250 direct wgrad for dense, stem, depthwise layers; numerics author-run only.

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

First-round review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: merge with develop, native build + ctest, byte identity ×4, goldens, PR tests, arch-domain test per flavor, shipped gfx1250 grid compile, CI at head. Flags whether a closer look is needed; not an approval or rejection.

… and llvm24 columns The raw.ptr.buffer.atomic.fadd.f32 declaration had arch-domain rows on the llvm23 column only, so test_every_decl_key_is_covered_on_every_wired_arch failed on an llvm20, llvm22 or llvm24 host. Add its canonical name and per-arch rows to the other three columns, measured with tools/gen_arch_domain.py --only raw.ptr.buffer.atomic.fadd.f32 on each flavor's toolchain (the llvm20 and llvm22 builds the columns were blessed with). llvm22 and llvm24: ok on every wired arch. llvm20: ok everywhere except gfx1250, which that compiler cannot target (target_unsupported), as for every other key in that column. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…mage tiles Where a block's dW tile does not fit the LDS epilogue, the wave32 wgrad grid falls back to the direct epilogue for the same tiling. For a 16-wave block with two tiles per wave and register-staged rows (tdm_stages=0) that binary spills: a block of more than 8 waves caps a wave at 256 VGPRs, and the register-staged multi-tile body needs more than that. The spec's register estimate counts only the accumulators and fragments, so it let them through, and the grid comment claimed the estimate drops every spilling combination. Add _wgrad_image_spills, a grid filter like the other *_spills predicates, that skips those tilings (32 binaries per dtype). Their TDM-staged fallbacks stay under the cap and are still cached. The spec's validity rules and both engines are unchanged. Every remaining 8- and 16-wave image job compiles without scratch. Test: test_w32_wgrad_grid_skips_the_spilling_16_wave_register_staged_tiles. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-11-01T22:12:49Z

Thanks for the review. The prior items were fixed at 1aa95f0: ABI bumped to 7 (develop is still 6), NaN-safe tolerance check, cachepolicy validated in both lowerers, and the docs now say callers must guard offsets.

N1: fixed in 40e2300. raw.ptr.buffer.atomic.fadd.f32 now has rows and a canonical entry on the llvm20, llvm22 and llvm24 columns, measured with gen_arch_domain.py; llvm20/llvm22 used their blessed compiler builds. The full llvm20/llvm22 columns reproduce under --check, and the coverage test passes on all four flavors.
N2: fixed in a7341a2. The wave32 wgrad grid no longer caches the register-staged (tdm_stages=0) multi-tile 16-wave direct-epilogue fallbacks, which spilled past the 256-VGPR cap of a >8-wave block. Their TDM-staged fallbacks are still cached. Every remaining 8/16-wave image job compiles without scratch. Covered by test_w32_wgrad_grid_skips_the_spilling_16_wave_register_staged_tiles.

cgmillette · 2026-11-01T23:51:03Z

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

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟡 Follow-up: all items fixed; body stale

Prior item	Status
therock-pr-bot (11-01) PR Description ❌	🟢 FIXED in body • body now has `JIRA ID : AICK-2878` (read) • bot re-check passes; `Not ready to Review` removed (read)
Copilot (11-01) wgrad spec grows, ABI not bumped	🟢 FIXED `8bd7ff3e` • header, `BINARY_ABI` and engine all report 7; spec 160 B (ran) • 160-byte canary trips on develop's struct (ran) • develop is 6; takes develop + 1 when it lands
Copilot (11-01) ×3: NaN passes dw test; lowerer cachepolicy; OOB-drop doc	🟢 FIXED `51715450`, `1aa95f09`, `d96f4402` • NaN output now fails the dw check on gfx1151 (ran) • both lowerers refuse 6 bad cachepolicy attrs (ran) • docs: callers must guard offsets (read)
Our 11-01 N1: arch-domain rows at llvm23 only	🟢 FIXED `40e2300f` • coverage test passes at llvm20/22/23/24 (ran) • llvm24 rows match our ROCm clang 24 measurement (ran) • llvm20/22 rows author-measured; not re-run here
Our 11-01 N2: 32 spilling 16-wave image binaries cached	🟢 FIXED `a7341a24` • image grid 544 jobs, 0 spills, fp16 + bf16 (ran) • kept TDM fallbacks: no scratch (ran) • new grid test passes (ran)
Bots / CI	🟡 1 failing, 1 queued • fails: gfx94X hipkernelprovider 2/2 — known golden-sha timeout (#14733) (read) • queued: Windows gfx110X sanity; 37 ✓ • merge: blocked; no approval yet
Our checks	🟢 Clean • merge vs `61b3ac0e`: clean (ran) • ctest 34 ✓ / 1 skip; byte identity ×4 GREEN; goldens ×4 OK (ran) • tests 280 ✓ / 64 skip; compile: no new spills; rules 1–3 clean (ran)

#	Sev	file:line	Finding
N3	🟡	PR body: Platform, Test Result	Body is stale after the follow-up commits • impact: squash commit says llvm23-only arch data, 1749 kernels/dtype; head has 1717 (ran) • fix: update those two lines

Utility: 🟢 Positive — gfx1250 direct wgrad for dense, stem, depthwise layers; numerics author-run only.

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, native build + ctest, byte identity ×4, goldens, PR tests, arch-domain tests per flavor, ABI counters, cachepolicy in both lowerers, NaN check on GPU, shipped gfx1250 grid compile, CI at head. Flags whether a closer look is needed; not an approval or rejection.

The ABI 8 layout canary for rocke_direct_conv_wgrad_spec_t sat in conv_direct_dw_dgrad_host.cpp, a dgrad test. Move it unchanged into conv_direct_wgrad_refusals_host.cpp, the host test of the wgrad validators, and restore the dgrad test to develop's text. List rocke_conv_direct_wgrad_refusals_host in the quick and FFM tiers of ROCKE_ENGINE_test_categories_external.yaml, next to the dgrad host test, so the provider tiers run it and the canary with it. test_cmake_test_registration.py failed on it being untiered. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…check The check that every direct wgrad refusal except the three boolean-flag ones names its value only looked for a digit, which 14 of the 17 texts from before the "(got ...)" rewording already had (">= 1", "wave32", "gfx1250"). Both tables now require the reason to end with "(got ...)". The 11 refusals that state a measured value inside the sentence instead (VGPR and LDS budgets, the threads cap, the depthwise dtype, channel and wave-size checks, the tdm wave count and the wave32 mfma_k gate) are listed with the value text each must carry. All 17 old texts fail the new check in both engines. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-11-04T13:13:13Z

Thanks for the follow-up. All three items are done, in new commits on top of 0e5ec03 (develop at 98e132f, no new develop merge needed):

N9: done in 6ec8661. The four "ok" tails this PR adds (the wave32 early return of rocke_direct_conv_wgrad_validate, the small_c and wave32 returns of rocke_direct_conv_wgrad_is_valid_spec, and the end of rocke_direct_depthwise_wgrad_is_valid_spec) now use develop's rocke_spec_set_reason(reason, reason_cap, "ok") instead of an inline strncpy. The text is unchanged. The seven ROCKE_DCONV_WGRAD_REJECT calls this PR adds stay on the local macro for now. The land order I'm working to is feat(rocke): gfx1250 direct dgrad for small groups at stride 2 and Toeplitz WMMA depthwise dgrad #14834, then this PR, then feat(rocke): gfx1250 WMMA direct fwd: 1-D, rectangular, stride-2 views, small-K, depthwise, small-group and dilated #14835. When feat(rocke): gfx1250 WMMA direct fwd: 1-D, rectangular, stride-2 views, small-K, depthwise, small-group and dilated #14835 merges develop, it moves those seven calls and the four local reason macros onto its shared ROCKE_SPEC_FAIL, which also replaces ROCKE_DCONV_WGRAD_REJECT.
wgrad canary: moved in ecfc11b. The ABI 8 layout canary for rocke_direct_conv_wgrad_spec_t is now in conv_direct_wgrad_refusals_host.cpp (text unchanged), and conv_direct_dw_dgrad_host.cpp is back to develop's text. I also added rocke_conv_direct_wgrad_refusals_host to the quick and FFM tiers of ROCKE_ENGINE_test_categories_external.yaml, next to the dgrad host test. Before that, the provider tiers did not list it, and develop's test_cmake_test_registration.py failed on it. Against feat(rocke): gfx1250 WMMA direct fwd: 1-D, rectangular, stride-2 views, small-K, depthwise, small-group and dilated #14835, conv_direct_dw_dgrad_host.cpp no longer conflicts.
"names a value" check: fixed in 41a887c. Both the Python and C++ tables now require the reason to end with (got ...). Eleven refusals state a measured value inside the sentence instead (the VGPR and LDS budgets, the thread cap, the depthwise dtype, channel and wave-size checks, the TDM wave count and the wave32 mfma_k gate). Those are listed in both tables, each with the value text it must contain. The three boolean-flag refusals are still exempt. All 17 old texts fail the new check in both engines. The old digit check passed 14 of them.

Gates: the C++ build and ctest -L host pass (46 tests, 1 environmental skip as before). Byte identity is GREEN at llvm20/22/23/24, and the goldens are unchanged. The refusal tests (test_direct_wgrad_refusals.py, rocke_conv_direct_wgrad_refusals_host) and the wgrad CPU tests pass, with no new failures against 0e5ec03. The emitted IR did not change.

cgmillette · 2026-11-04T15:01:08Z

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

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟡 Follow-up: fixes hold; minor nits left

Prior item	Status
Our 11-01–11-03 rows (N1–N8, ABI)	🟢 Unchanged • delta reopens none; IR and goldens unchanged (ran) • ABI: develop is 7; takes develop + 1 when it lands (ran) • 14835 takes 8 too: the later one renumbers
Our 11-04 N9: land order with 14835	🟢 FIXED `6ec8661c` • "ok" tails use develop's helper, which 14835 keeps (ran) • 7 old-macro calls: 14835 ports them, landing after this • fallback if 14835 lands first: canary bound now in refusals host test
Our 11-04 Code quality: canary home, value check	🟢 FIXED `ecfc11b9`, `41a887ca` • dgrad host test is develop's; 14835 conflicts 6 → 5 files (ran) • tier entry: develop's tier test failed at `0e5ec03c`, passes now (ran) • each engine catches a dropped tail or value; 17 old texts fail (ran)
Approvals	🟡 None yet • requested: rocke-lib, rocke-platform (code owners) (read) • required: 3 ✓; Multi-Arch CI Summary running (read) • ready for approval again once it passes
Bots / CI	🟡 Running • 40 ✓ · 4 running · 0 ✗: Windows gfx110X rocKE shards, 2 hipDNN tests • Linux gfx94X rocKE shards ✓, refusal host test now included (read: CI logs) • both superbuilds ✓; merge: blocked, review required
Our checks	🟢 Clean • merge vs `961dd763`: clean; one CMake test list shared with 14242, both kept (ran) • build + ctest 46 ✓ / 1 skip; byte identity ×4, goldens, differential, compile ✓ (ran) • tests 414 ✓ / 66 skip; installed 268 ✓; 2 files laneless, as on develop (ran)
Code quality	🟡 Minor • N9 tails, value check and canary move fixed (ran) • refusal texts mirrored per engine, never compared: one-engine drift passes (ran) • `_IntBuilder` copies develop's `rocke.tests.lane_maps` evaluator (ran)

#	Sev	file:line	Finding
N10	🟡	PR body, ABI: paragraph	Body still puts the layout canary in `conv_direct_dw_dgrad_host.cpp` • impact: the squash commits the old location to develop's history • fix: name `conv_direct_wgrad_refusals_host.cpp` there

Utility: 🟢 Positive — gfx1250 direct wgrad kernels; all three of last round's items fixed.

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 and 14835 (trial compile); native build + ctest, byte identity ×4, goldens, differential, PR tests (installed tree), compile, the refusal value checks in both engines (three mutations each), develop's tier registration test at both heads, CI logs, code quality. Flags whether a closer look is needed; not an approval or rejection.

…ke/gfx1250-wgrad-direct Develop now carries #14835, which replaces ROCKE_DCONV_WGRAD_REJECT with the shared ROCKE_SPEC_FAIL / rocke_spec_set_reasonf and takes ABI 8. - The seven ROCKE_DCONV_WGRAD_REJECT calls and the five ROCKE_DW_WGRAD_REJECT calls this branch adds now use ROCKE_SPEC_FAIL(reason, reason_cap, ...); ROCKE_DW_WGRAD_REJECT is removed. ROCKE_WGRAD_IMG_WHY, ROCKE_WGRAD_SC_WHY and ROCKE_DW_WGRAD_WHY keep their names and their return true, with a rocke_spec_set_reasonf body. The refusal texts are unchanged. - The wgrad spec growth is now ABI 9 (abi.h, BINARY_ABI in abi.py, the layout canary in conv_direct_wgrad_refusals_host.cpp). - direct_variants.__all__, benchmark_direct_conv.py (the wgrad chunk filter, then the kernel_map filter) and test_direct_kernel_cache.py keep both sides. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-11-04T19:49:56Z

Thanks for the follow-up.

N10 (body names the old canary location): fixed. The ABI: paragraph now names conv_direct_wgrad_refusals_host.cpp.

#14835 landed first, so this PR made the port from the overlap notice in its develop merge (10e926d):

The seven ROCKE_DCONV_WGRAD_REJECT calls and the five ROCKE_DW_WGRAD_REJECT calls now use the shared ROCKE_SPEC_FAIL(reason, reason_cap, ...), and ROCKE_DW_WGRAD_REJECT is removed.
ROCKE_WGRAD_IMG_WHY, ROCKE_WGRAD_SC_WHY and ROCKE_DW_WGRAD_WHY keep their names and their return true, with a rocke_spec_set_reasonf body.
The refusal texts are unchanged.
In the first conflict hunk, the rocke_wgrad_image_reason check stays, followed by develop's rocke_spec_set_reason(reason, reason_cap, "ok"). In the second, the ROCKE_WGRAD_WAVE_WHY and wave32 mfma_k checks stay; develop's ROCKE_WGRAD_WAVE64_WHY no longer has a user.
ABI: feat(rocke): gfx1250 WMMA direct fwd: 1-D, rectangular, stride-2 views, small-K, depthwise, small-group and dilated #14835's entry 8 is kept, and this PR's wgrad spec growth is now 9: ROCKE_ABI_VERSION, BINARY_ABI in abi.py, and the layout canary's bound in conv_direct_wgrad_refusals_host.cpp. The PR body says 9 as well.
Both sides are kept in direct_variants.py (__all__), benchmark_direct_conv.py (the wgrad chunk filter, then the kernel_map filter) and test_direct_kernel_cache.py.

Gates at 10e926d:

the C++ build and ctest -L host pass (48 tests, 1 environmental skip), including rocke_conv_direct_wgrad_refusals_host and rocke_conv_direct_dw_dgrad_host;
byte identity is GREEN at llvm20/22/23/24, and the representative-IR goldens are OK.

The PR's CPU tests show no new failures: 457 passed. There are 7 failures:

test_conv_abi.py::test_cache_applies_runtime_launch_limits is develop's [rocke] test_conv_abi.py::test_cache_applies_runtime_launch_limits fails on develop since #13782 (kernel admission still refuses the depthwise dgrad that #13782 enabled) #14643;
six test_conv_batch_and_miopen_parse.py cases need torch, which this CPU-only runner blocks.

The branch also merges cleanly with #14832 and #14834.

cgmillette · 2026-11-05T00:55:07Z

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

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟡 Follow-up: port holds; minor nits left

Prior item	Status
Our 11-01–11-04 rows (N1–N9, ABI)	🟢 Unchanged • the two develop merges reopen none: identity ×4, goldens, ctest ✓ (ran) • ABI: develop is 8; takes develop + 1 when it lands (read) • `abi.h`, `abi.py` and the canary bound agree at 9 (ran)
Our 11-04 overlap notice: port if 14835 lands first	🟢 FIXED `10e926d6` • 7 + 5 calls now `ROCKE_SPEC_FAIL`; neither old reject macro left (ran) • the port changes no refusal text: 60,000 specs × 11 arches vs `41a887ca` (ran) • `ROCKE_WGRAD_WAVE64_WHY` renamed in both engines, not dead (ran)
Our 11-04 N10: body names the old canary location	🟢 FIXED (body) • ABI: paragraph names `conv_direct_wgrad_refusals_host.cpp` (read)
Our 11-04 Code quality: texts never compared; `_IntBuilder` copy	🟡 OPEN • both files unchanged; no reply on either (read) • a Python-only reword still passes both tests (ran) • the lane-map test passes on develop's `_Evaluate` as is (ran)
Approvals	🟡 None yet • requested: rocke-lib, rocke-platform (code owners) (read) • required: 3 ✓; Multi-Arch CI Summary pending (read)
Bots / CI	🟡 Running • `077f3e43`: 42 ✓ · 2 running or queued · 0 ✗ • `10e926d6`: 4 math-libs ✗, both attempts: develop's hipCCL break (14976), pre-revert merge ref (read: CI logs) • now all 4 math-libs, both superbuilds, Linux rocKE shards ✓ (read: CI logs)
Our checks	🟢 Clean • merge vs `0c4056ee`: clean, 4 behind; provider tree = head's (ran) • `-Werror` build; ctest 47 ✓ / 1 skip; byte identity ×4; goldens; differential (ran) • tests 431 ✓; torch-only 13 ✓; installed 272 ✓ / 53 skip (ran)
Code quality	🟡 Minor • port onto develop's reason helpers done, one write left: N11 (ran) • engines' refusal texts still never compared: drift passes (ran) • `_IntBuilder` still copies develop's shared lane-map evaluator (ran)

#	Sev	file:line	Finding
N11	🟡	`conv_direct_grouped_spec_and_validate.cpp:4455`	One guarded `snprintf(reason, …)` left in the new depthwise `_validate` • impact: a local copy of the write 14835 made shared; the PR's other 19 use the helpers • fix: `rocke_spec_set_reasonf(reason, reason_cap, "DirectDepthwiseWgradSpec: %s", why);` (builds under `-Werror`)

Utility: 🟢 Positive — the port landed as asked, texts intact; Linux rocKE shards green.

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: the develop merges' conflict resolutions; a scratch merge with develop: native build + ctest, byte identity ×4, goldens, differential; both validators' refusal texts against the last head; PR tests (native engine, torch, installed tree); CI logs; code quality. Flags whether a closer look is needed; not an approval or rejection.

cgmillette

lgtm

…ke/gfx1250-wgrad-direct # Conflicts: # dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h

Copilot

🟡 Changes recommended

Public builders still accept a known-wrong kernel configuration, and depthwise atomics use an architecture-inappropriate cache policy, leaving critical correctness issues unresolved.

4 open issues

4 resolved since last review

_{Review effort: Max}

- Both small_c validators refuse uniform_wave with more than 8 waves, vec_x=4, sub_steps >= 4 and >= 4 wave tiles: that body spills its SGPR column tables into VGPR lanes and returns an all-zero dW. The cache grid already skipped it; direct builder callers could still reach it. Refusal case added to both tables. - The depthwise wgrad atomics carry CPOL_ATOMIC_CASCADE_DEV on gfx12 targets only and the default policy elsewhere, in both engines: the immediate is the gfx12 th/scope encoding, and gfx11 reads the same bits as glc/slc/dlc. The cache builds this kernel for every wave32 target. Tested per arch in both engines. - IRBuilder.buffer_atomic_fadd_f32 refuses a non-integer (or bool) cachepolicy instead of rounding it to a valid policy. - The image-per-block and small_c wgrad builders copy the validator's reason into the builder error, so rocke_direct_conv_wgrad_lower_to_llvm reports it; the C++ refusal test now checks it through that entry. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette · 2026-11-06T05:10:02Z

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

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🔴 Follow-up: golden not re-recorded; CI red

Prior item	Status
Our 11-01–11-04 rows (N1–N10, port onto develop's helpers)	🟢 Unchanged • `a08c89e` and two develop merges reopen none: identity ×4, ctest ✓ (ran) • reason writes still use develop's helpers, except N11 (read)
Our ABI row: takes develop + 1	🟡 OPEN • develop is 9; takes develop + 1 when it lands: `abi.h`, `abi.py` at 10 ✓ (read) • canary bound still `>= 9` (`conv_direct_wgrad_refusals_host.cpp:38`); body still says 9, develop 8 (read) • fix: bound `>= 10`; body "move to 10 (develop is at 9)"
Our 11-04 Code quality: texts never compared; `_IntBuilder` copy	🟡 OPEN • new uniform-wave refusal is a fourth copy of its text; a Python-only reword passes both tests (ran) • `_IntBuilder` unchanged; develop's exported `IntLaneBuilder` drops in: import it (ran)
Our 11-05 N11: guarded `snprintf(reason, …)` in depthwise `_validate`	🟡 OPEN • unchanged at `conv_direct_grouped_spec_and_validate.cpp:4472`; no reply (read)
Copilot (11-05) ×4: uniform-wave spill, depthwise policy, policy type, lost reason	🟢 FIXED `a08c89ee` • both validators refuse the spill spec; no cached small-C job lost (ran) • depthwise atomics: 20 on gfx1200/1201/1250, 0 on 8 gfx11 targets (ran) • type check and lowering reason tested in both engines; gfx950 IR moved: N12 (ran)
Approvals: cgmillette, AnaghaRaoAMD (11-05)	🟢 APPROVED `077f3e43` ×2 • older head; still counts on develop • not covering `a08c89e`'s four fixes or the ABI 10 move
Bots / CI	🔴 Failing • fails: all 4 rocKE shards — `rocke_differential_golden_ll`, `_canonical`: config 235 (N12); Multi-Arch CI Summary ✗ • Windows 2/2 also: `rocke_pytest` timeout, slow Azure runner (1/2 passed it), not the PR's • running: attempt 2 of the 4 shards, same merge ref; 39 ✓; merge: clean, 7 behind
Our checks	🔴 Differential ✗ • merge vs `10f12f0d`: clean; only [235] fails the golden; build + ctest 52 ✓ / 1 skip (ran) • with 14811's IR change (`be0e8acf`): byte identity ×4, representative golden ✓, tests 412 ✓ (ran) • head: tests 441 ✓, torch-only 6 ✓, unforced 119 ✓; installed 278 ✓ / 53 skip (ran)
Code quality	🟠 Should-fix • new Python cache-policy test runs in no CI lane: N13 (ran) • cascade policy: four TU-local C++ defines; one `ROCKE_CPOL_ATOMIC_CASCADE_DEV` in `ir.h` builds (ran) • still open: texts uncompared, `_IntBuilder`, N11 (ran)

#	Sev	file:line	Finding
N12	🔴	`differential/golden/llvm_gfx_all.json`	`a08c89e` moved config 235's IR (gfx950 depthwise policy 20 → 0); golden not re-recorded • impact: `rocke_differential_golden_ll` and `_canonical` fail on all 4 rocKE shards; the required Multi-Arch CI Summary failed • fix: `run_diff.py --mode ll --python-only --record-golden --only conv_direct_grouped --flavor llvm20,llvm22,llvm23,llvm24`, then `--mode ir --canonical`; moves only [235] (adds develop's unblessed 171–176)
N13	🟠	`platform/CMakeLists.txt:700`	New `test_direct_depthwise_wgrad_cachepolicy.py` runs in no CI lane • impact: the AOT cache builds this kernel in Python; nothing in CI checks its gfx11 policy, only the C++ twin runs (read) • fix: list it after `:700`, as TESTING.md asks of gradient host checks; an undone choice fails it; 4 ✓ installed (ran)

Utility: 🟡 Mixed — review fixes hold in both engines; golden re-record missed, CI red.

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: scratch merges with develop, 14811's IR change included: native build + ctest, byte identity ×4, goldens, differential (its cause traced with a policy shim, the re-record tried on a copy); the review fixes probed in both engines; PR tests (native engine, torch, installed tree, unforced LLVM); CI logs; code quality. Flags whether a closer look is needed; not an approval or rejection.

…ke/gfx1250-wgrad-direct

- Re-record the differential golden (ll for llvm20/22/23/24 and canonical IR): config 235 moved when the depthwise atomics took a per-arch cache policy; develop's unrecorded configs 171-176 are added with it. - Run test_direct_depthwise_wgrad_cachepolicy.py in the installed conv pytest lane. - Raise the wgrad spec layout canary to ROCKE_ABI_VERSION >= 10. - One ROCKE_CPOL_ATOMIC_CASCADE_DEV in rocke/ir.h replaces the four TU-local copies. - The depthwise wgrad validator writes its reason with rocke_spec_set_reasonf. - test_conv_tr_lane_map.py uses the shared IntLaneBuilder. - The refusal texts of the wgrad cases move to one table, library/tests/direct_wgrad_refusal_reasons.inc, that the C++ host test includes and the Python test parses, so a reword in one engine fails. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-11-06T13:01:40Z

Thanks for the follow-up. Everything from this round and the earlier open rows is fixed in 03ddb35, on top of a develop merge (e9cb4c1, clean):

N12 (golden not re-recorded): llvm_gfx_all.json is re-recorded with run_diff.py --mode ll --python-only --record-golden --only conv_direct_grouped --flavor llvm20,llvm22,llvm23,llvm24 and then --mode ir --canonical. Only config 235 moves (gfx950 depthwise policy 20 -> 0), and develop's unrecorded 171-176 come in with it. --check-golden passes in both modes.
N13: test_direct_depthwise_wgrad_cachepolicy.py is now listed in _ROCKE_LIBRARY_CONV_TEST_RELPATHS.
ABI row: the canary bound is now ROCKE_ABI_VERSION >= 10. The body says "move to 10 (develop is at 9)".
N11: the depthwise _validate writes its reason with rocke_spec_set_reasonf(reason, reason_cap, "DirectDepthwiseWgradSpec: %s", why).
_IntBuilder: test_conv_tr_lane_map.py now imports develop's IntLaneBuilder, and the local copy is deleted.
Cascade policy: one ROCKE_CPOL_ATOMIC_CASCADE_DEV in rocke/ir.h replaces the four TU-local defines. The host test's policy check uses it too.
Refusal texts never compared: the 51 texts move to one table, library/tests/direct_wgrad_refusal_reasons.inc (one ROCKE_WGRAD_REFUSAL(id, text) line per case).
- The C++ host test #includes the table, and the Python test parses it. Both drop their own text column.
- Each side also checks that every case has a text and every text has a case.
- Mutation checks: a Python-only reword (r_split must lie in) fails test_wgrad_refuses_with_reason[dw_r_split], and a C++-only reword (ws_replicas must be at least 1) fails img_wsr_lt1 in the host test.

Checks: byte identity GREEN at llvm20, llvm22, llvm23 and llvm24. IR parity golden OK. Build-tree ctest 53/53. The conv pytest lane's file list passes 1066, with 66 skipped. pre-commit is clean.

…ke/gfx1250-wgrad-direct # Conflicts: # dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/core_types.cpp # dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h # dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_opcode_values.cpp

jakpiase · 2026-11-06T14:29:07Z

Develop moved again: #14448 took ABI 10 and opcode 230. The develop merge 6a9aab8 resolves both conflicts by keeping develop's entries and putting this PR's after them:

ROCKE_ABI_VERSION (abi.h) and BINARY_ABI (abi.py) are 11. The history keeps develop's 10 entry and adds this PR's spec growth as 11. The layout canary in conv_direct_wgrad_refusals_host.cpp now requires >= 11.
ROCKE_OP_TILE_BUFFER_ATOMIC_FADD_F32 now follows ROCKE_OP_TILE_GPU_CHECK_FAILED in the enum, the name table and the purity table, and test_opcode_values.cpp pins it at 231.
The body is updated to say "move to 11 (develop is at 10)" and to give the opcode value.

Checks after the merge: byte identity GREEN at llvm20, llvm22, llvm23 and llvm24. IR parity and differential goldens (ll and canonical) OK. Build-tree ctest 58/58. The conv pytest lane plus test_buffer_atomic_fadd.py, test_trap.py and test_abi.py pass. pre-commit is clean.

cgmillette · 2026-11-06T19:35:31Z

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

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟠 Follow-up: conflicts with develop's 14822

Prior item	Status
Our 11-06 N12: golden not re-recorded	🟢 FIXED `03ddb35a` • only [235] moves, plus develop's unrecorded 171–176 (ran) • `--check-golden` passes ll ×4 and canonical; saved probe FIXED (ran)
Our 11-06 N13: cache-policy test in no CI lane	🟢 FIXED `03ddb35a` • listed in the conv lane (`platform/CMakeLists.txt:735`) (read) • 4 ✓ installed; passes on both gfx94X shards (ran, read)
Our ABI row: takes develop + 1	🟢 FIXED `6a9aab83` • `abi.h` keeps develop's 10 line, adds 11; `abi.py` 11 (ran) • canary `>= 11`; opcode 231 in enum, name and purity tables, pin (ran) • develop is 10; still takes develop + 1 when it lands
Our 11-04 Code quality: texts never compared; `_IntBuilder` copy	🟢 FIXED `03ddb35a` • one shared table; a one-engine reword fails the other test (ran) • develop's `IntLaneBuilder` imported; the copy is gone (read)
Our 11-05 N11: guarded `snprintf(reason, …)`	🟢 FIXED `03ddb35a` • `rocke_spec_set_reasonf` at `conv_direct_grouped_spec_and_validate.cpp:4470` (read)
Our 11-06 Code quality: four TU-local cascade-policy defines	🟢 FIXED `03ddb35a` • one `ROCKE_CPOL_ATOMIC_CASCADE_DEV` in `ir.h`; builders and host test use it (read)
Our 11-01–11-04 rows (N1–N10)	🟢 Unchanged • the delta reopens none: identity ×4, ctest ✓ at `6a9aab83`'s tree (ran)
Copilot (11-01, 11-05) ×8	🟢 FIXED, unchanged • 11-01 ×4 fixed on 11-01; 11-05 ×4 in `a08c89ee` • the delta touches none of them (read)
Approvals: cgmillette, AnaghaRaoAMD (11-05)	🟢 APPROVED `077f3e43` ×2 • older head; still counts on develop • not covering `a08c89ee`, `03ddb35a` or the ABI 11 move
Bots / CI	🟢 Green • 45 ✓, 52 skipped, 0 ✗ at `6a9aab83` (ran) • Windows attempt 1: `rocke_pytest` and Test step time limits on slow Azure runners, not the PR's; attempt 2 ✓ (read) • merge: dirty (N14); CI's merge ref predates 14822
Our checks	🟠 Conflict • merge vs `897ba70b`: 1 conflict, `direct_kernel_sweep.py` (N14) (ran) • resolved merge: conv lane + changed tests 1155 ✓ / 35 skip; ABI and golden probes FIXED (ran) • build + ctest 57 ✓, byte identity GREEN ×4, goldens ✓ at `6a9aab83`'s tree; develop since adds no C++ or IR change (ran) • rules 1–3 clean; the `read_text()` hit reads an ASCII table (ran)
Code quality	🟡 Minor • five open quality findings fixed; one-engine rewords fail the other test (ran) • table home, parser and `ir.h` macro fit (read) • `InstanceTests.cmake:124-125` keys on the table, no `else`: a moved table drops the host test silently; key on `../../library/kernels` (ran)

#	Sev	file:line	Finding
N14	🟠	`library/benchmarks/common/direct_kernel_sweep.py:653`	Conflicts with develop since 14822 (`97abc0b1`) • impact: GitHub cannot merge; keeping this PR's `if variant in _DGRAD_FPROP_SINGLE:` raises `NameError` (14822 drops that import): 58 tests fail (ran) • fix: keep this PR's two-stage wgrad block, drop develop's one-step `return`, take develop's `if direct_dgrad_kernarg_direction(variant) == "fwd":`; 1155 ✓ (ran)
N15	🟡	`platform/tests/core/test_opcode_values.cpp:266`	Open PR 14707 also takes opcode 231 and ABI 11 • impact: the two conflict in `ir.h`, `core_types.cpp`, `abi.h` and this test (ran) • in the planned order this PR lands first and keeps 231; 14707 moves to 232 and renumbers • if 14707 lands first, this PR moves to 232 and takes develop + 1

Utility: 🟢 Positive — every open row fixed and verified; one develop conflict to resolve

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

Follow-up review: first pass by qwen3.6 (local), verified by the orchestrator (Claude Opus 5.5). Checks: the develop conflict traced to its commit and resolved on a scratch merge (both one-sided resolutions as controls), the conv lane and changed tests on it (native engine, unforced LLVM), saved ABI and golden probes, the merge's ABI and opcode lines, a trial merge with 14707, CI logs, code quality. Flags whether a closer look is needed; not an approval or rejection.

…ke/gfx1250-wgrad-direct # Conflicts: # dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/direct_kernel_sweep.py

Key the CMake guard of rocke_conv_direct_wgrad_refusals_host on the sibling library tree, as the platform's library installs are, rather than on direct_wgrad_refusal_reasons.inc: a moved table now fails the build instead of dropping the test without a word. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-11-06T20:33:54Z

Thanks for the follow-up.

N14: fixed in the develop merge b45b342, with the resolution you gave. The two-stage wgrad block stays, develop's one-step return is dropped, and the dgrad branch takes develop's if direct_dgrad_kernarg_direction(variant) == "fwd":. _DGRAD_FPROP_SINGLE is gone from the file.
CMake guard (InstanceTests.cmake): fixed in 4921f11. The host test is now guarded on the sibling library/kernels tree, as the platform's library installs are, rather than on the table. A moved table now fails the build instead of dropping the test.
N15: noted. If the other PR lands first, this one moves to opcode 232 and develop + 1 for the ABI in its next develop merge.

Checks at 4921f11:

Byte identity is GREEN at llvm20, llvm22, llvm23 and llvm24.
IR parity and the differential goldens (ll ×4 and canonical) are OK.
Build-tree ctest passes 58/58, and rocke_conv_direct_wgrad_refusals_host is still registered.
The conv lane plus this PR's changed test files pass: 1155, with the torch-only parse tests run with torch.
pre-commit is clean.

cgmillette · 2026-11-07T00:27:59Z

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

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟠 Follow-up: conflicts with develop's 14907

Prior item	Status
Our 11-06 N14: conflicts with develop since 14822	🟢 FIXED `b45b3427` • resolved as proposed; `_DGRAD_FPROP_SINGLE` gone from the file (ran) • saved probe FIXED; plan tests ✓ (ran)
Our 11-06 Code quality: CMake guard keyed on the table	🟢 FIXED `4921f113` • keyed on `library/kernels`, like `platform/CMakeLists.txt:387` (read) • table moved: the build now fails, `fatal error: direct_wgrad_refusal_reasons.inc` (ran) • same move at `6a9aab83`: test dropped silently (ran)
Our ABI row: takes develop + 1	🟠 OPEN again, see N16 • develop is 11 since 14907; this PR still 11 (ran) • takes develop + 1 when it lands
Our 11-06 N15: 14707 also takes opcode 231 and ABI 11	🟡 OPEN, plan agreed • 14707 at `9f23cdc5` now takes ABI 12; 6 conflicting files (ran) • planned order: this PR first; 14707 renumbers • if 14707 lands first: opcode 232, develop + 1 (author)
Our 11-01–11-06 rows (N1–N13, quality)	🟢 FIXED, unchanged • no new push; differential goldens: 0 changed on the merge (ran) • cache-policy test runs in the installed conv lane (ran)
Copilot (11-01, 11-05) ×8	🟢 FIXED, unchanged • no new push touches them (read)
CongMa13 (11-03) unanswered-feedback reminder	🟢 Answered • author replied to every post since; 0 of 8 threads open (read)
Approvals: cgmillette, AnaghaRaoAMD (11-05)	🟢 APPROVED `077f3e43` ×2 • older head; still counts on develop • not covering `a08c89ee`, `03ddb35a`, the ABI move or the 11-06 merges
Bots / CI	🟡 Running • 42 ✓, 0 ✗; Windows gfx110X rocKE shards ✓ ×2 (read) • running: Linux gfx94X rocKE shards ×2, hipDNN samples • merge: conflicts; CI's merge ref predates 14907
Our checks	🟠 Conflict • merge vs `886472ee`: 2 files conflict (N16) (ran) • N16's resolution: goldens ✓ ×4 + canonical; 362 tests ✓ (ran) • build, ctest, byte identity: last clean merge `6a22815c` (ran) • rules 1–3 clean; the `read_text()` hit reads an ASCII table (ran)
Code quality	🟢 Holds • the guard matches the platform's copy-able library guard (read) • a moved table now fails the build, not the coverage (ran)

#	Sev	file:line	Finding
N16	🟠	`platform/cpp/include/rocke/abi.h:107`	Conflicts with develop since 14907 (`d1e9c5ff`) • impact: no merge; develop took ABI 11, this PR's number • fix: develop + 1 in `abi.h`, `abi.py` (merges silently at develop's 11), canary `conv_direct_wgrad_refusals_host.cpp:40` • fix: `benchmark_implicit_gemm_conv.py:4819`: keep the batch-slice filter, then develop's `_filter_wgrad_two_stage_candidates` • that resolution: goldens, tests ✓; saved probes FIXED (ran)

Utility: 🟢 Positive — prior items fixed; one develop merge and a land order left

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

Follow-up review: first pass by qwen3.6 (local), verified by the orchestrator (Claude Opus 5.5). Checks: the develop merge with 14907 resolved and checked (goldens, changed tests, ABI and port probes, syntax of the splice), the 11-06 develop merge's resolution, a moved-table mutant of the CMake guard, byte identity and ctest at the last clean merge, Windows CI logs for the refusals test, a trial merge with 14707, code quality. Flags whether a closer look is needed; not an approval or rejection.

jakpiase and others added 4 commits November 1, 2026 19:10

jakpiase requested review from a team as code owners November 1, 2026 19:23

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

therock-pr-bot Bot added the Not ready to Review label Nov 1, 2026

assistant-librarian Bot added the organization: ROCm label Nov 1, 2026

jakpiase requested a balanced review from Copilot November 1, 2026 19:51

Copilot started reviewing on behalf of jakpiase November 1, 2026 19:51

View session

Copilot AI reviewed Nov 1, 2026

View reviewed changes

jakpiase and others added 5 commits November 1, 2026 20:32

test(rocke): fail the depthwise wgrad check on a NaN error

5171545

A NaN rel_err compared false against `>= tol`, so non-finite depthwise wgrad output passed. Reject anything not below the tolerance, as the other wgrad helpers do. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette mentioned this pull request Nov 1, 2026

feat(rocke): gfx1250 WMMA direct fwd: 1-D, rectangular, stride-2 views, small-K, depthwise, small-group and dilated #14835

Merged

1 task done

jakpiase removed the request for review from a team November 1, 2026 21:45

therock-pr-bot Bot removed the Not ready to Review label Nov 1, 2026

jakpiase and others added 2 commits November 1, 2026 22:02

jakpiase and others added 2 commits November 4, 2026 13:13

jakpiase mentioned this pull request Nov 4, 2026

feat(rocke): async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad #14832

Merged

5 tasks done

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-wgrad-direct

077f3e4

cgmillette approved these changes Nov 5, 2026

View reviewed changes

AnaghaRaoAMD approved these changes Nov 5, 2026

View reviewed changes

tenpercent requested a max review from Copilot November 5, 2026 16:57

Copilot started reviewing on behalf of tenpercent November 5, 2026 16:58

View session

Merge remote-tracking branch 'origin/develop' into users/jakpiase/roc…

d550248

…ke/gfx1250-wgrad-direct # Conflicts: # dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/abi.h

Copilot AI reviewed Nov 5, 2026

View reviewed changes

jakpiase and others added 2 commits November 5, 2026 17:35

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-wgrad-direct

f19b203

jakpiase and others added 3 commits November 6, 2026 11:46

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-wgrad-direct

3029760

Merge remote-tracking branch 'origin/develop' into users/jakpiase/roc…

e9cb4c1

…ke/gfx1250-wgrad-direct

jefyang1 mentioned this pull request Nov 6, 2026

feat(rocke): gfx1250 fp8 GEMM bring-up and K-loop optimization (AICK-2347) #14707

Merged

1 task done

jakpiase and others added 2 commits November 6, 2026 20:33

Merge remote-tracking branch 'origin/develop' into users/jakpiase/roc…

b45b342

…ke/gfx1250-wgrad-direct # Conflicts: # dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/direct_kernel_sweep.py

Conversation

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

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Nov 1, 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!

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

Uh oh!

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!

cgmillette commented Nov 1, 2026

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase) — 🟠 First-round: arch-domain rows missing

Uh oh!

jakpiase commented Nov 1, 2026

Uh oh!

cgmillette commented Nov 1, 2026

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟡 Follow-up: all items fixed; body stale

Uh oh!

jakpiase commented Nov 4, 2026

Uh oh!

cgmillette commented Nov 4, 2026

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟡 Follow-up: fixes hold; minor nits left

Uh oh!

jakpiase commented Nov 4, 2026

Uh oh!

cgmillette commented Nov 5, 2026

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟡 Follow-up: port holds; minor nits left

Uh oh!

cgmillette left a comment

Choose a reason for hiding this comment

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!

cgmillette commented Nov 6, 2026

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🔴 Follow-up: golden not re-recorded; CI red

Uh oh!

jakpiase commented Nov 6, 2026

Uh oh!

jakpiase commented Nov 6, 2026

Uh oh!

cgmillette commented Nov 6, 2026

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟠 Follow-up: conflicts with develop's 14822

Uh oh!

jakpiase commented Nov 6, 2026

Uh oh!

cgmillette commented Nov 7, 2026

PR #14833 — gfx1250 direct wgrad kernels with two-stage reduction (jakpiase, AICK-2878) — 🟠 Follow-up: conflicts with develop's 14907

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

jakpiase commented Nov 1, 2026 •

edited

Loading

therock-pr-bot Bot commented Nov 1, 2026 •

edited

Loading

therock-pr-bot Bot commented Nov 1, 2026 •

edited

Loading