Add fp32 winograd for gfx12 - #5092
Open
pfultz2 wants to merge 26 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% -0.00%
===========================================
Files 623 623
Lines 32969 33097 +128
===========================================
+ Hits 30747 30866 +119
- Misses 2222 2231 +9 🚀 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
+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. |
CharlieL7
requested changes
Aug 6, 2026
CharlieL7
left a comment
CharlieL7
left a comment
Collaborator
There was a problem hiding this comment.
I would like to see more tests with non-randomized data to make sure this is accurate.
Comment on lines
+270
to
+272
| const int32_t w_u_stride = static_cast<int32_t>(w_str[0] * sizeof(float)); | ||
| const int32_t w_k_stride = static_cast<int32_t>(w_str[w_k_dim] * sizeof(float)); | ||
| const int32_t w_c_stride = static_cast<int32_t>(w_str[w_c_dim] * sizeof(float)); |
Collaborator
There was a problem hiding this comment.
These static_cast don't look necessary.
| // 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. | ||
| bool winograd_f23_fp32_profitable( |
Collaborator
There was a problem hiding this comment.
Can you explain more simply what this heuristic is?
Collaborator
Author
There was a problem hiding this comment.
Its a heuristic to pick winograd when it is faster than MLIR.
| // by the fp32 winograd heuristic, so it exercises the kernel by default. Odd | ||
| // spatial size exercises the boundary tiles (halo padding); the channel count is | ||
| // not a multiple of the per-lane output block so the partial-KO store path is | ||
| // covered. MIGRAPHX_DISABLE_WINOGRAD forces the default lowering for comparison. |
Collaborator
There was a problem hiding this comment.
How is the env variable set to compare?
Collaborator
Author
There was a problem hiding this comment.
You set MIGRAPHX_DISABLE_WINOGRAD=1 to get the baseline numbers.
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
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.