[HIP][iGEMM] limit the range of several CK HIP kernels (#2005)

ROCm · Mar 14, 2023 · f26f2fe · f26f2fe
1 parent 6ea823a
commit f26f2fe
Show file tree

Hide file tree

Showing 7 changed files with 20 additions and 0 deletions.
diff --git 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 <ck/library/tensor_operation_instance/gpu/convolution_backward_data.hpp>
 #endif
+#include <miopen/solver/implicitgemm_util.hpp>
 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<ck::half_t>(problem);

diff --git 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
@@ -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
@@ -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
@@ -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
@@ -34,6 +34,7 @@
 #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
 #include <ck/library/tensor_operation_instance/gpu/convolution_forward.hpp>
 #endif
+#include <miopen/solver/implicitgemm_util.hpp>
 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
@@ -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;