feat(rocke): async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad by jakpiase · Pull Request #14832 · ROCm/rocm-libraries

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

Motivation

The implicit-GEMM wgrad on gfx1250 (wave32 WMMA) ran a single-buffered K loop with no overlap, wrote 16-bit split-K partials as one zero-partner packed atomic per element, lowered the packed fp16 atomic as a fenced atomic, never split fp16 automatically, and had no two-stage (f32 workspace) split-K at all.

This PR gives the gfx1250 WMMA wgrad an async-to-LDS K loop, the cshuffle epilogue, an unfenced packed fp16 atomic, an fp16 split policy, and two-stage split-K, and caches the new variants on a geometry cover, keeping every kernel the cache built before. Every kernel change is mirrored in the C++ engine with byte-identical IR.

gfx1250 conv is served from the AOT cache, as on gfx950. This PR is AOT-cache-served; there is no dispatcher wiring.

Technical Details

Async-to-LDS K loop on WMMA (conv_implicit_gemm_wgrad.py, helpers/loads.py; C++ conv_implicit_gemm_wgrad.cpp, conv_implicit_gemm_conv_kloop_drivers.cpp, conv_implicit_gemm_conv_load_phase.cpp, helpers/loads.cpp)

async_dma on gfx1250 is a one-barrier ping-pong over global_load_async_to_lds (CoalescedTileLoader.load_async_to_lds / rocke_conv_emit_kloop_async_global_lds).
It reuses the synchronous K-outer loaders' slot partition and keeps the K-outer row pad of 8, since the copy writes a per-lane LDS address. The gfx9 buffer_load_lds path keeps its pad of 0.
The copy has no buffer-resource range check, and the wgrad descriptors rely on that check for the reduction tail and the free-axis overhang. So the loader ANDs an unsigned byte_offset < num_bytes (icmp ult, as the buffer resource checks it) into each chunk's valid and zero-fills failing chunks with an LDS store of the LDS tile's element type (fp16 or bf16).
It needs lds_k_outer and a 4-, 8- or 16-byte chunk on both operands.
Its chunk widths come from the channel counts, so both validators reject an explicit vector_size_a / vector_size_b on it with the same reason. The C++ validator now also checks those caps' range and that they divide the channel run, as Python does.
On the synchronous K loop those caps narrow the free-axis load widths in both engines: the C++ builder now applies them as Python's _wgrad_free_axis_vec does, and the rocke_engine binding passes both fields.
pipeline stays mem. unroll_k, prefetch_addr and chiplet_swizzle stay rejected on WMMA, as before.

cshuffle epilogue on WMMA

Both wgrad cshuffle epilogues (direct store and split-K atomic_store) read the WMMA fragment through CShuffleEpilogue.from_grid_op, the path the forward and dgrad WMMA epilogues already use. A 16-bit split-K partial is then written as real <2 x dtype> pairs.
Like MFMA's, that LDS pairing needs an even store vector. Both validators now require one for WMMA cshuffle split-K; the default WMMA epilogue still ignores it.

Packed fp16 atomic (core/lower_llvm.py; C++ core/lower_llvm/mem.cpp)

global_atomic_add_pk_f16 is now lowered as atomicrmw fadd <2 x half> syncscope("agent") monotonic with the device-local memory-model metadata, as the packed bf16 add already is.
The legacy llvm.amdgcn.global.atomic.fadd.v2f16 intrinsic it replaces is compiled as a fenced atomic, with a global_wb / global_inv pair and counter drains around every add. The new form selects the same global_atomic_pk_add_f16 with no fences.
The intrinsic's declaration has no other user, so it is dropped from both engines' declaration tables and from the intrinsic arch-domain columns.
The lowering is arch-independent, so the fp16 atomic kernels of other arches (MFMA fp16 split-K wgrad, fp16 atomic dgrad) change IR too.

fp16 split policy (helpers/split_k.py, Python only; the degree is a host-resolved kernarg with no C++ counterpart)

gfx1250 fp16 dW gets its own policy: fill to 1024 CTAs (2048 grouped), K slices down to 512, and an automatic degree capped at 97, the packed fp16 atomic's accuracy bound (atomic_rounding_degree_cap(10)). An override of the automatic degree (the resolver's override, which the benchmark's ROCKE_CONV_SPLIT_K maps onto) is still held to the 16-bit atomic limit of 128 on the packed atomic; a free reduction moves to the workspace above it instead. An explicit split_k is kept as given (only the gridDim.z budget bounds it).
This reverses the "fp16 cap 1" of the gfx1250 split-K policy. That cap was a symptom of the fenced packed fp16 atomic.
The fill and slice terms were fitted on a degree sweep of the cached split-K binaries. The cap of 97 bounds the worst-case rounding of the fp16 atomic sum: at 128, a worst-case set of partials exceeds the 5e-2 verification tolerance; at 97 it stays within it.
bf16 keeps its accuracy-bound cap, and fp32 keeps the general one.

Two-stage split-K on WMMA

The Stage 1 scratch epilogue decodes the WMMA C fragment through op.c_layout(), as the WMMA atomic epilogue does.
wgrad_two_stage_available is true on wave32 targets whose catalog marks float global atomics as validated (gfx1250).
The workspace reduction accumulates in fp32, so the split policy resolves its degree with the fp32 cap (128), not the fp16 atomic cap.
From the cache on wave32, bf16 also runs the two-stage kernels by default (--two-stage auto).

Benchmark run side (benchmark_implicit_gemm_conv.py)

Async-DMA cache entries are ranked by their recorded chunk widths (async_chunk_a / async_chunk_b), not as 8 wide, and the widest-vector pruning ranks async and synchronous kernels apart, so the async kernels are not pruned by the wider synchronous ones.

AOT (kernels/conv/aot_jobs.py)

The WMMA wgrad cache keeps every kernel it built before: the synchronous default-epilogue kernels, unsplit and packed-atomic split-K, on the whole tile ladder.
On 19 geometries (CACHE_WGRAD_WMMA_GEOMETRIES) it adds:
- the cshuffle epilogue, unsplit and packed-atomic split-K;
- two-stage split-K (ws_replicas = 8);
- the async_dma loop next to the plain loop for every non-two-stage configuration, built once per geometry with its chunk widths recorded in the identity.
The async loop is offered only on targets with global_load_async_to_lds. count_jobs follows the same rule; CPU tests check that the default grid is unchanged, that the new variants stay on the cover, and that the count matches the job walk on gfx1250, gfx1151 and gfx1201.
gfx1250 implicit wgrad: every previously cached identity is kept, plus 912 new jobs per dtype. gfx950 is unchanged.

Docs: README_conv_implicit_gemm_wgrad.md, dsl_docs/instances/convolution.md, SUPPORT_MATRIX.md.

Test Plan

