Skip to content
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

[SWDEV-306318][WORKAROUND][gfx908] Limit WA PR:1619 by only ignore 1x1 conv where C is not 8x #1675

Merged
merged 1 commit into from
Aug 5, 2022
Merged

[SWDEV-306318][WORKAROUND][gfx908] Limit WA PR:1619 by only ignore 1x1 conv where C is not 8x #1675

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
10 changes: 5 additions & 5 deletions src/solver/conv_asm_implicit_gemm_gtc_fwd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1504,11 +1504,6 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlops::IsApplicable(const ConvolutionConte
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS{}))
return false;

#if WORKAROUND_SWDEV_306318
if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS{}))
return false;
#endif

const auto device_name = ctx.GetStream().GetDeviceName();
if(device_name != "gfx908")
return false;
Expand Down Expand Up @@ -1536,6 +1531,11 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlops::IsApplicable(const ConvolutionConte
return false;
}

#if WORKAROUND_SWDEV_306318
if((ctx.kernel_size_h == 1) && (ctx.kernel_size_w == 1) && (ctx.n_inputs % 8 != 0))
return false;
Copy link
Contributor

@atamazov atamazov Feb 19, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is recommended to insert

if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS{}))

just above return false. This would allow explicit enabling the solver for 1x1 C%8 != 0 configs for triaging, without rebuilding the library.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@atamazov Do you mean with the current commits, 1x1 c%8 !=0 configs could not be tested, even after the kernels being improved? and by inserting
if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS{}))
we are able to test 1x1 c%8 !=0 by using
export MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS=1 ?

So my understand is that the env var MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS should have three states:

  1. Disabled ( defined from env and have false value)
  2. Enabled ( defined from env and have true value)
  3. Neither Disabled nor Enabled ( not defined from env)

right?

Copy link
Contributor

@atamazov atamazov Feb 21, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@qianfengz Yes, and in the latter case both IsEnabled and IsDisabled return false. This is useful in cases like this one.

#endif

const auto target = ctx.GetStream().GetTargetProperties();
if(target.Xnack() && *target.Xnack())
return false;
Expand Down