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

[HOTFIX][SWDEV-326460] check MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC in IsApplicable() #1451

Merged
merged 2 commits into from
Mar 8, 2022
Merged

[HOTFIX][SWDEV-326460] check MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC in IsApplicable() #1451

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0d6d34d
check MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC in IsApplicable()
carlushuang Mar 5, 2022
e37c4b0
fix clang format issue
junliume Mar 7, 2022
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
4 changes: 2 additions & 2 deletions include/miopen/miopen.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2436,8 +2436,8 @@ miopenBatchNormalizationForwardTraining(miopenHandle_t handle,
*
* If either estimatedMean, or estimatedVariance are null pointers then the values for the mean and
* variance will be calculated from input data and this calculated mean and variance will be used
* to update input values.
* If variance is zero and epsilon is also zero, this function outputs NAN values. Input espilon
* to update input values.
* If variance is zero and epsilon is also zero, this function outputs NAN values. Input espilon
* value should always be non zero positive value.
*
* @param handle MIOpen handle (input)
Expand Down
10 changes: 3 additions & 7 deletions src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -782,13 +782,6 @@ bool PerformanceConfigAsmImplicitGemmGTCBwdXdlopsNHWC::IsValid(const Convolution
if(ctx.IsFp16() && gemm_k_global_split != 0 && vector_store != 1)
return false;

// An internal rule when we create the kernel_param_list. There always will be a config with
// gemm_k_global_split=0 in front of a almost the same config, with only gemm_k_global_split=1,
// so it will never choose the *=1 one. And indeed in the HeuristicInit() we never wish to find
// a config with gemm_k_global_split=1. So it is safe here to disable *=1 case
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}) && gemm_k_global_split != 0)
return false;

const auto group = ctx.group_counts;
const auto k = ctx.n_inputs;
const auto c = ctx.n_outputs;
Expand Down Expand Up @@ -893,6 +886,9 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsApplicable(const ConvolutionC
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC{}))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

const auto device_name = ctx.GetStream().GetDeviceName();
if((device_name != "gfx908") && (device_name != "gfx90a"))
return false;
Expand Down
6 changes: 3 additions & 3 deletions src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -655,9 +655,6 @@ bool PerformanceConfigAsmImplicitGemmGTCFwdXdlopsNHWC::IsValid(const Convolution
if(ctx.IsFp16() && gemm_k_global_split != 0 && vector_store != 1)
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}) && gemm_k_global_split != 0)
return false;

const auto& c = ctx.n_inputs;
const auto& k = ctx.n_outputs;
const auto& group = ctx.group_counts;
Expand Down Expand Up @@ -826,6 +823,9 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable(const ConvolutionC
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC{}))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

const auto device_name = ctx.GetStream().GetDeviceName();
if((device_name != "gfx908") && (device_name != "gfx90a"))
return false;
Expand Down
6 changes: 3 additions & 3 deletions src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -727,9 +727,6 @@ bool PerformanceConfigAsmImplicitGemmGTCWrwXdlopsNHWC::IsValid(const Convolution
if(ctx.IsFp16() && tensor_b_thread_lengths[3] != 1 && gemm_k_global_split != 0 && vector_store != 1)
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}) && gemm_k_global_split != 0)
return false;

const auto& k = ctx.n_inputs;
const auto& c = ctx.n_outputs;
const auto& y = ctx.kernel_size_h;
Expand Down Expand Up @@ -801,6 +798,9 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsApplicable(const ConvolutionC
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC{}))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

const auto device_name = ctx.GetStream().GetDeviceName();
if((device_name != "gfx908") && (device_name != "gfx90a"))
return false;
Expand Down
3 changes: 3 additions & 0 deletions src/solver/conv_asm_implicit_gemm_wrw_gtc_dynamic_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -810,6 +810,9 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlops::IsApplicable(const ConvolutionConte
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS{}))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

const auto device_name = ctx.GetStream().GetDeviceName();
if(device_name != "gfx908")
return false;
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -638,6 +638,8 @@ bool ConvHipImplicitGemmBwdDataV1R1::IsApplicable(const ConvolutionContext& ctx)
{
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!ctx.use_hip_kernels)
return false;
if(!ctx.IsLayoutDefault())
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_hip_implicit_gemm_bwd_v1r1_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -751,6 +751,8 @@ bool ConvHipImplicitGemmBwdDataV1R1Xdlops::IsApplicable(const ConvolutionContext
#endif
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

if(!IsComposableKernelSupportedHardware(ctx))
return false;
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -749,6 +749,8 @@ bool ConvHipImplicitGemmBwdDataV4R1::IsApplicable(const ConvolutionContext& ctx)

if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
shaojiewang marked this conversation as resolved.
Show resolved Hide resolved

if(!IsComposableKernelSupportedHardware(ctx))
return false;
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_hip_implicit_gemm_bwd_v4r1_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -824,6 +824,8 @@ bool ConvHipImplicitGemmBwdDataV4R1Xdlops::IsApplicable(const ConvolutionContext
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS{}))
return false;
#endif
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!IsComposableKernelSupportedHardware(ctx))
return false;
if(!ctx.direction.IsBackwardData())
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,8 @@ bool ConvHipImplicitGemmV4R1Fwd::IsApplicable(const ConvolutionContext& ctx) con
return false;
if(!IsComposableKernelSupportedHardware(ctx))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!ctx.direction.IsForward())
return false;
if(!ctx.use_hip_kernels)
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -586,6 +586,8 @@ bool ConvHipImplicitGemmV4R4Fwd::IsApplicable(const ConvolutionContext& ctx) con
{
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!ctx.use_hip_kernels)
return false;
if(!ctx.IsLayoutDefault())
Expand Down
3 changes: 3 additions & 0 deletions src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -967,6 +967,9 @@ bool ConvHipImplicitGemmForwardV4R4Xdlops::IsApplicable(const ConvolutionContext
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_XDLOPS{}))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

