[WORKAROUND][OCL][MI100][MI200] Disable MIOpenGEMM convolutions (W/A …

…for #1315). Disable iGemm ASM GTC XDLOPS NCHW convolutions (W/A for #1317) (#1321)

* W/A for #1315. Disable MIOpenGEMM convolutions for xDLOPs GPUs (MI100/MI200) && OpenCL BE
* W/A for #1317. Disable iGemm ASM GTC XDLOPS convolutions for NCHW configs && OCL BE (keep them enabled for NHWC)
* [Jenkins] Add Fp32 Full tests stages for Opencl BE && MI100/MI200
* [NFC] Fix comments related to WORKAROUND_MIOPENGEMM_SINCE_ROCM41

Jenkinsfile

@@ -789,6 +789,26 @@ pipeline {
                         buildHipClangJobAndReboot(setup_flags: Bf16_flags + Full_test, build_install: "true", gpu_arch: "gfx90a:xnack-")
                     }
                 }
+                stage('Fp32 OpenCL All gfx908') {
+                    when {
+                        beforeAgent true
+                        expression { params.TARGET_GFX908 && params.DATATYPE_FP32 }
+                    }
+                    agent{ label rocmnode("vega") }
+                    steps{
+                        buildHipClangJobAndReboot(compiler: 'g++', setup_flags: Full_test, gpu_arch: "gfx908")
+                    }
+                }
+                stage('Fp32 OpenCL Install All gfx90a') {
+                    when {
+                        beforeAgent true
+                        expression { params.TARGET_GFX90A && params.DATATYPE_FP32 }
+                    }
+                    agent{ label rocmnode("vega") }
+                    steps{
+                        buildHipClangJobAndReboot(compiler: 'g++', setup_flags: Full_test, build_install: "true", gpu_arch: "gfx90a:xnack-")
+                    }
+                }
             }
         }
 

src/CMakeLists.txt

@@ -109,6 +109,7 @@ set( MIOpen_Source
     solver/gemm.cpp
     solver/gemm_bwd.cpp
     solver/gemm_wrw.cpp
+    solver/gemm_common.cpp
     dropout.cpp
     dropout_api.cpp
     readonlyramdb.cpp
@@ -192,6 +193,7 @@ set( MIOpen_Source
     include/miopen/batched_transpose_sol.hpp
     include/miopen/magic_div.hpp
     include/miopen/util_sol.hpp
+    include/miopen/solver/gemm_common.hpp
     md_graph.cpp
     mdg_expr.cpp
     conv/invokers/gcn_asm_1x1u.cpp

src/include/miopen/conv/asm_implicit_gemm.hpp

@@ -25,12 +25,17 @@
  *******************************************************************************/
 #ifndef CK_ASM_IMPLICITGEMM_HPP_
 #define CK_ASM_IMPLICITGEMM_HPP_
+
+#include <miopen/config.h>
+
 #include <string>
 #include <ostream>
 #include <tuple>
 #include <vector>
 #include <limits>
 
+#define WORKAROUND_ISSUE_1317 (MIOPEN_BACKEND_OPENCL)
+
 namespace miopen {
 
 namespace solver {

src/include/miopen/rocm_features.hpp

@@ -55,12 +55,9 @@
 ///  --input 1, 4, 4, 161, 700 --weights 1, 4, 3, 11, 11
 ///   --pads_strides_dilations 3 3 3 2 2 2 4 4 4 --trans_output_pads 0 0 0
 ///
-/// W/A is in effect only when MIOpenGEMM is used (OCL BE) abd includes:
-/// - Disabling GEMM for failing configs.
-/// - Adding Naive Solvers. Naive solvers are inteded for use as backup for
-///   Immediate Mode Fallback when GEMM is disabled.
-/// - Note: When MIOpenGEMM is not in use, Naive Solvers are disabled. This minimizes
-///   impact of the W/A to the HIP backend.
+/// W/A is in effect only when MIOpenGEMM is used (OCL BE) and disables
+/// GEMM for the failing configs. When this happens, Naive solvers
+/// are used as backup on the Immediate Mode Fallback path.
 #define WORKAROUND_MIOPENGEMM_SINCE_ROCM41 \
     (MIOPEN_USE_MIOPENGEMM && (HIP_PACKAGE_VERSION_FLAT >= 4001000000ULL))
 

src/include/miopen/solver/gemm_common.hpp

@@ -0,0 +1,44 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2021 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#ifndef GUARD_SOLVER_GEMM_COMMON_HPP_
+#define GUARD_SOLVER_GEMM_COMMON_HPP_
+
+#include <miopen/execution_context.hpp>
+
+namespace miopen {
+namespace conv {
+namespace solver {
+namespace gemm {
+
+bool IsWorkaroundIssue1315(const miopen::ExecutionContext& ctx);
+
+} // namespace gemm
+} // namespace solver
+} // namespace conv
+} // namespace miopen
+
+#endif

src/solver/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp

@@ -867,6 +867,11 @@ ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::Search(const ConvolutionContext& ctx
 
 bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsApplicable(const ConvolutionContext& ctx) const
 {
+#if WORKAROUND_ISSUE_1317
+    if(ctx.IsLayoutDefault())
+        if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC{}))
+            return false;
+#endif
     if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_BWD_GTC_XDLOPS_NHWC{}))
         return false;
 

src/solver/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp

@@ -792,6 +792,11 @@ ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::GetWorkspaceSize(const ConvolutionCo
 
 bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable(const ConvolutionContext& ctx) const
 {
+#if WORKAROUND_ISSUE_1317
+    if(ctx.IsLayoutDefault())
+        if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC{}))
+            return false;
+#endif
     if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_FWD_GTC_XDLOPS_NHWC{}))
         return false;
 

