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.

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

Conversation

carlushuang
Copy link
Contributor

@carlushuang carlushuang commented Aug 5, 2022

#1619 originally try to fix SWDEV-306318, but introduced large performance regression.
This PR try to narrow down the non-applicable range, only when 1x1 conv, C is not multiple of 8, this solver is not applicable. This is true for both fp32 and fp16

cc @atamazov

Verified

This commit was created on GitHub.com and signed with GitHub’s verified signature. The key has expired.
Copy link
Contributor

@DrizztDoUrden DrizztDoUrden left a comment

Choose a reason for hiding this comment

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

lgtm

@junliume junliume changed the title WA PR:1619 by only ignore 1x1 conv where C is not 8x [SWDEV-306318][WORKAROUND][gfx908] Limit WA PR:1619 by only ignore 1x1 conv where C is not 8x Aug 5, 2022
@junliume junliume merged commit ea3c408 into develop Aug 5, 2022
@junliume junliume deleted the wa_1619_perf_regre branch October 10, 2022 23:47
@@ -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

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

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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

5 participants