Byte identity (platform/tools/check_byte_identity.py) at llvm20, llvm22, llvm23 and llvm24, each with a fresh build root.
Parity configs 55-70 of conv_implicit_gemm_wgrad_emit.{py,c} cover:
- the async K loop: split-K fp32 dW with partial N tiles, 2x2 warps, and bf16 A/B (config 67, the bf16 zero-fill);
- the cshuffle epilogue: direct fp16, split-K bf16, and async at C=24;
- two-stage: fp16 dW with an odd dW row (C=3), bf16 dW on K-outer 2x2 warps, and fp16 dW over the async loop with one scratch slab;
- explicit vector_size_a / vector_size_b caps (4 / 2) on the synchronous K loop, K-outer and M-outer (configs 69-70);
- rejects both engines must agree on: async at C=3, cshuffle split-K on gfx1201, cshuffle split-K with an odd store vector, two-stage on gfx1151, and an explicit vector_size_a on the async loop (config 68).
rocke_wgrad_streamk_host (ctest -L host): the C++ validator's vector_size_a / vector_size_b reasons, checked against the Python text.
test_conv_wgrad_streamk.py: through the rocke_engine binding, capped synchronous kernels lower byte-identically, a cap wider than the auto width (6 at kpg 12) keeps the auto width, and the cap rejects give Python's reasons.
Representative-IR harness: new gfx1250 async, cshuffle, two-stage and capped K-outer cases.
Golden re-blessed for the new keys and for the fp16-atomic IR change only.
CPU tests:
- test_conv_wgrad_correctness.py: validator gates, two-stage availability, cshuffle store vector;
- test_conv_wgrad_prefetch_addr.py, test_conv_split_k.py, the split-policy dispatch tests (including the fp16 cap of 97 on worst-case partials);
- test_coalesced_load_async_to_lds.py: the async range check is unsigned, and the zero-fill follows the tile type;
- test_conv_abi.py: async kernels survive the widest-vector pruning, and count_jobs matches the walk per target;
- test_kernel_label.py, test_conv_aot_grid_limits.py, the arch-domain tests.
gfx1250 GPU numeric tests against the CPU reference:
- TestWgradWmmaCShuffle;
- TestWgradTwoStageWmmaNumerics: an odd-dW-row stem and a 2x2 stride-2 downsample, fp16 / bf16, plain / K-outer / async;
- the async odd-slice split-K tail;
- the gfx1250 split-K resolver test, whose fp16 cases now run at the resolved degree.
A fresh gfx1250 wgrad cache built with --compile-all --directions wgrad, then --run-from-cache --verify on fp16 / bf16 split-K async, cshuffle and two-stage cases.

Test Result

Byte identity: GREEN at llvm20, llvm22, llvm23 and llvm24.
Representative-IR golden and the differential golden: OK in every flavor.
CPU tests: no new failures against develop.
Fresh gfx1250 wgrad cache, built from this branch into an empty directory: every implicit wgrad job builds, 0 compile failures (the rejected jobs are the same ones develop rejects).
gfx1250 GPU tests (test_conv_wgrad_correctness.py, test_conv_split_k.py, test_conv_wgrad_prefetch_addr.py), compared with the same files on develop: no new failures. The two-stage cases that develop could not build on WMMA now pass, and every new WMMA case passes (cshuffle, cshuffle + async, two-stage plain / K-outer / async, the async odd-slice tail, the split-K resolver). The remaining failures are develop's MFMA-only cases, which fail the same way on develop on gfx1250.
--run-from-cache --verify smoke from the fresh cache, fp16 and bf16 (3x3, 2x2 stride 2, 1x1, a 3-output-channel layer): every verified kernel passes, covering cshuffle, the async loop, packed-atomic split-K and two-stage split-K.
The fp16 atomic change on gfx942 / gfx950 is covered by byte identity and the golden; it was not run on those GPUs here.

Submission Checklist

Mirrored in the C++ engine; byte identity GREEN at llvm20 / llvm22 / llvm23 / llvm24.
Parity configs for every new axis and dtype; golden re-blessed (new keys plus the fp16-atomic keys only).
AOT grid and run side updated; AOT-cache-served, no dispatcher wiring.
Docs updated.
gfx1250 GPU numerics and the cache --verify smoke.

JIRA ID : AICK-2880

…MA wgrad The WMMA wgrad K loop was single-buffered with no overlap: load, barrier, WMMA, barrier. On gfx1250, async_dma now runs a one-barrier ping-pong over global_load_async_to_lds (CoalescedTileLoader.load_async_to_lds): - it reuses the synchronous K-outer loaders for the slot partition and keeps the K-outer row pad of 8, since the copy writes a per-lane LDS address (the gfx9 buffer_load_lds path keeps its pad of 0); - each phase waits s_wait_asynccnt 0 and takes one LDS barrier, then issues the next tile while the current tile's WMMAs run; - the copy has no buffer-resource range check, which the wgrad descriptors rely on for the reduction tail and the free-axis overhang, so the loader ANDs byte_offset < num_bytes into each chunk's valid predicate and zero-fills failing chunks with an LDS store; - it needs lds_k_outer and a 4-, 8- or 16-byte chunk on both operands. Admit epilogue="cshuffle" on WMMA wgrad. Both wgrad cshuffle epilogues (direct store and split-K atomic_store) read the WMMA fragment through CShuffleEpilogue.from_grid_op, the path the forward and dgrad WMMA epilogues already use, so a 16-bit split-K partial is written as genuine <2 x dtype> pairs instead of one zero-partner packed atomic per element. Like MFMA's, that LDS pairing needs an even store vector, which both engines' validators now require for WMMA cshuffle split-K (the default WMMA epilogue still ignores it). pipeline stays "mem", and unroll_k / prefetch_addr / chiplet_swizzle stay rejected on WMMA, as before. Both engines are mirrored (rocke_conv_emit_kloop_async_global_lds and rocke_coalesced_tile_loader_load_async_to_lds in C++). Parity configs 55-62 cover the async loop (fp32 dW split-K, partial N tiles, 2x2 warps), the C=3 async reject, and the cshuffle epilogue (direct fp16, split-K bf16, with async at C=24, the gfx1201 reject, and the odd store-vector reject). Config 53 (gfx1250 bf16 split-K with cshuffle) is now admitted by both engines. Two new representative-IR cases (async, cshuffle) are blessed; no existing hash changes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The legacy llvm.amdgcn.global.atomic.fadd.v2f16 intrinsic is lowered as a fenced atomic: a global_wb / global_inv pair (buffer_wbl2 / buffer_inv on gfx950) and full counter drains around every single add, which serialised the split-K epilogues. Emit the generic agent-scope monotonic atomicrmw fadd <2 x half> with the device-local memory-model metadata, as the packed bf16 add already does; the backend selects the same global_atomic_pk_add_f16 without the fences. Mirrored in the C++ engine (lower_llvm/mem.cpp). The intrinsic declaration has no other user, so it is dropped from both engines' declaration tables and from the intrinsic arch-domain columns. The HIP lowering keeps __builtin_amdgcn_global_atomic_fadd_v2f16, which clang already emits unfenced. Every kernel with a packed fp16 atomic changes its IR: the fp16 split-K wgrad configs and the fp16 atomic dgrad epilogue configs. The parity golden and the representative-IR golden are re-blessed for exactly those keys. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Give gfx1250 fp16 dW its own wgrad split policy: fill to 1024 CTAs (2048 grouped) with K slices down to 512, up to degree 128. bf16 keeps its accuracy-bound policy and fp32 the general one. This reverses the "fp16 cap 1" of the gfx1250 split-K policy. That cap was a symptom of the fenced packed fp16 atomic: with the atomic lowered as a monotonic atomicrmw (previous commit), deeper fp16 splits pay off. The policy was fitted on a degree sweep of the cached split-K binaries, and no fp16 degree failed the 5e-2 verification there. Python only: the C++ engine has no split-policy counterpart (the degree is a kernarg resolved by the host). Update the policy tests and the gfx1250 GPU split-K resolver test, whose fp16 cases now run at the resolved degree. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…k widths An async-DMA binary records no A / B load width (the copy picks its chunk), so the widest-vector filter of the from-cache sweep added None and aborted. Rank such an entry by the chunk widths recorded in its identity (async_chunk_a / async_chunk_b), as a synchronous kernel with the same load widths; fall back to 8 only when none is recorded. Ranking every async binary as 8 wide let it outrank, and so filter out, the widest synchronous kernels of a problem whose channel run caps the loads below 8. Benchmark / run side only; no kernel IR change. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The two-stage scratch-store epilogue decoded the accumulator only through the MFMA atom, so wave32 had no two-stage path: a 16-bit dW with an odd row (a 3-channel stem) could not be split at all, and bf16 split-K stopped at the packed atomic's accuracy cap. Decode the WMMA C fragment through op.c_layout() in the scratch epilogue, as the WMMA atomic epilogue does, and enable two-stage on wave32 targets whose catalog marks float global atomics as validated (gfx1250; the f32 scratch adds are global float atomics). The workspace reduction accumulates in fp32, so the split policy resolves its degree with the fp32 cap. From the cache on wave32, bf16 also runs the two-stage kernels by default (--two-stage auto), next to the packed-atomic ones. The MFMA path's IR is unchanged. Mirrored in the C++ engine (availability gate and the WMMA scratch epilogue). Parity configs 63-66: fp16 dW with an odd dW row (C=3), bf16 dW on 64x64x64 2x2 warps K-outer, fp16 dW over the async K loop with one scratch slab, and the gfx1151 reject. One new representative-IR case. CPU tests cover the availability predicate and the gfx1151 reject; a new gfx1250 GPU test runs two-stage on an odd-dW-row stem and a 2x2 stride-2 downsample, fp16 / bf16, plain, K-outer and async K loops. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

