[WORKAROUND] Disable DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 kernel until fixed in CK

src/include/miopen/solver/implicitgemm_ck_util.hpp

@@ -175,7 +175,8 @@ typename ConvPtrsType::iterator FindConvPtrByID(ConvPtrsType& conv_ptrs,
 
 template <typename DeviceOpType,
           typename CKArgsType,
-          typename ProblemDescriptionType = miopen::conv::ProblemDescription>
+          bool perform_WORKAROUND_ISSUE_3661 = false,
+          typename ProblemDescriptionType    = miopen::conv::ProblemDescription>
 std::vector<std::string> FillValidKernelsIDs(const ProblemDescriptionType& problem)
 {
     const auto args      = CKArgsType{problem};
@@ -186,8 +187,19 @@ std::vector<std::string> FillValidKernelsIDs(const ProblemDescriptionType& probl
     valid_kernels.reserve(conv_ptrs.size());
     for(size_t idx = 0; idx < conv_ptrs.size(); ++idx)
     {
-        if(args.IsSupportedBy(conv_ptrs[idx]))
-            valid_kernels.emplace_back(std::move(conv_ptrs[idx]->GetTypeString()));
+        if constexpr(perform_WORKAROUND_ISSUE_3661)
+        {
+            std::string typeString = conv_ptrs[idx]->GetTypeString();
+            // cppcheck-suppress stlIfStrFind; No string::starts_with() in c++17
+            if(typeString.find("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3") != 0 &&
+               args.IsSupportedBy(conv_ptrs[idx]))
+                valid_kernels.emplace_back(std::move(typeString));
+        }
+        else
+        {
+            if(args.IsSupportedBy(conv_ptrs[idx]))
+                valid_kernels.emplace_back(conv_ptrs[idx]->GetTypeString());
+        }
     }
     assert(!valid_kernels.empty());
     return valid_kernels;
@@ -263,14 +275,26 @@ bool IsCKArgsSupported(const ProblemDescriptionType& problem, const std::string&
 
 template <typename DeviceOpType,
           typename CKArgsType,
-          typename ProblemDescriptionType = miopen::conv::ProblemDescription>
+          bool perform_WORKAROUND_ISSUE_3661 = false,
+          typename ProblemDescriptionType    = miopen::conv::ProblemDescription>
 bool IsCKApplicable(const ProblemDescriptionType& problem)
 {
     const auto args = CKArgsType{problem};
 
     const auto ptrs = DeviceOpType::GetInstances();
-    return std::any_of(
-        ptrs.begin(), ptrs.end(), [&args](auto& ptr) { return args.IsSupportedBy(ptr); });
+
+    return std::any_of(ptrs.begin(), ptrs.end(), [&args](auto& ptr) {
+        if constexpr(perform_WORKAROUND_ISSUE_3661)
+        {
+            return ptr->GetTypeString().find("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3") !=
+                       0 &&
+                   args.IsSupportedBy(ptr);
+        }
+        else
+        {
+            return args.IsSupportedBy(ptr);
+        }
+    });
 }
 
 #define WORKAROUND_CK_ISSUE_1184 1

src/solver/conv/conv_hip_implicit_gemm_3d_grouped_fwd_xdlops.cpp

@@ -41,6 +41,11 @@
 #include <miopen/solver/implicitgemm_ck_util.hpp>
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_3D_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS)
 
+#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
+// Disable DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 until it is fixed in CK
+#define WORKAROUND_ISSUE_3661 1
+#endif
+
 namespace miopen {
 namespace solver {
 namespace conv {
@@ -349,15 +354,28 @@ void PerformanceConfigHipImplicitGemm3DGroupFwdXdlops::Init(const ProblemDescrip
     {
     case BILINEAR:
         valid_kernels =
+#if WORKAROUND_ISSUE_3661
+            FillValidKernelsIDs<DeviceOpGFwdBilinearPtrs<DataType>, CKArgs<DataType>, true>(
+                problem);
+#else
             FillValidKernelsIDs<DeviceOpGFwdBilinearPtrs<DataType>, CKArgs<DataType>>(problem);
+#endif
         break;
     case SCALE:
         valid_kernels =
+#if WORKAROUND_ISSUE_3661
+            FillValidKernelsIDs<DeviceOpGFwdScalePtrs<DataType>, CKArgs<DataType>, true>(problem);
+#else
             FillValidKernelsIDs<DeviceOpGFwdScalePtrs<DataType>, CKArgs<DataType>>(problem);
+#endif
         break;
     default:
         valid_kernels =
+#if WORKAROUND_ISSUE_3661
+            FillValidKernelsIDs<DeviceOpGFwdDefaultPtrs<DataType>, CKArgs<DataType>, true>(problem);
+#else
             FillValidKernelsIDs<DeviceOpGFwdDefaultPtrs<DataType>, CKArgs<DataType>>(problem);
+#endif
         break;
     }
     index     = 0;
@@ -389,9 +407,17 @@ bool ConvHipImplicitGemm3DGroupFwdXdlops::CheckCKApplicability(
     switch(problem.GetAlphaBetaCase())
     {
     case BILINEAR:
+#if WORKAROUND_ISSUE_3661
+        return IsCKApplicable<DeviceOpGFwdBilinearPtrs<DataType>, CKArgs<DataType>, true>(problem);
+    case SCALE:
+        return IsCKApplicable<DeviceOpGFwdScalePtrs<DataType>, CKArgs<DataType>, true>(problem);
+    default:
+        return IsCKApplicable<DeviceOpGFwdDefaultPtrs<DataType>, CKArgs<DataType>, true>(problem);
+#else
         return IsCKApplicable<DeviceOpGFwdBilinearPtrs<DataType>, CKArgs<DataType>>(problem);
     case SCALE: return IsCKApplicable<DeviceOpGFwdScalePtrs<DataType>, CKArgs<DataType>>(problem);
     default: return IsCKApplicable<DeviceOpGFwdDefaultPtrs<DataType>, CKArgs<DataType>>(problem);
+#endif
     }
 }
 #endif

src/solver/conv/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp

@@ -41,6 +41,11 @@
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS)
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_GROUP_CONV_IMPLICIT_GEMM_HIP_FWD_XDLOPS_AI_HEUR)
 
+#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
+// Disable DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 until it is fixed in CK
+#define WORKAROUND_ISSUE_3661 1
+#endif
+
 namespace miopen {
 namespace solver {
 namespace conv {
@@ -185,7 +190,11 @@ void PerformanceConfigHipImplicitGemmGroupFwdXdlops::Init(
     const ProblemDescription& problem) // should be parameterized with execution context
 {
     if(valid_kernels.empty())
+#if WORKAROUND_ISSUE_3661
+        valid_kernels = FillValidKernelsIDs<DeviceOpGFwdPtrs<DataType>, CKArgs, true>(problem);
+#else
         valid_kernels = FillValidKernelsIDs<DeviceOpGFwdPtrs<DataType>, CKArgs>(problem);
+#endif
     index     = 0;
     kernel_id = valid_kernels[index];
 }
@@ -201,7 +210,11 @@ template <typename DataType>
 bool ConvHipImplicitGemmGroupFwdXdlops::CheckCKApplicability(
     const ProblemDescription& problem) const
 {
+#if WORKAROUND_ISSUE_3661
+    return IsCKApplicable<DeviceOpGFwdPtrs<DataType>, CKArgs, true>(problem);
+#else
     return IsCKApplicable<DeviceOpGFwdPtrs<DataType>, CKArgs>(problem);
+#endif
 }
 
 #if MIOPEN_ENABLE_AI_KERNEL_TUNING
@@ -326,8 +339,12 @@ template <typename DataType>
 bool PerformanceConfigHipImplicitGemmGroupFwdXdlops::RunParameterPredictionModel(
     const ExecutionContext& ctx, const ProblemDescription& problem)
 {
+#if WORKAROUND_ISSUE_3661
+    valid_kernels = FillValidKernelsIDs<DeviceOpGFwdPtrs<DataType>, CKArgs, true>(problem);
+#else
     valid_kernels = FillValidKernelsIDs<DeviceOpGFwdPtrs<DataType>, CKArgs>(
         problem); // filter valid_kernel ID's
+#endif
     static const std::string& arch = ctx.GetStream().GetDeviceName();
     if(arch == "gfx90a")
         InitHeuristicKernelIDs("DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle");

test/gtest/group_conv.hpp

@@ -335,7 +335,7 @@ struct GroupConvTestFix
     void RunSolverImpl(const ConvTensorsType& tensors, const ProblemDescription& problem)
     {
 
-        std::cout << conv_config << std::endl;
+        MIOPEN_LOG_I(conv_config);
         auto&& handle = get_handle();
 
         Solver solv{};