Repository navigation
feat(rocke): shape-generic (AOT) Winograd forward transforms for gfx1250 - #14791
Conversation
The Winograd F(m,3) transform kernels (K1 filter, K2a prescale max, K2
input, K4 output) bake the problem extents in, so every convolution shape
needs its own binaries and the pipeline cannot be served from an AOT cache.
conv_winograd_fwd_rt builds the same stages, workspace layouts and
arithmetic with every extent as a kernel argument (divisions by runtime
extents through a host-computed magic number), so one binary serves every
problem its knobs fit. The K3 batched GEMMs already take their extents at
run time.
- kernels/common/conv_winograd_fwd_rt.py: per-stage specs (RtFilterSpec,
RtPrescaleSpec, RtInputSpec, RtOutputSpec), stage_launch (grid, block and
scalar kernargs), the AOT grid (aot_stage_specs) and the default
per-problem pick (default_stage_specs). K2 keeps the register path (with
a fused prescale over a runtime lane count) and the LDS region.
- C++ engine port with byte-identical IR; the shared Winograd math helpers
move into rocke_wino (no IR change to the existing families).
- conv_winograd_fwd_rt_emit.{py,c} parity pair and a representative IR
golden shard.
- Dispatch: plan kind unfused2d_rt and the opt-in candidate
conv_winograd_fwd_aot (algorithm="winograd_aot") over the selector's 3x3
plans; algorithm="auto" never picks it.
- benchmarks/common/benchmark_winograd_aot.py: --compile-all builds the
grid into a cache, --run-from-cache assembles each case's pipeline from
it, picking each stage's binary by timing and verifying the pipeline.
- Tests: CPU (kernels, dispatch) and on-GPU numerics; docs and support
matrix.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
There was a problem hiding this comment.
🟡 Changes recommended
Validation gaps allow unsupported geometry and non-power-of-two LDS lane counts, leading to incompatible kernel launches and missing or duplicate stores.
6 open issues
Reject nonpositive and non-power-of-two LDS lane counts · New Validate geometry in stage_launch before returning launch arguments · New Validate chunks as a positive integer before planning · New Reject filter_rows values below 1 in the Python validator · New Reject block sizes above the 1024-thread limit · New Include new test files in the installed CTest lists · New
What changed in this PR
Adds shape-generic Winograd forward transforms for gfx1250, allowing compiled binaries to be reused across convolution shapes through an opt-in AOT dispatch route while preserving the existing route.
Changes:
- Introduces runtime-extent transform specs, launch arguments and AOT binary configurations.
- Adds
winograd_aotdispatch and a cache compilation, tuning and verification benchmark. - Shares C++ transform helpers and adds parity, golden, CPU and GPU coverage with supporting documentation.
Paths are relative to dnn-providers/hip-kernel-provider/rocke/.
| File | Description |
|---|---|
platform/tests/instances/rocke_ir_parity_harness.py |
Registers representative runtime-transform IR cases. |
platform/tests/golden/rocke_representative_ir_sha256/conv_winograd_fwd_rt.json |
Adds LLVM-flavor golden hashes. |
platform/python/rocke/instances/SUPPORT_MATRIX.md |
Documents AOT Winograd availability. |
platform/dsl_docs/instances/index.md |
Indexes the new transform specs. |
platform/dsl_docs/instances/convolution.md |
Explains runtime transforms and dispatch. |
platform/cpp/instances/common/conv_winograd_fwd.cpp |
Delegates transform math to shared helpers. |
platform/cpp/instances/common/conv_winograd_common.cpp |
Houses reusable transform and prescale helpers. |
platform/cpp/include/rocke/instance_conv_winograd_internal.h |
Declares shared C++ helpers. |
platform/cpp/include/rocke/instance_conv_winograd_fwd_rt.h |
Defines the runtime-transform C API. |
library/tests/test_conv_winograd_fwd_rt_gpu.py |
Adds end-to-end numeric and binary-reuse tests. |
library/tests/test_conv_winograd_fwd_rt_cpu.py |
Tests division, specs, AOT configurations and launches. |
library/tests/test_builder_signature_contract.py |
Registers defaulted runtime-transform specs. |
library/tests/parity/conv_winograd_fwd_rt_emit.py |
Adds Python parity configurations and emission. |
library/tests/parity/conv_winograd_fwd_rt_emit.c |
Adds matching C-side parity emission. |
library/tests/dispatch/test_conv_winograd_aot_dispatch.py |
Tests routing, launch arguments and batch splitting. |
library/kernels/common/conv_winograd_select.py |
Labels runtime-transform plans. |
library/kernels/common/conv_winograd_fwd_rt.py |
Implements shape-generic transforms and launch geometry. |
library/dispatch/grouped_convolution.py |
Registers the opt-in AOT candidate. |
library/dispatch/conv_winograd_plan.py |
Assembles runtime-transform launch pipelines. |
library/benchmarks/common/benchmark_winograd_aot.py |
Compiles caches and benchmarks assembled pipelines. |
Review effort: Balanced
Review follow-ups for the shape-generic Winograd forward transforms: - Both engines refuse an LDS lane count that is not a positive power of two (the lane mapping is a shift and mask), an LDS region of zero width, a filter_rows below 1 and a K2 / K4 block above 1024 threads. - stage_launch runs RtGeometry.check itself, so a non-3x3, strided or oversized problem is refused there and not only by the callers. - The Winograd launch plans refuse a batch split that is not a positive integer instead of normalizing it. - Six new rejected parity configs; the new CPU, dispatch and GPU tests are in the installed CTest lists. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #14791 — shape-generic (AOT) Winograd forward transforms for gfx1250 (jakpiase, AICK-2884) — 🟡 First-round: wide-C fused K2 spills, untested
Utility: 🟢 Positive — AOT-cacheable 3x3 Winograd transforms behind an opt-in route; wide-C untested on GPU. 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, build + ctest, byte identity, goldens, PR tests, gfx1250 compile and AOT-grid spill sweep, CI at head. Flags whether a closer look is needed; not an approval or rejection. |
The default shape-generic K2 with the fused prescale ran at block 1024 for anything over 256 channel lanes. A 1024-thread workgroup caps a wave at 128 VGPRs, so the wider binaries spilled, and at 512 lanes the block-1024 K2 was slower than the block-512 one the JIT pipeline uses (block = max(block, lanes)). - The AOT grid adds fused K2 binaries at block 512. - The default is the smallest block of 256 / 512 / 1024 that holds all lanes; tiles over 36 bins (F(6x4), F(4x6)) take the split prescale past 512 lanes, where their block-1024 K2 spills most and the split is faster. - CPU test of the block choice; GPU cases for bf16 C = 256, 2048 and 4096 (the fused max over more than one wave, blocks 512 / 1024, and the split). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
|
Re N1 (wide-C fused K2 spills): fixed in 0f78ac0.
|
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #14791 — shape-generic (AOT) Winograd forward transforms for gfx1250 (jakpiase, AICK-2884) — 🔴 Follow-up: llvm24 alias breaks goldens
Utility: 🟢 Positive — AOT-cacheable 3x3 Winograd transforms behind an opt-in route; wide-C spills trimmed. 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: range-diff since our review, merge with today's develop, build + ctest, byte identity ×4 flavors, golden gate and a scratch re-record, differential golden, PR tests, AOT-grid spill sweep with default picks, CI logs at head. Flags whether a closer look is needed; not an approval or rejection. |
…ke/gfx1250-winograd-aot
develop now records LLVM 24 as its own golden flavor instead of an alias of LLVM 23. Re-record the conv_winograd_fwd_rt shard after merging develop: drop the llvm24 -> llvm23 alias and add the six llvm24 cases. The llvm20/22/23 hashes are unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
|
Re N1 (llvm24 alias in the new golden shard): fixed in 7fe0c22. I merged develop (ff9284d) and re-recorded the |
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #14791 — shape-generic (AOT) Winograd forward transforms for gfx1250 (jakpiase, AICK-2884) — 🟡 Follow-up: alias fixed, CI running
Net-new: none. Utility: 🟢 Positive — AOT-cacheable 3x3 Winograd transforms behind an opt-in route; LLVM 24 golden added. 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: delta since our review (per-file patches, merge resolution), merge with today's develop, build + ctest, byte identity ×4 flavors, golden gate and a scratch re-record, differential golden, PR tests, compile, CI at head. Flags whether a closer look is needed; not an approval or rejection. |
benchmark_conv_compare.py now runs benchmark_winograd_aot.py as a third engine beside the direct-conv and implicit-GEMM sweeps. It runs by default on a --run-from-cache comparison whose directory holds the Winograd cache, and --skip-winograd leaves it out. A JIT comparison, a non-fwd --direction or a missing cache skip it with an info line. - case_results: "winograd" is an engine, with its own merged and CSV columns. The winner is the fastest engine that measured the case, and the speedup is over the next best. - benchmark_winograd_aot.py reads the compare inputs: --miopen-cmd, the shape flags, and --dtype / --direction overrides. It skips the cases the winograd_aot route refuses before it loads the cache. It times the assembled pipeline with --timing staged|block. --results writes case_results rows, --skip-measured resumes (it replaces --resume), and it prints the standard Best: TFLOPS line. A case with no fitting pipeline exits 0; one whose pipelines all fail verification exits 1. - run_conv_shards.py --script winograd shards the sweep. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
|
@cgmillette thanks for the follow-up. Agreed: the llvm24 alias fix (7fe0c22) is in, and there were no net-new findings to address. @bartekxk your
Testing: new host-side tests in |
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #14791 — shape-generic (AOT) Winograd forward transforms for gfx1250 (jakpiase, AICK-2884) — 🟡 Follow-up: compare integration verified, CI running
Net-new: none. Utility: 🟢 Positive — opt-in AOT Winograd for gfx1250, now benchmarked beside direct and implicit-GEMM. 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: delta since our review, merge with today's develop, build + ctest, byte identity ×4 flavors, golden gate and the saved alias check, differential golden, PR tests (source tree and installed layout), a host-side run of the compare, sweep and shard paths with the GPU stubbed, compile, CI at head. Flags whether a closer look is needed; not an approval or rejection. |
TheRock Submodule Bump ActivityNewest first
|



