Add fp32 winograd for gfx12 - #5092
Open
pfultz2 wants to merge 22 commits into
Open
Conversation
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #5092 +/- ##
========================================
Coverage 93.26% 93.26%
========================================
Files 623 623
Lines 32969 32969
========================================
Hits 30747 30747
Misses 2222 2222 🚀 New features to boost your workflow:
|
pfultz2
requested review from
CharlieL7,
TedThemistokleous,
kahmed10 and
shivadbhavsar
July 23, 2026 18:32
Regressions detected 🔴 * No develop baseline was found for this PR's branch point; compared against the latest available develop run instead. |
|
Contributor
There was a problem hiding this comment.
Pull request overview
This PR extends the GPU Winograd F(2,3) convolution path on gfx12 to support fp32 by adding a new FMA/DPP-based kernel, wiring it through prefuse + JIT compilation, and adding a verification test to exercise the new selection/weight-transform logic.
Changes:
- Add a gfx12 fp32 Winograd F(2,3) kernel (
winograd_conv_fp32.hpp) plus JIT wiring to compile/tune it. - Extend
prefuse_opsWinograd matching/heuristics and add fp32-specific host-side weight transforms (full-U, S-store, v-innermost). - Factor gfx12 “OOB raw buffer load returns 0” helpers into a shared kernel header and update the existing fp16 WMMA kernel to use it.
- Add a new
test/verifycase intended to exercise the fp32 Winograd path.
Reviewed changes
Copilot reviewed 6 out of 6 changed files in this pull request and generated 6 comments.
Show a summary per file
| File | Description |
|---|---|
| test/verify/test_conv_3x3_winograd_fp32.cpp | New verify test to exercise fp32 Winograd selection and edge cases |
| src/targets/gpu/prefuse_ops.cpp | fp32 Winograd matcher enablement, fp32 heuristics, and fp32 weight-literal transforms (full-U / S-store / vinner) |
| src/targets/gpu/kernels/include/migraphx/kernels/winograd_conv.hpp | Refactor to use shared OOB buffer-load helpers |
| src/targets/gpu/kernels/include/migraphx/kernels/winograd_conv_fp32.hpp | New fp32 FMA/DPP Winograd kernel implementation |
| src/targets/gpu/kernels/include/migraphx/kernels/buffer_load.hpp | New shared gfx12 OOB-tolerant raw buffer-load utilities |
| src/targets/gpu/jit/winograd_conv.cpp | JIT: add fp32 kernel codegen path, tuning configs, and integer_divide_ceil usage |
Comment on lines
+46
to
+51
| mm->add_instruction( | ||
| migraphx::make_op("convolution", | ||
| {{"padding", {1, 1}}, {"stride", {1, 1}}, {"dilation", {1, 1}}}), | ||
| x, | ||
| w); | ||
| return p; |
Comment on lines
+168
to
+182
| const auto nw = v.get("nw", std::size_t{4}); | ||
| const auto ko = v.get("ko", std::size_t{8}); | ||
| const auto tiles = v.get("tiles", std::size_t{2}); | ||
| // sk = within-WG channel-split factor (must divide nw). sk>1 has nw/sk | ||
| // NT-groups whose sk waves split the channel contraction. | ||
| const auto sk = v.get("sk", std::size_t{1}); | ||
| // pipe = software-pipeline the input transform into the FMA loop (true) vs | ||
| // the simple transform-then-FMA path (false). pipe costs a second live | ||
| // v_reg, so the tuner only offers it on small (non-spilling) solutions. | ||
| const bool pipe = v.get("pipe", false); | ||
| // cu = channel-unroll (weight load width): 4=b128, 2=b64, 1=b32. Smaller | ||
| // cu shrinks the pipelined double-buffer (finer-grained pipeline) at the | ||
| // cost of more, narrower weight loads. | ||
| const auto cu = v.get("cu", std::size_t{4}); | ||
| // Weight layout (chosen by prefuse, read back from the literal's shape): |
Comment on lines
+84
to
+90
| static winograd_fp32_weight_layout winograd_fp32_weight_layout_of(const shape& w) | ||
| { | ||
| if(w.lens().at(0) == 3) | ||
| return winograd_fp32_weight_layout::sstore; | ||
| return w.lens().at(1) == 4 ? winograd_fp32_weight_layout::full_u | ||
| : winograd_fp32_weight_layout::full_u_vinner; | ||
| } |
Comment on lines
+150
to
+153
| static_assert(KO >= 1, "KO must be >= 1"); | ||
| static_assert(TILES >= 1, "TILES must be >= 1"); | ||
| static_assert(SK >= 1 and SK <= NW and (NW % SK) == 0, "SK must divide NW"); | ||
| static_assert(CU == 1 or CU == 2 or CU == 4, "CU must be 1, 2, or 4"); |
Comment on lines
+324
to
+346
| // BOTTLENECK (measured by address-isolation on 256->256@64, replacing a load's | ||
| // offset with a constant so it hits one cached line): full 0.698ms; with the | ||
| // INPUT load coalesced 0.229; with the WEIGHT load coalesced 0.239; with BOTH | ||
| // coalesced = pure compute 0.089ms. So the fused COMPUTE is ~3x FASTER than MLIR | ||
| // (0.089 vs 0.276) -- the whole gap is the two SCATTERED loads (both read | ||
| // lane==v_col: input W-columns are C-strided; weight U v-slices are K*C apart), | ||
| // which THRASH the cache super-linearly (0.698 >> 0.089+0.14+0.15). Traffic is | ||
| // ~23 GB/s (<<peak) so it is cache-miss latency/thrashing, not bandwidth. | ||
| // | ||
| // FIX SHIPPED for the WEIGHT scatter (the host-controllable one): U is laid out | ||
| // v-innermost [u,k,c,v] (see prefuse compute_winograd_weights_f23_fp32 vinner) | ||
| // so the 4 v_col lanes read consecutive floats -> the weight load coalesces, | ||
| // relieving the thrash. GATED to out_c>=128 && in_c<=out_c: its strided (b32) | ||
| // channel load adds issue overhead that regresses small/cached-weight shapes. | ||
| // Net +4.5% geomean vs the scattered path (0.878->0.918x MLIR), memory-bound | ||
| // configs -61%->-47..-56%. The INPUT scatter (17MB, uncacheable, no layout | ||
| // freedom) is the dominant residual and has no clean in-kernel fix -- real | ||
| // input coalescing (lane==channel rewrite, LDS spatial-blocking) was measured | ||
| // NET-NEUTRAL-to-LOSS, and a full fused implicit-GEMM (scratchpad/ | ||
| // winograd_conv_fp32_gemm.hpp) helps memory-bound (~0.69x) but is a big | ||
| // aggregate loss (0.48x, wrecks small shapes) -- the 16 winograd positions cap | ||
| // its arithmetic intensity. Beating MLIR outright would need a MULTI-kernel | ||
| // winograd (transform kernels + a library batched GEMM on materialized V/M). |
Comment on lines
+33
to
+46
| // gfx12 buffer-resource word 3 (from composable_kernel): makes raw buffer loads | ||
| // return 0 for out-of-range byte offsets, so bounds/halo checks collapse to an | ||
| // offset select against a sentinel instead of a per-load branch. | ||
| constexpr uint32_t oob_buffer_rsrc_word3 = 0x31004000; | ||
|
|
||
| // Build an out-of-bounds-tolerant buffer descriptor for a read-only pointer (the | ||
| // resource is only ever loaded from; the const_cast is required by the builtin's | ||
| // non-const pointer parameter). | ||
| template <class T> | ||
| __device__ inline __amdgpu_buffer_rsrc_t make_oob_buffer_rsrc(const T* p, uint32_t byte_count) | ||
| { | ||
| auto* base = const_cast<T*>(p); // NOLINT(cppcoreguidelines-pro-type-const-cast) | ||
| return __builtin_amdgcn_make_buffer_rsrc(base, 0, byte_count, oob_buffer_rsrc_word3); | ||
| } |
Collaborator
|
Failing CI |
kahmed10
reviewed
Jul 30, 2026
Comment on lines
+318
to
+346
| // Full block transform (all CU channels). For NHWC (channels innermost, | ||
| // stride 1) the CU channels are contiguous, so load them with one b128 per | ||
| // (tile, a-row) -- 4x fewer input loads than the per-channel b32 path -- then | ||
| // apply the same per-channel v/u transform. (NCHW channels are H*W apart, so | ||
| // it stays per-channel.) | ||
| // | ||
| // BOTTLENECK (measured by address-isolation on 256->256@64, replacing a load's | ||
| // offset with a constant so it hits one cached line): full 0.698ms; with the | ||
| // INPUT load coalesced 0.229; with the WEIGHT load coalesced 0.239; with BOTH | ||
| // coalesced = pure compute 0.089ms. So the fused COMPUTE is ~3x FASTER than MLIR | ||
| // (0.089 vs 0.276) -- the whole gap is the two SCATTERED loads (both read | ||
| // lane==v_col: input W-columns are C-strided; weight U v-slices are K*C apart), | ||
| // which THRASH the cache super-linearly (0.698 >> 0.089+0.14+0.15). Traffic is | ||
| // ~23 GB/s (<<peak) so it is cache-miss latency/thrashing, not bandwidth. | ||
| // | ||
| // FIX SHIPPED for the WEIGHT scatter (the host-controllable one): U is laid out | ||
| // v-innermost [u,k,c,v] (see prefuse compute_winograd_weights_f23_fp32 vinner) | ||
| // so the 4 v_col lanes read consecutive floats -> the weight load coalesces, | ||
| // relieving the thrash. GATED to out_c>=128 && in_c<=out_c: its strided (b32) | ||
| // channel load adds issue overhead that regresses small/cached-weight shapes. | ||
| // Net +4.5% geomean vs the scattered path (0.878->0.918x MLIR), memory-bound | ||
| // configs -61%->-47..-56%. The INPUT scatter (17MB, uncacheable, no layout | ||
| // freedom) is the dominant residual and has no clean in-kernel fix -- real | ||
| // input coalescing (lane==channel rewrite, LDS spatial-blocking) was measured | ||
| // NET-NEUTRAL-to-LOSS, and a full fused implicit-GEMM (scratchpad/ | ||
| // winograd_conv_fp32_gemm.hpp) helps memory-bound (~0.69x) but is a big | ||
| // aggregate loss (0.48x, wrecks small shapes) -- the 16 winograd positions cap | ||
| // its arithmetic intensity. Beating MLIR outright would need a MULTI-kernel | ||
| // winograd (transform kernels + a library batched GEMM on materialized V/M). |
Collaborator
There was a problem hiding this comment.
I don't think we should be commenting on the exact perf differences. General observations are fine.
kahmed10
reviewed
Jul 30, 2026
Comment on lines
+673
to
+678
| // Measured (gfx1201 fp32, exhaustive-tune, tight-interleaved) shapes in the | ||
| // spatial-16..64 high-channel band where S-store beats full U by >=5%. The band's | ||
| // S-vs-U win/loss is micro-architecturally NON-MONOTONIC -- 512->512@16 wins but | ||
| // 515->512@16 loses 1.5x; 192->191@64 wins but 192->192@64 loses; 768->383@32 | ||
| // wins but 384->384@32 loses -- so a smooth rule can't separate them without | ||
| // regressing real full-U winners. Hence a measured table, like the fp16 path. |
kahmed10
reviewed
Jul 30, 2026
Comment on lines
+730
to
+751
| // Heuristic for when the fp32 FMA/DPP F(2,3) winograd kernel beats the default | ||
| // (rocMLIR implicit-GEMM) lowering on gfx12. Derived from a 3x3/pad-1/stride-1 | ||
| // sweep of real-model shapes (tools/bench_conv.py, exhaustive-tune) with the | ||
| // kernel's own weight-store selection (S-store / v-inner) active. Structure | ||
| // mirrors the fp16 winograd_f23_profitable, but the thresholds differ: the fp32 | ||
| // kernel has 2.25x fewer MACs than MLIR yet a heavier input/weight scatter, so it | ||
| // wins the compute-bound low/mid-channel shapes and loses the memory-bandwidth- | ||
| // bound high-channel large-spatial ones. | ||
| // - NHWC: rocMLIR's channels-last GEMM reads the input fully coalesced and wins | ||
| // almost everywhere; the winograd kernel's C-strided input scatter only pays | ||
| // off at tiny spatial + high channels (measured geomean ~0.92x overall, wins | ||
| // only at spatial<=16, min_ch>=256). So NHWC is gated to that narrow zone. | ||
| // - NCHW: winograd wins broadly (count-weighted ~1.25x on the measured set). | ||
| // Excluded regions: | ||
| // * C*K >= 700k: bandwidth-bound big GEMMs MLIR owns (1280-channel convs). | ||
| // * min(C,K) >= 224: only small spatial (<=32) wins (2.25x fewer MACs); | ||
| // mid/large spatial is input/output-transform + weight-expansion bound. | ||
| // * min(C,K) >= 128 at spatial >= 128: transform-bound, loses. | ||
| // * spatial >= 512 with out_ch > 32: the 4x input-tile re-read dominates | ||
| // unless the output fits a single KO block. | ||
| // Output-collapse (out_ch <= 3) and tiny-input (in_ch < 16) shapes are handled | ||
| // layout-independently up front. MIGRAPHX_ENABLE/DISABLE_WINOGRAD override it. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Technical Details
Changelog Category
Add a
CHANGELOG.mdentry for any option other thanNot ApplicableFollow the LLVM AI Tool Use Policy for contributions using AI.