Skip to content

feat(miopen): Add gfx120x support to Winograd Rage #6439

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
k-artem merged 22 commits into from
Jul 17, 2026
Merged

feat(miopen): Add gfx120x support to Winograd Rage #6439

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
22 commits
Select commit Hold shift + click to select a range
154b76b
[miopen][winograd][gfx12] Add gfx12 support to Winograd Rage
scerzh Apr 9, 2026
5ea85dd
[miopen][winograd] Address review commnents and const the strings
scerzh Apr 15, 2026
551b5af
Merge remote-tracking branch 'origin/develop' into users/scerzh/winog…
scerzh Apr 15, 2026
b8d1a9a
Fix compile errors
scerzh Apr 15, 2026
b797676
Fix clang-format
scerzh Apr 15, 2026
84b6415
Merge branch 'develop' into users/scerzh/winograd-rage-add-gfx12
k-artem Apr 17, 2026
b017701
Merge remote-tracking branch 'origin/develop' into users/scerzh/winog…
scerzh Apr 20, 2026
57519a3
Address PR comments
scerzh Apr 23, 2026
320a452
Merge remote-tracking branch 'origin/develop' into users/scerzh/winog…
scerzh Apr 23, 2026
b2220c0
Merge branch 'develop' into users/scerzh/winograd-rage-add-gfx12
k-artem Apr 23, 2026
9a8d6f0
Merge remote-tracking branch 'origin/develop' into users/scerzh/winog…
scerzh Apr 28, 2026
fefbef8
Remove Shader Model for gfx942
scerzh Apr 27, 2026
b90d9da
Merge branch 'develop' into users/scerzh/winograd-rage-add-gfx12
k-artem May 6, 2026
9856749
aligned checks inside WinoRage code after 1250 was added
k-artem Jun 18, 2026
4b251f8
Merge branch 'develop' into users/scerzh/winograd-rage-add-gfx12
k-artem Jun 18, 2026
897cd81
Merge remote-tracking branch 'origin/develop' into users/scerzh/winog…
scerzh Jun 30, 2026
79a71ca
[miopen][winograd][gfx12] Re-register TransposedConvWinoRageRxS and a…
scerzh Jun 29, 2026
f9307da
Merge branch 'develop' into users/scerzh/winograd-rage-add-gfx12
k-artem Jul 1, 2026
37d5691
ci: retrigger policy check
scerzh Jul 8, 2026
4065b75
Merge branch 'develop' into users/scerzh/winograd-rage-add-gfx12
k-artem Jul 13, 2026
7475341
Merge branch 'develop' into users/scerzh/winograd-rage-add-gfx12
k-artem Jul 15, 2026
4661d52
Merge branch 'develop' into users/scerzh/winograd-rage-add-gfx12
scerzh Jul 16, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions projects/miopen/docs/reference/env_variables.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -364,6 +364,11 @@ For more information, see :doc:`Logging and debugging <../how-to/debug-log>`.
- | 0: Disable
| 1: Enable

* - | ``MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3``
| Controls Winograd Rage RxS F(2,3) solution.
- | 0: Disable
| 1: Enable

Multi-pass Winograd solution control
====================================