Motivation
The Winograd F(m,3) forward pipeline on gfx1250 builds its transform kernels (K1 filter, K2a prescale max, K2 input, K4 output) per problem: the extents are baked in, so every convolution shape needs its own binaries and the pipeline cannot be served from an AOT kernel cache next to the direct and implicit-GEMM instances. The K3 batched GEMMs already take their extents at run time.
ISSUE ID : AICK-2884
Technical Details
Shape-generic transforms:
kernels/common/conv_winograd_fwd_rt.pyconv_winograd_fwd. Every problem extent is a kernel argument, so one binary serves every problem its knobs fit.magic_div), exact for 31-bit indices.RtFilterSpec,RtPrescaleSpec,RtInputSpec,RtOutputSpec. The compile-time knobs are the tile, the dtypes, the channel vector width, the workgroup size, the K2 register run or LDS region, and the prescale. Every field is defaulted.stage_launchreturns a stage's grid, block and scalar kernargs.aot_stage_specs()is the AOT grid.default_stage_specs(problem, tile, dtype)is the default per-problem pick: the JIT pipeline's tuned knobs, narrowed to the problem's channel extents.Both engines
cpp/instances/common/conv_winograd_fwd_rt.cpp(instance_conv_winograd_fwd_rt.h), byte-identical IR.bt_1d,at_1d,filter_matrix_int, the prescale bit helpers) move fromconv_winograd_fwd.cppintorocke_wino. There is no IR change to the existing families.conv_winograd_fwd_rt_emit.{py,c}, 50 configs:conv_winograd_fwd_rt.Dispatch
unfused2d_rtindispatch/conv_winograd_plan.py. Each transform step is named by its binary's kernel name (its AOT identity) and carries the problem in its kernargs.conv_winograd_fwd_aot(algorithm="winograd_aot"). It runs the selector's 3x3 plans on the shape-generic binaries.algorithm="auto"never picks it, and the existingconv_winograd_fwdroute is unchanged.Benchmark
benchmarks/common/benchmark_winograd_aot.py:--compile-allbuilds the grid into a cache.--run-from-cacheassembles each case's pipeline from it. For every tile it picks each stage's binary by timing, verifies the assembled pipeline against the reference convolution, and times the chain (--timing staged|block).--miopen-file/--miopen-cmd/ shape flags,--dtype/--direction,--results(case_resultsrows, enginewinograd),--skip-measured, and the standardBest:line. Cases thewinograd_aotroute refuses are skipped with the reason.benchmark_conv_compare.pyruns it as a third engine after direct-conv and implicit-GEMM. It runs by default on a--run-from-cachecomparison whose directory also holds the Winograd cache, and--skip-winogradturns it off.--resultsmerges the three engines on one row (winograd_*columns; the winner and speedup are taken over every engine that measured the case).run_conv_shards.py --script winogradshards the Winograd sweep.Test Plan
tools/check_byte_identity.py --only winograd) at llvm20, llvm22, llvm23 and llvm24.tests/test_rocke_ci_static.py(IR parity coverage and golden).library/tests/test_conv_winograd_fwd_rt_cpu.py:magic_divexactness;stage_launchkernargs.library/tests/dispatch/test_conv_winograd_aot_dispatch.py: route selection, launch-plan steps and kernargs, batch split, refusals.library/tests/test_run_conv_shards.py: the three-engine merge, CSV and summary; compare forwarding the Winograd args and skipping it (no cache, JIT, non-fwd,--skip-winograd); its exit status; the AOT script's case reading and refusals.benchmark_conv_compare.py --run-from-cache --verify --resultsrun; a mixed case file (-F 0, 1x1, stride 2); block timing;--skip-measuredresume;run_conv_shards.py --script winograd.test_conv_winograd_dispatch.py,test_grouped_conv_dispatch.py,test_library_layering.py,test_builder_signature_contract.py.library/tests/test_conv_winograd_fwd_rt_gpu.pycompares against the conv reference (max|out - ref| / max(max|ref|, 1) < 5e-2) with poisoned workspaces. It covers:pre-commit runon the changed files.Test Result
conv_winograd_fwd_rt: 34 configs identical, 16 rejected by both engines.test_run_conv_shards.py88 passed; together with the Winograd CPU, dispatch, layering and signature tests, 182 passed.Submission Checklist