… async variants On wave32 (WMMA) the implicit-GEMM wgrad cache keeps every kernel it built before: the synchronous default-epilogue kernels, unsplit and packed-atomic split-K, on the whole tile ladder. On top of that, on 19 (tile_m, tile_n, tile_k, warp_m, warp_n) geometries (CACHE_WGRAD_WMMA_GEOMETRIES), it adds: - the cshuffle epilogue, unsplit and packed-atomic split-K; - two-stage (f32 workspace, ws_replicas = 8) split-K on the default epilogue, which the two-stage availability on gfx1250 brings in; - the async_dma K loop (global_load_async_to_lds) next to the plain loop for every non-two-stage configuration, built once per geometry with A / B widths left to the copy and its chunk widths recorded in the identity (async_chunk_a / async_chunk_b). Each cover geometry was the fastest of at least one reference wgrad layer, and together they hold the fastest measured time of every layer. count_jobs follows the same rule (and now counts the async loop), and a CPU test pins it: the default grid is unchanged and the new variants stay on the cover. The gfx950 grid is unchanged. The gfx1250 GPU split-K resolver test takes its 2x2-warp tiles from the cover and only the synchronous packed-atomic binaries. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

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

✅ 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:55Z

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

Copilot

🟡 Changes recommended

Unresolved critical defects can cause out-of-bounds async copies near the 2 GiB boundary and fp16 accumulation errors exceeding the accuracy tolerance.

5 open issues

What changed in this PR

This PR extends rocKE’s AOT-cache-served gfx1250 WMMA weight-gradient kernels with asynchronous loading, cshuffle epilogues, and two-stage split-K, with corresponding C++ implementations.

Changes:

Adds async global-to-LDS ping-pong loading and WMMA fragment handling for cshuffle and workspace epilogues.
Updates packed fp16 atomic lowering, automatic split-K policies, cache generation, and benchmark selection.
Expands validation, numerical tests, parity cases, IR goldens, and documentation.

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

File	Description
`platform/tests/instances/rocke_ir_parity_harness.py`	Adds representative async, cshuffle, and two-stage cases.
`platform/tests/instances/differential/golden/llvm_gfx_all.json`	Updates differential IR hashes.
`platform/tests/golden/rocke_representative_ir_sha256/conv_wgrad.json`	Records new variants and fp16 atomic IR changes.
`platform/tests/golden/rocke_representative_ir_sha256/conv_wgrad_two_stage.json`	Adds gfx1250 workspace-epilogue goldens.
`platform/tests/dispatch/dispatch_tests/gemm/test_split_k_wgrad.py`	Tests the revised fp16 split policy.
`platform/python/rocke/instances/SUPPORT_MATRIX.md`	Documents expanded WMMA wgrad support.
`platform/python/rocke/helpers/split_k.py`	Introduces a gfx1250 fp16-specific policy.
`platform/python/rocke/helpers/loads.py`	Adds asynchronous global-to-LDS tile loading.
`platform/python/rocke/core/lower_llvm.py`	Lowers packed fp16 adds through `atomicrmw`.
`platform/python/rocke/core/ir.py`	Updates packed fp16 atomic documentation.
`platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm24.json`	Removes the unused legacy atomic intrinsic entry.
`platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json`	Removes the unused legacy atomic intrinsic entry.
`platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm22.json`	Removes the unused legacy atomic intrinsic entry.
`platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json`	Removes the unused legacy atomic intrinsic entry.
`platform/dsl_docs/instances/convolution.md`	Explains WMMA loading, padding, and epilogues.
`platform/cpp/instances/common/conv_implicit_gemm_wgrad.cpp`	Implements C++ WMMA validation and epilogue changes.
`platform/cpp/instances/common/conv_implicit_gemm_conv_load_phase.cpp`	Routes async copies through the new loader.
`platform/cpp/instances/common/conv_implicit_gemm_conv_kloop_drivers.cpp`	Adds the async ping-pong K-loop driver.
`platform/cpp/include/rocke/instance_conv_implicit_gemm_internal.h`	Declares async context fields and the driver.
`platform/cpp/include/rocke/helper_rocke.helpers.loads.h`	Declares the async tile-loading helper.
`platform/cpp/helpers/loads.cpp`	Implements async copies and zero-fill handling.
`platform/cpp/core/lower_llvm/mem.cpp`	Mirrors packed fp16 `atomicrmw` lowering.
`platform/cpp/core/lower_llvm/data.cpp`	Removes the legacy atomic declaration.
`library/tests/test_conv_wgrad_correctness.py`	Adds WMMA validation and numerical coverage.
`library/tests/test_conv_split_k.py`	Updates automatic fp16 split expectations.
`library/tests/test_conv_abi.py`	Checks cache geometry coverage and job counts.
`library/tests/parity/conv_implicit_gemm_wgrad_emit.py`	Adds Python emission and rejection cases.
`library/tests/parity/conv_implicit_gemm_wgrad_emit.c`	Adds matching C++ engine parity configurations.
`library/kernels/conv/aot_jobs.py`	Generates new variants on the WMMA geometry cover.
`library/kernels/common/README_conv_implicit_gemm_wgrad.md`	Documents new kernels and split-K behavior.
`library/kernels/common/conv_split_k.py`	Updates reduction eligibility and workspace degree selection.
`library/kernels/common/conv_implicit_gemm_wgrad.py`	Implements Python WMMA async and epilogue support.
`library/benchmarks/common/benchmark_implicit_gemm_conv.py`	Updates workspace selection and async-width ranking.

_{Review effort: Balanced}