Expand Down
5 changes: 5 additions & 0 deletions projects/miopen/src/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -514,8 +514,11 @@ if( MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN_BACKEND STREQUAL "HIP" OR MIOPEN_B
kernels/winograd/Conv_Winograd_Fury_v2_X_X_X_metadata.inc
kernels/winograd/Conv_Winograd_Rage_v4_6_0_gfx94x_fp16_fp32acc_f2x3_stride1.inc
kernels/winograd/Conv_Winograd_Rage_v4_6_0_metadata.inc
kernels/winograd/Conv_Winograd_Rage_v4_6_1_gfx12x_fp16_fp32acc_f2x3_stride1.inc
kernels/winograd/Conv_Winograd_Rage_v4_7_0_gfx94x_fp16_fp32acc_f2x3_stride1.inc
kernels/winograd/Conv_Winograd_Rage_v4_7_0_metadata.inc
kernels/winograd/Conv_Winograd_Rage_v4_9_1_gfx12x_fp16_fp32acc_f2x3_stride1.inc
kernels/winograd/Conv_Winograd_Rage_v4_X_X_metadata.inc
kernels/xform_bidirect_winograd_code.inc
kernels/xform_data_filter.inc
kernels/xform_kd_cov2.inc
Expand Down Expand Up @@ -658,7 +661,9 @@ if( MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN_BACKEND STREQUAL "HIP" OR MIOPEN_B
kernels/winograd/Conv_Winograd_Fury_v4_6_0_fp16_fp32acc_f2x3_c16_stride1.s
kernels/winograd/Conv_Winograd_Fury_v4_6_0_fp16_fp32acc_f2x3_c32_stride1.s
kernels/winograd/Conv_Winograd_Rage_v4_6_0_fp16_fp32acc_f2x3_stride1.s
kernels/winograd/Conv_Winograd_Rage_v4_6_1_fp16_fp32acc_f2x3_stride1.s
kernels/winograd/Conv_Winograd_Rage_v4_7_0_fp16_fp32acc_f2x3_stride1.s
kernels/winograd/Conv_Winograd_Rage_v4_9_1_fp16_fp32acc_f2x3_stride1.s
kernels/xform_data.s
kernels/xform_filter.s
kernels/xform_bidirect_winograd_data.s
Expand Down
2 changes: 2 additions & 0 deletions projects/miopen/src/env_debug.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,7 @@ MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X4)
MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5)
MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6)
MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1)
MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3)
MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS)
MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW)
MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_DIRECT)
Expand Down Expand Up @@ -187,6 +188,7 @@ const LibEnvVar& FindEnvVariable(std::string_view name)
{MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5.GetName(), MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X5},
{MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6.GetName(), MIOPEN_DEBUG_AMD_WINOGRAD_MPASS_F3X6},
{MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1.GetName(), MIOPEN_DEBUG_AMD_WINOGRAD_RXS_F2X3_G1},
{MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3.GetName(), MIOPEN_DEBUG_AMD_WINOGRAD_RAGE_RXS_F2X3},
{MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS.GetName(), MIOPEN_DEBUG_CHECK_SUB_BUFFER_OOB_MEMORY_ACCESS},
{MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW.GetName(), MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW},
{MIOPEN_DEBUG_CONV_DIRECT.GetName(), MIOPEN_DEBUG_CONV_DIRECT},
Expand Down
2 changes: 1 addition & 1 deletion projects/miopen/src/kernels/winograd/Conv_Winograd_Rage_v4_6_0_metadata.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -179,7 +179,7 @@ __fp16_overflow = 0
.amdhsa_dx10_clamp __dx10_clamp
.amdhsa_ieee_mode __ieee_mode
.amdhsa_fp16_overflow __fp16_overflow
/* \todo add gfx12 support with these directives
/* \todo add gfx120 support with these directives
.amdhsa_wavefront_size32 0
.amdhsa_workgroup_processor_mode 0
*/
Expand Down
14 changes: 14 additions & 0 deletions projects/miopen/src/kernels/winograd/Conv_Winograd_Rage_v4_6_1_fp16_fp32acc_f2x3_stride1.s
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,14 @@
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT

.include "Conv_Winograd_Rage_v4_X_X_metadata.inc"

.if (.amdgcn.gfx_generation_number == 12 && .amdgcn.gfx_generation_minor == 0)
KERNEL_PROLOG 4_6_1, _fp16_fp32acc_f2x3_stride1
.noaltmacro
.include "Conv_Winograd_Rage_v4_6_1_gfx12x_fp16_fp32acc_f2x3_stride1.inc"
.altmacro
KERNEL_EPILOG 80, 250, 0, 384, 4_6_1, _fp16_fp32acc_f2x3_stride1
.else
.error "Unsupported gfx version"
.endif
Loading
Loading