[HOTFIX][CK] Do not call solver member functions in Invokers of ConvH…

…ipImplicitGemmFwd/BwdXdlops (#1972)
ROCm · Feb 9, 2023 · 1aff40c · 1aff40c
1 parent f6feb84
commit 1aff40c
Show file tree

Hide file tree

Showing 3 changed files with 16 additions and 20 deletions.
diff --git a/src/include/miopen/solver.hpp b/src/include/miopen/solver.hpp
@@ -5709,11 +5709,6 @@ struct ConvHipImplicitGemmFwdXdlops final
 
     template <typename DataType>
     bool CheckCKApplicability(const ProblemDescription&) const;
-    template <typename DataType>
-    void RunCKSolution(const Handle& handle,
-                       const AnyInvokeParams& primitive_parameters,
-                       const ProblemDescription& problem,
-                       const PerformanceConfigHipImplicitGemmFwdXdlops& config) const;
 };
 
 struct PerformanceConfigHipImplicitGemmBwdXdlops
@@ -5821,11 +5816,6 @@ struct ConvHipImplicitGemmBwdXdlops final
 
     template <typename DataType>
     bool CheckCKApplicability(const ProblemDescription&) const;
-    template <typename DataType>
-    void RunCKSolution(const Handle& handle,
-                       const AnyInvokeParams& primitive_parameters,
-                       const ProblemDescription& problem,
-                       const PerformanceConfigHipImplicitGemmBwdXdlops& config) const;
 };
 
 struct AnySolver;

diff --git a/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_bwd_data_xdlops.cpp
@@ -194,12 +194,13 @@ bool ConvHipImplicitGemmBwdXdlops::CheckCKApplicability(const ProblemDescription
     return false;
 }
 
+namespace {
+
 template <typename DataType>
-void ConvHipImplicitGemmBwdXdlops::RunCKSolution(
-    const Handle& handle,
-    const AnyInvokeParams& primitive_parameters,
-    const ProblemDescription& problem,
-    const PerformanceConfigHipImplicitGemmBwdXdlops& config) const
+void RunCKSolution(const Handle& handle,
+                   const AnyInvokeParams& primitive_parameters,
+                   const ProblemDescription& problem,
+                   const PerformanceConfigHipImplicitGemmBwdXdlops& config)
 {
     const auto args      = CKArgsBwd{problem};
     const auto conv_ptrs = DeviceOpBwdPtrs<DataType>::GetInstances();
@@ -245,6 +246,8 @@ void ConvHipImplicitGemmBwdXdlops::RunCKSolution(
         handle.AccumKernelTime(elapsed_time);
     }
 }
+
+} // namespace
 #endif
 
 void PerformanceConfigHipImplicitGemmBwdXdlops::HeuristicInit(const ProblemDescription& problem)

diff --git a/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_fwd_xdlops.cpp
@@ -182,12 +182,13 @@ bool ConvHipImplicitGemmFwdXdlops::CheckCKApplicability(const ProblemDescription
     return false;
 }
 
+namespace {
+
 template <typename DataType>
-void ConvHipImplicitGemmFwdXdlops::RunCKSolution(
-    const Handle& handle,
-    const AnyInvokeParams& primitive_parameters,
-    const ProblemDescription& problem,
-    const PerformanceConfigHipImplicitGemmFwdXdlops& config) const
+void RunCKSolution(const Handle& handle,
+                   const AnyInvokeParams& primitive_parameters,
+                   const ProblemDescription& problem,
+                   const PerformanceConfigHipImplicitGemmFwdXdlops& config)
 {
     const auto args      = CKArgs{problem};
     const auto conv_ptrs = DeviceOpPtrs<DataType>::GetInstances();
@@ -224,6 +225,8 @@ void ConvHipImplicitGemmFwdXdlops::RunCKSolution(
         handle.AccumKernelTime(elapsed_time);
     }
 }
+
+} // namespace
 #endif
 
 void PerformanceConfigHipImplicitGemmFwdXdlops::HeuristicInit(const ProblemDescription& problem)