…n the job count count_jobs counted the async_dma wgrad variants on every wave32 target, but _wgrad_jobs drops them on gfx1151 and gfx1201, which have no global_load_async_to_lds. The count then exceeded what the walk can yield. _wgrad_k_loops now takes the target and offers the async loop only where target.memory.has_async_global_lds holds; both the walk and the count use it. Test: count_jobs equals the wgrad walk on gfx1250, gfx1151 and gfx1201, and only gfx1250 yields async jobs. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The from-cache sweep keeps only the widest vector copy of each kernel kind. An async-DMA wgrad binary records the chunk widths its copies use, while a synchronous binary records its width caps, which its loader can narrow at build time. On the gfx1250 cover geometry 32x64x32 (2x4 warps) the async loop records A=4 / B=8 and the synchronous copy A=8 / B=8, so with both in one kind every async kernel was dropped before it was benchmarked. The kind now includes async_dma, so each loop is pruned to its own widest copy. Benchmark / run side only; no kernel IR change. Test: on that geometry the async cshuffle kernel survives the pruning and the synchronous kernels are still cut to their widest copy. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…c's accuracy bound The gfx1250 fp16 wgrad split policy went up to degree 128. Each packed fp16 atomic add rounds the running sum in dW, so valid inputs can miss the 5e-2 tolerance at that degree: with one 1.0 partial followed by 127 partials of 2**-11 the fp16 sum stays at 1.0, a 5.8% error. The policy now caps fp16 at atomic_rounding_degree_cap(10) = 97, the same accuracy bound bf16 uses with 7 fraction bits. The cap binds only the packed-atomic reduction: the f32 workspace degree is already resolved as fp32 dW and keeps the general cap of 128. Explicit and overridden degrees are unchanged. No kernel IR or AOT job changes; an automatic fp16 dW degree on gfx1250 that would exceed 97 is now 97. Tests: the worst-case fp16 partials exceed the tolerance at 128 and stay within it at the automatic degree (the pointwise N=64, 32x32, C=K=16 shape on a 64x128x64 tile); the policy and resolver expectations move from 128 to 97, and the workspace reduction of the same kind of shape still resolves to 128. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…c LDS loader CoalescedTileLoader.load_async_to_lds makes the buffer-resource range check itself, since global_load_async_to_lds has none. Two fixes, in both engines: - The check compared the i32 byte offset with num_bytes signed. A padded reduction tail near 2 GiB can give the offset 0x80000000, which is negative as an i32 and so passed the check, issuing an out-of-bounds copy instead of a zero-fill. The compare is now unsigned (icmp ult), as the resource's is. - The C++ mirror zero-filled out-of-range chunks with an fp16 vector whatever the operands, while Python uses the loader's element type, so bf16-input async kernels diverged between the engines. The C++ loader carries no element type, so the zero now takes the destination LDS tile's (f16 or bf16; anything else is an error, as this port loads 16-bit elements). New parity config 67 (case 55 with bf16 operands) covers the bf16 zero-fill in both engines. The async wgrad goldens move with the compare: configs 55, 56, 60 and 65 and the new 67 in the differential golden (re-recording the family also adds the gfx1250 configs 49-51, which only the llvm24 section had recorded), and the async case in the representative conv_wgrad shard. Test: test_coalesced_load_async_to_lds.py checks the unsigned compare and the element type of the zero store for fp16 and bf16 tiles. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

cgmillette · 2026-11-01T20:33:35Z

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

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase) — 🟡 First-round: minor fixes, CI running

Prior item	Status
Reviewers	🟢 None yet • no human reviews or threads (read)
therock-pr-bot (11-01) PR Description check ❌	🟡 OPEN • body has no JIRA/ISSUE ID or closing keyword (read) • author adds one; bot then drops "Not ready to Review"
Copilot (11-01) ×2 high: fp16 degree past accuracy bound; signed async bound check	🟡 OPEN, valid • fp16 degree up to 128 > `atomic_rounding_degree_cap(10)` = 97; worst case 5.8%, typical ~1% (ran: probe) • signed check misfires only at byte offsets ≥ 2^31 (read) • fix: cap fp16 at 97; `cmp_ult` in both engines
Copilot (11-01) ×3: async pruned; count_jobs; C++ bf16 zero-fill	🟡 OPEN, valid • async never ranked on 2 of 19 cover geometries (ran) • `count_jobs` +76 on gfx1151/gfx1201; progress only (ran) • C++ zero-fill is f16 in bf16 tiles; no parity case (read)
Bots / CI	🟡 Running • fails: Multi-Arch CI Summary ×2 — runs cancelled by label re-runs; therock-pr-bot — body • running: Multi-Arch math-libs builds (gfx94X, gfx110X); 41 ✓
Our checks	🟢 Clean • merge vs `fce4a134`: clean, 0 behind; ctest 34 ✓ / 1 skip (ran) • engine-wide fp16 atomic: byte identity GREEN, 24 families × 4 flavors; goldens OK; ISA unfenced (ran) • tests 468 ✓ / 66 skip, 1 ✗ = develop issue 14643; new AOT grid 0 spills; gfx1151 GPU 40 ✓ (ran)

Net-new: none.

Utility: 🟢 Positive — gfx1250 wgrad gains async K loop, cshuffle, two-stage; fp16 atomics unfenced.

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 (every family, 4 flavors), goldens, PR tests, compile + new AOT grid, fp16-atomic ISA, gfx1151 GPU run, Copilot items, CI at head. Flags whether a closer look is needed; not an approval or rejection.

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

Thanks. The five Copilot items listed there are fixed at cc12d0f: the fp16 auto cap is 97 (6e49bda); the async check is unsigned and the zero-fill uses the tile type (cc12d0f); async pruning (4d5e200); count_jobs (d02b387). The earlier Multi-Arch failures came from a cancelled run that was superseded by a re-run on the same head.

cgmillette · 2026-11-01T22:41:17Z

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

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: fixes verified, CI running

Prior item	Status
Reviewers	🟢 None yet • no human reviews yet (read)
therock-pr-bot (11-01) PR Description check ❌	🟢 FIXED • body now carries AICK-2880; bot check ✓ (read) • bot removed "Not ready to Review"
Copilot (11-01) ×2 high: fp16 degree past accuracy bound; signed async bound check	🟢 FIXED `6e49bdaa`, `cc12d0fb` • fp16 auto degree ≤ 97 on 24,646 sampled shapes; was 128 (ran) • `icmp ult` in both engines; IR delta is that compare only (ran) • new loader test fails on the old compare (ran)
Copilot (11-01) ×2: count_jobs; C++ bf16 zero-fill	🟢 FIXED `d02b3875`, `cc12d0fb` • count_jobs = walk on gfx1151/gfx1201; was +76 (ran) • bf16 config 67: bf16 zero stores, byte-identical in both engines (ran)
Copilot (11-01): async kernels pruned before benchmarking	🟢 Mostly `4d5e2006` • async ranked apart from sync; the PR's test passes (ran) • full sweeps: 32x64x32 (2x4), 32x256x32 (1x8) async still lose to 8-wide async (ran)
Bots / CI	🟡 Running • fails: Multi-Arch CI Summary — cancelled run, superseded by label re-run (read) • running: Multi-Arch math-libs (gfx94X, gfx110X), tests not started; 30 ✓ • merge: blocked (no approval yet; checks running)
Our checks	🟢 Clean • merge vs `61b3ac0e`: clean, 1 behind; build + ctest 34 ✓ / 1 skip (ran) • byte identity GREEN, 24 families × 4 flavors; goldens OK; rules 1–3 clean (ran) • tests 476 ✓ / 66 skip, 1 ✗ = develop issue 14643; compile 0 spills (ran)