src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp

@@ -795,6 +795,11 @@ ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::Search(const ConvolutionContext& ctx
 
 bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsApplicable(const ConvolutionContext& ctx) const
 {
+#if WORKAROUND_ISSUE_1317
+    if(ctx.IsLayoutDefault())
+        if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC{}))
+            return false;
+#endif
     if(miopen::IsDisabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_ASM_WRW_GTC_XDLOPS_NHWC{}))
         return false;
 

src/solver/gemm.cpp

@@ -32,6 +32,7 @@
 #include <miopen/handle.hpp>
 #include <miopen/kernel.hpp>
 #include <miopen/rocm_features.hpp>
+#include <miopen/solver/gemm_common.hpp>
 #include <miopen/tensor.hpp>
 #include <miopen/tensor_ops.hpp>
 #include <miopen/util.hpp>
@@ -77,17 +78,20 @@ static inline bool IsAnyBufferFp16(const TensorDescriptor& xDesc,
 }
 #endif
 
-bool GemmFwdBase::IsApplicable(const ExecutionContext&,
+bool GemmFwdBase::IsApplicable(const ExecutionContext& ctx,
                                const conv::ProblemDescription& problem) const
 {
 #if MIOPEN_USE_GEMM
+    if(conv::solver::gemm::IsWorkaroundIssue1315(ctx))
+        return false;
     const auto& xDesc = problem.GetIn();
     const auto& wDesc = problem.GetWeights();
     const auto& yDesc = problem.GetOut();
     return problem.GetDirection() == conv::Direction::Forward && problem.IsLayoutDefault() &&
            !(IsAnyBufferBF16(xDesc, yDesc, wDesc) && !IsBf16Supported) &&
            !(IsAnyBufferFp16(xDesc, yDesc, wDesc) && !IsFp16Supported);
 #else
+    std::ignore = ctx;
     std::ignore = problem;
     return false;
 #endif

src/solver/gemm_bwd.cpp

@@ -31,6 +31,7 @@
 #include <miopen/gemm_v2.hpp>
 #include <miopen/handle.hpp>
 #include <miopen/kernel.hpp>
+#include <miopen/solver/gemm_common.hpp>
 #include <miopen/tensor.hpp>
 #include <miopen/tensor_ops.hpp>
 #include <miopen/util.hpp>
@@ -91,17 +92,20 @@ SlowdownFactor(int n_oper, const double oper_factor, const double multiple_oper_
         return 1.0;
 }
 
-bool GemmBwdBase::IsApplicable(const ExecutionContext&,
+bool GemmBwdBase::IsApplicable(const ExecutionContext& ctx,
                                const conv::ProblemDescription& problem) const
 {
 #if MIOPEN_USE_GEMM
+    if(conv::solver::gemm::IsWorkaroundIssue1315(ctx))
+        return false;
     const auto& dyDesc = problem.GetIn();
     const auto& wDesc  = problem.GetWeights();
     const auto& dxDesc = problem.GetOut();
     return problem.GetDirection() == conv::Direction::BackwardData && problem.IsLayoutDefault() &&
            !(IsAnyBufferBF16(dxDesc, dyDesc, wDesc) && !IsBf16Supported) &&
            !(IsAnyBufferFp16(dxDesc, dyDesc, wDesc) && !IsFp16Supported);
 #else
+    std::ignore = ctx;
     std::ignore = problem;
     return false;
 #endif

src/solver/gemm_common.cpp

@@ -0,0 +1,59 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2021 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+
+#include <miopen/config.h>
+#include <miopen/solver/gemm_common.hpp>
+
+#include <tuple> // std::ignore
+
+/// This W/A disables all GEMM convolution solvers for xDLOPs
+/// targets when MIOpenGEMM is used (OCL BE). More info at
+/// https://github.com/ROCmSoftwarePlatform/MIOpen/issues/1315.
+///
+/// W/A affects ROCm releases starting from 4.5 and also
+/// pre-5.0 Mainline HIP builds, e.g. 9148.
+#define WORKAROUND_ISSUE_1315 (MIOPEN_USE_MIOPENGEMM && (HIP_PACKAGE_VERSION_FLAT >= 4004000000ULL))
+
+namespace miopen {
+namespace conv {
+namespace solver {
+namespace gemm {
+
+bool IsWorkaroundIssue1315(const miopen::ExecutionContext& ctx)
+{
+#if WORKAROUND_ISSUE_1315
+    const auto device = ctx.GetStream().GetTargetProperties().Name();
+    return (device == "gfx908") || (device == "gfx90a");
+#else
+    std::ignore = ctx;
+    return false;
+#endif
+}
+
+} // namespace gemm
+} // namespace solver
+} // namespace conv
+} // namespace miopen

src/solver/gemm_wrw.cpp

@@ -3,6 +3,7 @@
 #include <miopen/conv/wrw_invoke_params.hpp>
 #include <miopen/errors.hpp>
 #include <miopen/gemm_v2.hpp>
+#include <miopen/solver/gemm_common.hpp>
 #include <miopen/tensor_ops.hpp>
 #include <miopen/util.hpp>
 
@@ -60,10 +61,12 @@ SlowdownFactor(int n_oper, const double oper_factor, const double multiple_oper_
         return 1.0;
 }
 
-bool GemmWrwBase::IsApplicable(const ExecutionContext&,
+bool GemmWrwBase::IsApplicable(const ExecutionContext& ctx,
                                const conv::ProblemDescription& problem) const
 {
 #if MIOPEN_USE_GEMM
+    if(conv::solver::gemm::IsWorkaroundIssue1315(ctx))
+        return false;
     const auto& dyDesc = problem.GetIn();
     const auto& dwDesc = problem.GetWeights();
     const auto& xDesc  = problem.GetOut();
@@ -72,6 +75,7 @@ bool GemmWrwBase::IsApplicable(const ExecutionContext&,
            !(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsBF16PathValid) &&
            !(IsAnyBufferFp16(xDesc, dyDesc, dwDesc) && !IsFp16Supported);
 #else
+    std::ignore = ctx;
     std::ignore = problem;
     return false;
 #endif