if(!ctx.use_hip_kernels)
return false;

Expand Down
3 changes: 3 additions & 0 deletions src/solver/conv_hip_implicit_gemm_fwd_v4r4_xdlops_padded_gemm.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1042,6 +1042,9 @@ bool ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm::IsApplicable(
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R4_PADDED_GEMM_XDLOPS{}))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

if(!ctx.use_hip_kernels)
return false;

Expand Down
3 changes: 3 additions & 0 deletions src/solver/conv_hip_implicit_gemm_fwd_v4r5_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -997,6 +997,9 @@ bool ConvHipImplicitGemmForwardV4R5Xdlops::IsApplicable(const ConvolutionContext
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R5_XDLOPS{}))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

if(!ctx.use_hip_kernels)
return false;

Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -590,6 +590,8 @@ bool ConvHipImplicitGemmV4R4WrW::IsApplicable(const ConvolutionContext& ctx) con
{
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!ctx.use_hip_kernels)
return false;
if(!ctx.IsLayoutDefault())
Expand Down
3 changes: 3 additions & 0 deletions src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1036,6 +1036,9 @@ bool ConvHipImplicitGemmWrwV4R4Xdlops::IsApplicable(const ConvolutionContext& ct
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_WRW_V4R4_XDLOPS{}))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

if(!ctx.use_hip_kernels)
return false;

Expand Down
3 changes: 3 additions & 0 deletions src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops_padded_gemm.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1110,6 +1110,9 @@ bool ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm::IsApplicable(const Convolutio
if(!IsComposableKernelSupportedHardware(ctx))
return false;

if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;

if(!IsXdlopsSupport(ctx))
return false;

Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_mlir_igemm_bwd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,6 +42,8 @@ bool ConvMlirIgemmBwd::IsApplicable(const ConvolutionContext& ctx) const
#if MIOPEN_USE_MLIR
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!ctx.direction.IsBackwardData())
return false;
if(!IsComposableKernelSupportedHardware(ctx))
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_mlir_igemm_bwd_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,8 @@ bool ConvMlirIgemmBwdXdlops::IsApplicable(const ConvolutionContext& ctx) const
#if MIOPEN_USE_MLIR
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD_XDLOPS{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!IsXdlopsSupport(ctx))
return false;
if(!ctx.direction.IsBackwardData())
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_mlir_igemm_fwd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -160,6 +160,8 @@ bool ConvMlirIgemmFwd::IsApplicable(const ConvolutionContext& ctx) const
#if MIOPEN_USE_MLIR
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_MLIR_IGEMM_FWD{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!ctx.direction.IsForward())
return false;
if(!IsComposableKernelSupportedHardware(ctx))
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_mlir_igemm_fwd_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,6 +50,8 @@ bool ConvMlirIgemmFwdXdlops::IsApplicable(const ConvolutionContext& ctx) const
#if MIOPEN_USE_MLIR
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_MLIR_IGEMM_FWD_XDLOPS{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!IsXdlopsSupport(ctx))
return false;
if(!ctx.direction.IsForward())
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_mlir_igemm_wrw.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,8 @@ bool ConvMlirIgemmWrW::IsApplicable(const ConvolutionContext& ctx) const
#if MIOPEN_USE_MLIR
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_MLIR_IGEMM_WRW{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!ctx.direction.IsBackwardWrW())
return false;
if(!IsComposableKernelSupportedHardware(ctx))
Expand Down
2 changes: 2 additions & 0 deletions src/solver/conv_mlir_igemm_wrw_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,6 +44,8 @@ bool ConvMlirIgemmWrWXdlops::IsApplicable(const ConvolutionContext& ctx) const
#if MIOPEN_USE_MLIR
if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_MLIR_IGEMM_WRW_XDLOPS{}))
return false;
if(miopen::IsEnabled(MIOPEN_DEBUG_CONVOLUTION_DETERMINISTIC{}))
return false;
if(!IsXdlopsSupport(ctx))
return false;
if(!ctx.direction.IsBackwardWrW())
Expand Down