diff --git a/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp index 91284c2194..6c55310ab6 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp @@ -34,6 +34,7 @@ #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL #include #endif +#include MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_XDLOPS) namespace miopen { @@ -367,6 +368,8 @@ bool ConvHipImplicitGemmBwdXdlops::IsApplicable(const ConvolutionContext& ctx, return false; if(arch == "gfx90a" && problem.conv_problem.IsGfx90aFp16altRequired()) return false; + if(!IsIndexRangeLargeEnough(problem)) + return false; switch(problem.conv_problem.GetInDataType()) { case miopenHalf: return CheckCKApplicability(problem); diff --git a/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp b/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp index f3daccd977..012f98e742 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp @@ -650,6 +650,8 @@ bool ConvHipImplicitGemmBwdDataV1R1::IsApplicable(const ConvolutionContext& ctx, return false; if(problem.group_counts != 1) return false; + if(!IsIndexRangeLargeEnough(problem)) + return false; #if WORKAROUND_ISSUE_309 if(problem.IsBfp16()) return false; diff --git a/src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp b/src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp index edebdccff1..dae4a6e7a9 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp @@ -759,6 +759,10 @@ bool ConvHipImplicitGemmBwdDataV4R1::IsApplicable(const ConvolutionContext& ctx, { return false; } + + if(!IsIndexRangeLargeEnough(problem)) + return false; + int gemm_m = 0; int gemm_n = 0; diff --git a/src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp b/src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp index 311d01be32..e8c78cab3e 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_v4r1.cpp @@ -57,6 +57,8 @@ bool ConvHipImplicitGemmV4R1Fwd::IsApplicable(const ConvolutionContext& ctx, return false; if(!problem.IsFp32() && !problem.IsFp16() && !problem.IsBfp16()) return false; + if(!IsIndexRangeLargeEnough(problem)) + return false; if(!problem.IsLayoutDefault()) return false; if(ctx.GetStream().GetDeviceName() == "gfx90a" && @@ -95,6 +97,8 @@ bool ConvHipImplicitGemmV4R1WrW::IsApplicable(const ConvolutionContext& ctx, return false; if(!problem.IsFp32() && !problem.IsFp16() && !problem.IsBfp16()) return false; + if(!IsIndexRangeLargeEnough(problem)) + return false; if(!problem.IsLayoutDefault()) return false; if(ctx.GetStream().GetDeviceName() == "gfx90a" && diff --git a/src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp b/src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp index 0149e64e5b..3b2badcd21 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp @@ -593,6 +593,8 @@ bool ConvHipImplicitGemmV4R4Fwd::IsApplicable(const ConvolutionContext& ctx, return false; if(problem.group_counts != 1) return false; + if(!IsIndexRangeLargeEnough(problem)) + return false; int gemm_m = 0; int gemm_n = 0; diff --git a/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp index e4599a1e6b..4a7f8f6979 100644 --- a/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp +++ b/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp @@ -34,6 +34,7 @@ #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL #include #endif +#include MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS) namespace miopen { @@ -347,6 +348,8 @@ bool ConvHipImplicitGemmFwdXdlops::IsApplicable(const ConvolutionContext& ctx, return false; if(arch == "gfx90a" && problem.conv_problem.IsGfx90aFp16altRequired()) return false; + if(!IsIndexRangeLargeEnough(problem)) + return false; if(!problem.IsLayoutNHWC()) return false; switch(problem.conv_problem.GetInDataType()) diff --git a/src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp b/src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp index f22665b91d..24f3e8c86b 100644 --- a/src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp +++ b/src/solver/conv_hip_implicit_gemm_wrw_v4r4.cpp @@ -596,6 +596,8 @@ bool ConvHipImplicitGemmV4R4WrW::IsApplicable(const ConvolutionContext& ctx, return false; if(problem.group_counts != 1) return false; + if(!IsIndexRangeLargeEnough(problem)) + return false; int gemm_m = 0; int gemm_n = 0;