#	Sev	file:line	Finding
N1	🟡	PR body	Body (the squash message) still says fp16 "degree up to 128" • impact: develop's commit message contradicts the 97 cap • fix: say 97, and parity configs 55-67

Utility: 🟢 Positive — gfx1250 wgrad gains async K loop, cshuffle, two-stage; review fixes verified.

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 (every family, 4 flavors), goldens, PR tests, compile, fp16 degree sweep, async pruning and job-count probes, async IR diff, CI at head. Flags whether a closer look is needed; not an approval or rejection.

…mplicit

…ke/gfx1250-wgrad-wmma-implicit

jakpiase · 2026-11-03T09:39:39Z

Thanks for the follow-up.

N1: fixed in the PR body: the gfx1250 fp16 policy now says the automatic degree is capped at 97 (the packed fp16 atomic's accuracy bound; explicit split_k and overrides stay limited to 128, and the f32 workspace keeps the fp32 cap), and the Test Plan lists parity configs 55-67, with 67 covering the bf16 async zero-fill. The body also now describes the unsigned async range check, the async/sync pruning split and the per-target count_jobs gate, and lists byte identity at llvm24.

Develop is merged into the branch (a3467a8, no conflicts). After the merge: byte identity GREEN at llvm20/22/23/24, the goldens are OK, ctest -L host passes, and the CPU tests show no new failures.

…ke/gfx1250-wgrad-wmma-implicit

CongMa13

One actionable issue is reported inline.

…ke/gfx1250-wgrad-wmma-implicit # Conflicts: # dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm20.json # dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm22.json # dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm23.json # dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/arch/data/intrinsic_arch_domain.llvm24.json

…loop in C++ Python is_valid_wgrad_spec rejects an explicit vector_size_a/b when async_dma is set, because the async K loop derives its chunk widths from the channel counts alone, and range- and divisibility-checks the caps otherwise. The C++ validator had none of these checks, so a C spec with a cap on the async path passed validation and the builder ignored it. Port the Python loop into rocke_implicit_gemm_conv_wgrad_is_valid_spec at the same point in the gate order, with the same reason strings: reject a cap under async_dma, then require 1 <= cap <= 8 (4 for fp32 A) and that the cap divides the operand's stride-1 channel run (kpg for A, cpg for B). Tests: rocke_wgrad_streamk_host pins the Python reasons for the async rejection (A and B), the range and divisibility rejections, and the specs both engines accept; parity config 68 (config 55 with vector_size_a=4) is a both-engines-reject case. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-11-03T22:52:19Z

Thanks for the review.

CongMa13 inline (C++ async chunk widths ignore vector_size_a/b): fixed in 480216e: the C++ wgrad validator now rejects an explicit vector_size_a / vector_size_b under async_dma with Python's reason, at the same point in the gate order, and also ports Python's range and divisibility checks on those caps, so both engines accept the same specs. Covered by rocke_wgrad_streamk_host (reasons checked against the Python text) and parity config 68 (config 55 with vector_size_a=4, which both engines reject).

Develop is merged into the branch (8564ead). The only conflict was the four intrinsic_arch_domain.llvm*.json files: develop extended the global.atomic.fadd.v2f16 column, which this PR removes along with the intrinsic, so the column stays removed. After the merge and the fix: byte identity GREEN at llvm20/22/23/24, the goldens are OK, ctest -L host passes, and the CPU tests show no new failures.

…mplicit

cgmillette · 2026-11-04T10:31:38Z

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

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: fix verified, two minors open

Prior item	Status
CongMa13 (11-03) C++ async chunk widths ignore vector_size_a/b	🟢 FIXED `480216eb` • C++ validator = Python on 4,500 cap specs, reasons too (ran) • 5 of 5 mutants fail `rocke_wgrad_streamk_host`, which every rocKE CI shard runs (ran) • C++ sync loop still ignores caps: N2
Ours (11-01) N1 body: fp16 "degree up to 128"	🟡 OPEN • now right: 97 auto cap, configs 55–68, 912 jobs, 19 geometries (ran) • wrong: an explicit split_k is not held to 128; 200 stays 200 (ran) • left: say only the override is held to 128
Copilot (11-01) ×4: fp16 cap, unsigned check, count_jobs, zero-fill	🟢 FIXED `6e49bdaa` `cc12d0fb` `d02b3875` • cap probe FIXED; loader and count tests, byte identity pass (ran)
Copilot (11-01) async kernels pruned before benchmarking	🟢 Mostly `4d5e2006` • unchanged since 11-01; the PR's pruning test passes (ran)
therock-pr-bot (11-01) PR Description check	🟢 FIXED • bot ✓ at `8638d310` (read)
Reviewers	🟢 None yet • no approval; CongMa13 commented 11-03 (read)
Bots / CI	🟡 Running • running: Multi-Arch build stages, both hipDNN superbuilds; 19 ✓ • `8cc4d654` red was develop's fp64 doc test (14998); this run has its fix (ran) • merge: blocked, no approval yet
Our checks	🟢 Clean • merge vs `98e132f8`: clean, 0 behind; rocKE tree = checked merge (ran) • build + ctest 45 ✓ / 1 skip; byte identity GREEN ×4 flavors; goldens OK (ran) • tests 478 ✓, 1 ✗ = develop issue 14643; 0 spills; 2 changed test files in no CI lane (develop's) (ran)
Code quality	🟡 Minor • N2: C++ validates A/B caps its builder ignores (ran) • fp16 cap 97 has one home; the policy's 128 is unread (ran) • no open-PR overlaps (ran); cshuffle choice copied again (read)

#	Sev	file:line	Finding
N2	🟡	`platform/cpp/instances/common/conv_implicit_gemm_wgrad.cpp:3112`	C++ admits sync A/B caps its builder ignores • impact: caps 4/2 change Python's IR, not C++'s (ran); the binding drops both, so `rocke_engine` accepts all 6 new rejects (ran) • fix: apply the cap in `wgrad_free_axis_vec` (the builder gap predates this PR); read both fields in the binding; add a sync-cap parity config

Utility: 🟢 Positive — CongMa13's fix verified by mutation; body and C++ cap gaps found.

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 (every family, 4 flavors), goldens, PR tests (source and installed), compile, two-engine validator sweep and mutation test, split-K resolver and AOT grid probes, arch-domain tests, CI logs at both heads, code quality. Flags whether a closer look is needed; not an approval or rejection.

Python caps the synchronous free-axis load widths of the implicit-GEMM wgrad with vector_size_a / vector_size_b (_wgrad_free_axis_vec takes the smaller of the auto width and the cap). The C++ builder computed the auto width inline and ignored both fields, so a capped spec built different IR in the two engines, and the pybind binding dropped both fields, so rocke_engine neither built the capped kernel nor ran the cap checks the C++ validator already has. - wgrad_free_axis_vec takes the cap, as Python does, and the synchronous loader setup (M-outer and K-outer) uses it in place of its inline copy. The async availability check still passes no cap, as _wmma_async_vecs does; both validators reject a cap there. - conv_wgrad_build_spec reads vector_size_a and vector_size_b. Tests: parity configs 69-70 (config 55 on the synchronous K loop with caps 4/2, K-outer and M-outer); both fail against the previous C++ builder. A new gfx1250 K-outer harness case with the same caps. test_conv_wgrad_streamk checks through the binding that the capped kernels lower byte-identically and that the validator reasons agree on the async, range and divisibility rejects. Golden: new keys only. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

jakpiase · 2026-11-04T13:19:17Z

Thanks for the follow-up.

N2 (C++ admits sync A/B caps its builder ignores): fixed in 8ac37b6. wgrad_free_axis_vec now takes the explicit cap and returns the smaller of the auto width and the cap, as Python's _wgrad_free_axis_vec does. The synchronous loader setup (M-outer and K-outer) uses it in place of its inline copy of the auto width. The async availability check still passes no cap, as _wmma_async_vecs does; both validators reject a cap on the async loop. conv_wgrad_build_spec in the binding now reads vector_size_a and vector_size_b, so rocke_engine builds the capped kernel and runs the C++ cap checks. Coverage:

parity configs 69-70: config 55 on the synchronous K loop with caps 4/2, K-outer (69) and M-outer (70). Both fail against the previous C++ builder and are GREEN now;
a new gfx1250 K-outer harness case with the same caps;
test_conv_wgrad_streamk.py: through the binding, the capped kernels (gfx1250 K-outer and M-outer, gfx950 stream-K; caps 4/2, 4 only, 2 only) lower byte-identically to Python, and the validator gives Python's result for the async, range and divisibility rejects. The new tests fail with the previous binding.

The existing configs' hashes are unchanged; the golden diff only adds the new keys.

N1 (body): an explicit split_k is not held to 128, as you found. The resolver keeps it as given, bounded only by the gridDim.z budget. Only an override is held to 128 on the packed atomic (with a free reduction it moves to the workspace above 128), and the automatic fp16 degree is capped at 97. The body now says that. It also lists parity configs 55-70 and the binding test in the Test Plan.

The branch merges cleanly with develop (0 commits behind). At 8ac37b6: byte identity GREEN at llvm20/22/23/24, the representative-IR and differential goldens are OK, ctest -L host passes (45 passed, 1 skipped), and the CPU wgrad tests show no new failures against the previous head.

cgmillette

lgtm

cgmillette · 2026-11-04T14:11:29Z

No conflicts in the current merge pool. Targeted for next merge wave. Please prioritize needed approval.

cgmillette · 2026-11-04T14:42:02Z

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

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: fixes verified, one test nit

Prior item	Status
Ours (11-04) N2 C++ admits sync A/B caps its builder ignores	🟢 FIXED `8ac37b6b` • old builder turns parity 69/70 red; so does dropping either cap (ran) • 2,345 capped specs agree in both engines; old builder ignored 1,200 (ran) • new binding tests fail on the old builder and binding; CI skips them (no engine) (ran)
Ours (11-01) N1 body: fp16 "degree up to 128"	🟢 FIXED (body, 11-04) • resolver at head: explicit 200 kept, override held to 128, auto 97 (ran) • body now says exactly that; Test Plan lists 55–70 (read)
CongMa13 (11-03) C++ async chunk widths ignore vector_size_a/b	🟢 FIXED `480216eb` • async cap rejects match Python, via the binding too (ran)
Copilot (11-01) ×4: fp16 cap, unsigned check, count_jobs, zero-fill	🟢 FIXED `6e49bdaa` `cc12d0fb` `d02b3875` • files unchanged since; cap probe FIXED, 97 ≤ 97 (ran)
Copilot (11-01) async kernels pruned before benchmarking	🟢 Mostly `4d5e2006` • unchanged since 11-01 (read)
therock-pr-bot (11-01) PR Description check	🟢 FIXED • bot comment: all checks passed, 13:24Z (read)
Approvals	🟢 APPROVED `8ac37b6b` (cgmillette, 11-04) • ready for code-owner approval; N3 is a minor nit • still required: rocke-lib review of `library/` (read)
Bots / CI	🟡 Running • running: rocKE shards 1–2/2 on gfx94X; Windows math-libs gfx110X; 31 ✓ • therock-pr-bot cancelled at the push; its twin run passed (read) • merge: blocked: code-owner review, Multi-Arch CI pending (read)
Our checks	🟢 Clean • merge vs `9e2661d9`: clean, 2 behind; same rocKE tree as checked (ran) • ctest 45 ✓ / 1 skip; byte identity GREEN ×4, every family; goldens OK, keys only added (ran) • tests 511 ✓, 1 ✗ = develop issue 14643, at llvm22 and llvm24; installed 68 ✓ / 76 skip (ran) • `test_conv_abi.py`, `test_conv_split_k.py`: no CI lane runs them (develop's) (ran)
Code quality	🟡 Minor • N2 fixed: one A/B width helper per engine (ran) • N3: the min is untested (ran) • new comment cites alias `_free_axis_vec` (`.py:2601`); API standard rule 5: delete it (read)

#	Sev	file:line	Finding
N3	🟡	`platform/cpp/instances/common/conv_implicit_gemm_wgrad.cpp:612`	No test pins "the smaller of auto width and cap" • impact: a cap-as-width mutant passes all 33 tests and 71 configs; cap 6 at kpg 12 then ≠ Python (ran) • fix: subtest `self._wmma(problem=replace(self._P, K=12), vector_size_a=6)`: C++ IR == Python

Utility: 🟢 Positive — C++ cap fix verified by mutation and sweep; one test gap found.

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 (every family, 4 flavors), goldens and their diffs, PR tests (source, installed, unforced llvm24), compile, builder and binding mutation tests, two-engine cap sweep, split-K resolver probe, CI at head, code quality. Flags whether a closer look is needed; not an approval or rejection.

A cap wider than the auto free-axis width (cap 6 at kpg 12, auto 4) must leave the width at 4. The new binding test checks that both engines lower that spec to the uncapped kernel's IR, on the K-outer and M-outer loops. Also drop the _free_axis_vec alias of _wgrad_free_axis_vec in the builder and name _wgrad_free_axis_vec in the C++ comment that cited the alias. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…ke/gfx1250-wgrad-wmma-implicit

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

Thanks for the follow-up.

N3 (no test pins "the smaller of auto width and cap"): fixed in b02e4e2. TestLoadWidthCapsDualEngine.test_cap_wider_than_auto_width_keeps_auto_width builds config 55's spec at K=12 (kpg 12, which autos to 4) with vector_size_a=6 on the K-outer and M-outer synchronous loops. The spec is valid. Python lowers it to the same IR as the uncapped spec, and the C++ engine (through the binding) lowers it to that same IR. The cap-as-width mutant of wgrad_free_axis_vec fails both subtests.

Alias: the _free_axis_vec = _wgrad_free_axis_vec alias in the builder is gone. Its four call sites call _wgrad_free_axis_vec directly, and the C++ comment names _wgrad_free_axis_vec. The emitted IR is unchanged.

Develop is merged into the branch (838c729, no conflicts; it now includes #14835). At the new head:

byte identity is GREEN at llvm20/22/23/24;
the representative-IR goldens are OK;
ctest -L host passes (47 tests, 1 environmental skip);
the wgrad CPU tests pass (476 passed). The one failure 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.

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

…mplicit

cgmillette · 2026-11-05T02:36:01Z

Automated overlap notice (rocKE merge monitor)

Hi @jakpiase, this PR and #14907 each replace develop's packed fp16 global atomic, the llvm.amdgcn.global.atomic.fadd.v2f16 call, with an agent-scope atomicrmw fadd <2 x half> in both engines, in two different forms. Row N7 of our 5 November follow-up review of #14907 flagged it, and you have approved #14907; this notice records the resolution and adds rebase steps checked on the merge.

The two copies, here at b4a7d2a8f2 and in #14907 at 40f4ec1139 (paths under dnn-providers/hip-kernel-provider/rocke/):

Here: the bodies of _op_memref_global_atomic_add_pk_f16 (platform/python/rocke/core/lower_llvm.py:3600) and op_memref_global_atomic_add_pk_f16 (platform/cpp/core/lower_llvm/mem.cpp:394), each with syncscope("agent") monotonic written out (:3621, :415).
fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907: one packed-add helper per engine, with the element type and the scope as parameters: _lower_pk_atomic_add (lower_llvm.py:3542) and lower_pk_atomic_add (mem.cpp:381), which the bf16 and fp16 ops call (:3589-3593, :415-423), the fp16 one with "agent". It also takes the index operand's own width and the op's ordering.
The op's docstring in platform/python/rocke/core/ir.py (here :2437, fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907 :2436).

On every wgrad parity config the two forms emit the same IR: the differential hashes below carry over unchanged.

Owner: #14907: the fix lands before the feature (cgmillette's decision), and its helper is already the one home for the packed bf16 and fp16 adds. Resolution: this PR drops its fp16 body and calls #14907's helper. Your removal of the old intrinsic's declarations (data.cpp, the four intrinsic_arch_domain files) is not in #14907 and stays: after the merge neither engine calls it.

Suggested order: #14907 first, then this PR, as recorded. Today's merge plan (develop c6542db5aa) still lands this PR in wave 2 and #14907 in wave 3, after it; if this PR does land first, the last section applies. In the recorded order #14907 does nothing for this overlap, and this PR makes the changes below.

Rebase, after #14907 lands (checked with git merge-tree of this PR's b4a7d2a8f2 onto develop c6542db5aa merged with #14907's 40f4ec1139: 8 files conflict, all from #14907, since this PR merges cleanly with develop alone):

lower_llvm.py, mem.cpp and ir.py, one conflict each: take fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907's side (the helper call, its docstring).
library/kernels/common/conv_implicit_gemm_wgrad.py and platform/cpp/instances/common/conv_implicit_gemm_wgrad.cpp, one conflict each: keep both, fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907's register-prefetch K loop first, then your elif _wmma_async: branch and your else if(ctx.async_global_lds) call.
library/benchmarks/common/benchmark_implicit_gemm_conv.py, 2 conflicts (the cached two-stage rule and its help text): take fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907's side. From the cache it times every two-stage kernel unless --two-stage never, which already includes the wave32 bf16 ones your rule adds.
Renumber your wgrad parity configs 55-70 to 57-72, since fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907 takes 55-56. Both emitters pick fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907's 55-56 first (_PFA_BASES in _spec, make_cfg in the C emitter), so without it your 55 and 56, the gfx1250 async K loop at C=24 and at 64x64x64, silently leave the gates while both engines still agree (the C emitter merges without a conflict). In library/tests/parity/conv_implicit_gemm_wgrad_emit.py and .c, shift every id test, dict key, case label and comment of your configs; in the .py header take your 52-54 text, fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907's 55-56 paragraph, then your list renumbered, and make the END comment 0-72. Also the two comments that cite "the parity emitter's case 55" (library/tests/test_conv_wgrad_streamk.py:724, platform/tests/instances/wgrad_streamk_host.cpp:477).
platform/tests/instances/differential/golden/llvm_gfx_all.json, 5 conflicts (one per section): take fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907's 55 and 56, and move your conv_implicit_gemm_wgrad entries 55-70 up by two, hashes unchanged; your 53 stays. Re-recording the family on the merge gives exactly that: from platform/tests/instances/differential, python run_diff.py --mode ll --python-only --record-golden --only conv_implicit_gemm_wgrad --flavor llvm20,llvm22,llvm23,llvm24, then python run_diff.py --mode ir --canonical --python-only --record-golden --only conv_implicit_gemm_wgrad.

Resolved this way (on develop fe12dc0b98; the two develop commits since touch neither PR's files): the engine and rocke_wgrad_streamk_host build, and the host test passes; both engines agree on all 73 wgrad configs at all four flavors; the differential golden check passes at all four flavors and canonical; the representative conv_wgrad and conv_wgrad_two_stage shards re-record unchanged; and both PRs' changed tests with the arch-domain tests pass (632 passed, 70 skipped), but for test_conv_abi.py::test_cache_applies_runtime_launch_limits, which fails the same way on develop.

If this PR lands first (today's plan order), take #14907's text now in place of your fp16 body: its _lower_pk_atomic_add with the two ops (lower_llvm.py:3542-3593), its lower_pk_atomic_add with the comment above it and the two ops (mem.cpp:363-423), and its ir.py docstring (:2436-2437); keep your configs at 55-70. Checked: #14907's rebase then has no conflict in lower_llvm.py, mem.cpp or ir.py and keeps one helper per engine (5 other files still conflict). Landing second, #14907 then renumbers its two configs to 71-72 in both emitters and the golden.

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

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

cgmillette · 2026-11-05T10:11:28Z

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

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: N3 fixed; rebase on 14842

Prior item	Status
Ours (11-04) N3: no test pins the min of auto width and cap	🟢 FIXED `b02e4e28` • new test fails a no-min mutant in either engine (ran) • CI skips it whole: `_engine()` comes first (`test_conv_wgrad_streamk.py:780`) (ran) • fix: call `_engine()` only in the last assertion; Python half then runs (ran)
Ours (11-05) overlap notice: fp16 packed add, owner 14907	🟡 OPEN, amended • decision stands; 14907 now `05a74c73`, merged with develop (read) • this PR keeps its fp16 copy until 14907 lands (read) • amend "take 14907's text now": from `05a74c73`; `40f4ec11` lacks 14842's guard (ran)
Ours N1–N2, CongMa13, Copilot ×5, PR Bot (11-01 to 11-04)	🟢 FIXED (as at follow-up 3); Copilot pruning Mostly `4d5e2006` • only N3's 3 files changed since `8ac37b6b` (read) • fp16 cap probe FIXED, 97 ≤ 97; parity 69–70 GREEN on the merge (ran)
cgmillette (11-04)	🟢 APPROVED `8ac37b6b` • older head; still counts on develop • not covered: `b02e4e28` and the 14842 rebase • rocke-lib code-owner review still requested (read)
Bots / CI	🟢 Green • 45 ✓ · 51 skipped; required checks ✓ (read) • attempt 1 ✗: Windows math-libs, TheRock artifact tool crash (hipthreads); attempt 2 ✓ (read) • CI predates 14929, 14834, 14842, 15097; merge: dirty (2 files)
Our checks	🟡 Conflict • merge vs `8b6ada32`: 2 conflicts, both from 14842 (`f28ed8f0`): N4 (ran) • resolved per N4: ctest 46 ✓ / 1 skip; byte identity GREEN ×4, all families; goldens clean (ran) • tests (+14842's, 14834's), both engines: 1394 ✓, 1 ✗ (develop issue 14643), llvm22 and llvm24 (ran) • installed 68 ✓ / 77 skip; compile: 0 failures, 0 newly rejected, 0 spills (ran)
Code quality	🟡 Minor • 9 new rejects hand-write develop's `ROCKE_SPEC_FAIL` (`.cpp:657`, 8 more) (read) • 8 new `Optional`/`Tuple`/`List` hints; PYTHON_STYLE §3 (`.py:1368`) (read) • fp16 packed add: own copy until 14907 lands (overlap row)

#	Sev	file:line	Finding
N4	🟡	`platform/python/rocke/core/lower_llvm.py:3612`, `platform/cpp/core/lower_llvm/mem.cpp:406`	Both conflicts are 14842's guard inside the fp16 packed add • impact: the PR's side drops 14842's constant-space refusal; CI misses the C++ loss (ran) • fix: keep 14842's `require_global_header_ptr` line; drop its `need("global.atomic.fadd.v2f16")` (ran)

Utility: 🟢 Positive — N3 pinned in both engines; 14842 conflict resolved and checked.

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 (resolved), build + ctest, byte identity (every family, 4 flavors), goldens, differential golden, tests on the merge (both engines at llvm22 and llvm24, and without the engine), installed tests, compile, N3 mutants in both engines, conflict-resolution probe, CI at head, code quality. Flags whether a closer look is needed; not an approval or rejection.

…ke/gfx1250-wgrad-wmma-implicit # Conflicts: # dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp # dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py

aledudek

LGTM

From what is see the only problem left are merge conflicts, but they arise every second with our blazingly fast development pace

cgmillette · 2026-11-05T17:11:30Z

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

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: N4 fixed; rebase after 14907

Prior item	Status
Ours (11-05) N4: 14842's guard inside the fp16 packed add	🟢 FIXED `1f9d175d` • both engines keep 14842's guard, drop only `need` (`lower_llvm.py:3653`, `mem.cpp:429`) (ran) • one home: develop's helper in each engine; probe passes, controls OK (ran) • fp16 split-K still emits the packed `atomicrmw` (configs 20, 51) (ran)
Ours (11-05) overlap notice: fp16 packed add, owner 14907	🟡 OPEN • 14907 lands first, auto-merge on; jakpiase rebases then (read) • copy here ignores `ordering`, which the bf16 sibling reads (read) • notice's steps hold, same 8 files; 14907's text keeps 14842's guard and the IR (ran)
Ours N1–N3, CongMa13, Copilot ×5, PR Bot (11-01 to 11-04)	🟢 FIXED (as at follow-up 4); Copilot pruning mostly fixed `4d5e2006` • only `mem.cpp` changed since `b4a7d2a8` (ran) • N3's tests still skip whole in CI: `_engine()` first (`test_conv_wgrad_streamk.py:780`) (read)
cgmillette (11-04)	🟢 APPROVED `8ac37b6b` • older head; still counts on develop • approved behaviour unchanged by the merge; new since: `b02e4e28` (test, alias rename), 14842's guard (ran) • review decision now APPROVED (read)
aledudek (11-05)	🟢 APPROVED `1f9d175d` • at head • "only merge conflicts left": none with develop; 14907 forces one rebase (ran)
Bots / CI	🟡 Running • 39 ✓ · 49 skipped · 1 queued: Windows gfx110X sanity, since 13:28Z (read) • gfx94X rocKE shards 87/87 ✓ each; no Windows rocKE shard yet (read) • merge: clean; blocked, Multi-Arch summary not posted; CI predates 15024, 14843 (read)
Our checks	🟢 Clean • merge vs `240c8b0d` (has 15024, 14843): clean, 0 conflicts (ran) • ctest 48 ✓ / 1 skip; byte identity GREEN ×4, all 24 families (ran) • goldens and differential golden clean; compile: 0 failures, 0 spills (ran) • tests 514 ✓, 1 ✗ (develop issue 14643), llvm22 and llvm24; installed 68 ✓ / 77 skip (ran)
Code quality	🟡 Minor • N4 fixed: 14842's guard kept, one call per engine (read) • fp16 packed add: own copy until 14907 lands (ran) • 9 hand-written rejects, 8 legacy hints unchanged (read)

No net-new findings. The quality lane's points are statuses of the overlap notice, N3 and follow-up 4's code quality row.

Utility: 🟢 Positive — 14842's guard kept in both engines; approved behaviour unchanged since 8ac37b6b.

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 (every family, 4 flavors), goldens, differential golden, tests on the merge (llvm22 and llvm24), installed tests, compile, N4 probe, fp16 split-K IR, rebase trial with 14907, diff against the approved head, CI at head, code quality. Flags whether a closer look is needed; not an approval or rejection.

jakpiase and others added 6 commits November 1, 2026 18:17

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 4 commits November 1, 2026 20:23

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

jakpiase and others added 2 commits November 2, 2026 23:34

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-wgrad-wmma-i…

a69123a

…mplicit

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

a3467a8

…ke/gfx1250-wgrad-wmma-implicit

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

45494bf

…ke/gfx1250-wgrad-wmma-implicit

CongMa13 reviewed Nov 3, 2026

View reviewed changes

Comment thread ...oviders/hip-kernel-provider/rocke/platform/cpp/instances/common/conv_implicit_gemm_wgrad.cpp Resolved

jakpiase and others added 2 commits November 3, 2026 22:18

jakpiase requested a review from CongMa13 November 3, 2026 22:53

jakpiase added 2 commits November 4, 2026 02:15

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-wgrad-wmma-i…

8cc4d65

…mplicit

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-wgrad-wmma-i…

8638d31

…mplicit

cgmillette approved these changes Nov 4, 2026

View reviewed changes

jakpiase and others added 2 commits November 4, 2026 19:47

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

838c729

…ke/gfx1250-wgrad-wmma-implicit

This was referenced Nov 4, 2026

feat(rocke): gfx1250 direct wgrad kernels with two-stage reduction #14833

Open

fix(rocke/wgrad): stop fencing the packed fp16 split-K atomic #14907

Merged

Merge branch 'develop' into users/jakpiase/rocke/gfx1250-wgrad-wmma-i…

b4a7d2a

…mplicit

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

1f9d175

…ke/gfx1250-wgrad-wmma-implicit # Conflicts: # dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/mem.cpp # dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py

aledudek approved these changes Nov 5, 2026

View reviewed changes

cgmillette enabled auto-merge (squash) November 5, 2026 17:49

cgmillette merged commit 584c490 into develop Nov 6, 2026
237 of 242 checks passed

cgmillette deleted the users/jakpiase/rocke/gfx1250-wgrad-wmma-implicit branch November 6, 2026 01:22

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

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!

Uh oh!

cgmillette commented Nov 1, 2026

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase) — 🟡 First-round: minor fixes, CI running

Uh oh!

jakpiase commented Nov 1, 2026

Uh oh!

cgmillette commented Nov 1, 2026

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: fixes verified, CI running

Uh oh!

jakpiase commented Nov 3, 2026

Uh oh!

CongMa13 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

jakpiase commented Nov 3, 2026

Uh oh!

cgmillette commented Nov 4, 2026

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: fix verified, two minors open

Uh oh!

jakpiase commented Nov 4, 2026

Uh oh!

cgmillette left a comment

Choose a reason for hiding this comment

Uh oh!

cgmillette commented Nov 4, 2026

Uh oh!

cgmillette commented Nov 4, 2026

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: fixes verified, one test nit

Uh oh!

jakpiase commented Nov 4, 2026

Uh oh!

cgmillette commented Nov 5, 2026

Uh oh!

cgmillette commented Nov 5, 2026

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: N3 fixed; rebase on 14842

Uh oh!

aledudek left a comment

Choose a reason for hiding this comment

Uh oh!

cgmillette commented Nov 5, 2026

PR #14832 — async K loop, cshuffle, fp16 split-K and two-stage split-K for gfx1250 WMMA wgrad (jakpiase, AICK-2880) — 🟡 Follow-up: N4 fixed; rebase after 14907

Uh oh!

Uh oh!

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