Remove MIOpenGEMM and MIOpenTensile leftovers

README.md

@@ -32,9 +32,6 @@ python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
   * HIP - 
     * HIP and HCC libraries and header files.
   * OpenCL - OpenCL libraries and header files.
-* [MIOpenGEMM](https://github.com/ROCmSoftwarePlatform/MIOpenGEMM) - enable various functionalities including transposed and dilated convolutions. 
-  * This is optional on the HIP backend, and required on the OpenCL backend.
-  * Users can enable this library using the cmake configuration flag `-DMIOPEN_USE_MIOPENGEMM=On`, which is enabled by default when OpenCL backend is chosen.
 * [ROCm cmake](https://github.com/RadeonOpenCompute/rocm-cmake) - provide cmake modules for common build tasks needed for the ROCM software stack.
 * [Half](http://half.sourceforge.net/) - IEEE 754-based half-precision floating point library
 * [Boost](http://www.boost.org/) 
@@ -101,8 +98,6 @@ This prefix can used to specify the dependency path during the configuration pha
 
 * MIOpen's HIP backend uses [rocBLAS](https://github.com/ROCmSoftwarePlatform/rocBLAS) by default. Users can install rocBLAS minimum release by using `apt-get install rocblas`. To disable using rocBLAS set the configuration flag `-DMIOPEN_USE_ROCBLAS=Off`. rocBLAS is *not* available for the OpenCL backend.
 
-* MIOpen's OpenCL backend uses [MIOpenGEMM](https://github.com/ROCmSoftwarePlatform/MIOpenGEMM) by default. Users can install MIOpenGEMM minimum release by using `apt-get install miopengemm`.
-
 ## Building MIOpen from source
 
 ### Configuring with cmake

docs/DebugAndLogging.md

@@ -172,9 +172,7 @@ Additionally, using environment variable "MIOPEN_GEMM_ENFORCE_BACKEND", can over
 both MIOpenGEMM and rocBlas depending on the input configuration:
 
 * `MIOPEN_GEMM_ENFORCE_BACKEND=1`, use rocBLAS if enabled
-* `MIOPEN_GEMM_ENFORCE_BACKEND=2`, use MIOpenGEMM for FP32, use rocBLAS for FP16 if enabled
 * `MIOPEN_GEMM_ENFORCE_BACKEND=3`, no gemm will be called
-* `MIOPEN_GEMM_ENFORCE_BACKEND=4`, use MIOpenTensile for FP32, use rocBLAS for FP16 if enabled
 * `MIOPEN_GEMM_ENFORCE_BACKEND=<any other value>`, use default behavior
 
 To disable using rocBlas entirely, set the configuration flag `-DMIOPEN_USE_ROCBLAS=Off` during MIOpen configuration.

docs/install.md

@@ -6,9 +6,6 @@
   * HIP - 
     * HIP and HCC libraries and header files.
   * OpenCL - OpenCL libraries and header files.
-* [MIOpenGEMM](https://github.com/ROCmSoftwarePlatform/MIOpenGEMM) - enable various functionalities including transposed and dilated convolutions. 
-  * This is optional on the HIP backend, and required on the OpenCL backend.
-  * Users can enable this library using the cmake configuration flag `-DMIOPEN_USE_MIOPENGEMM=On`, which is enabled by default when OpenCL backend is chosen.
 * [ROCm cmake](https://github.com/RadeonOpenCompute/rocm-cmake) - provide cmake modules for common build tasks needed for the ROCM software stack.
 * [Half](http://half.sourceforge.net/) - IEEE 754-based half-precision floating point library
 * [Boost](http://www.boost.org/) 
@@ -73,5 +70,3 @@ cmake -P install_deps.cmake --minimum --prefix /root/MIOpen/install_dir
 This prefix can used to specify the dependency path during the configuration phase using the `CMAKE_PREFIX_PATH`.
 
 * MIOpen's HIP backend uses [rocBLAS](https://github.com/ROCmSoftwarePlatform/rocBLAS) by default. Users can install rocBLAS minimum release by using `apt-get install rocblas`. To disable using rocBLAS set the configuration flag `-DMIOPEN_USE_ROCBLAS=Off`. rocBLAS is *not* available for the OpenCL backend.
-
-* MIOpen's OpenCL backend uses [MIOpenGEMM](https://github.com/ROCmSoftwarePlatform/MIOpenGEMM) by default. Users can install MIOpenGEMM minimum release by using `apt-get install miopengemm`.

include/miopen/config.h.in

@@ -30,7 +30,6 @@
 #cmakedefine01 MIOPEN_BACKEND_HIP
 #cmakedefine01 MIOPEN_MODE_NOGPU
 #cmakedefine01 MIOPEN_USE_MIOPENTENSILE
-#cmakedefine01 MIOPEN_USE_MIOPENGEMM
 #cmakedefine01 MIOPEN_USE_ROCBLAS
 #cmakedefine01 MIOPEN_BUILD_DEV
 #cmakedefine01 MIOPEN_GPU_SYNC
@@ -86,7 +85,7 @@
 #cmakedefine MIOPEN_OFFLOADBUNDLER_BIN "@MIOPEN_OFFLOADBUNDLER_BIN@"
 #cmakedefine MIOPEN_CACHE_DIR "@MIOPEN_CACHE_DIR@"
 
-#define MIOPEN_USE_GEMM (MIOPEN_USE_MIOPENTENSILE || MIOPEN_USE_MIOPENGEMM || MIOPEN_USE_ROCBLAS)
+#define MIOPEN_USE_GEMM (MIOPEN_USE_MIOPENTENSILE || MIOPEN_USE_ROCBLAS)
 
 // Usage of "defined" operator within macro expansion is undefined behavior,
 // so "defined(NDEBUG)" cannot be used there... unlike the following macro:

src/CMakeLists.txt

@@ -237,7 +237,6 @@ set( MIOpen_Source
     solver/fft.cpp
     solver/gemm.cpp
     solver/gemm_bwd.cpp
-    solver/gemm_common.cpp
     solver/gemm_wrw.cpp
     solver/pooling/forward2d.cpp
     solver/pooling/forwardNaive.cpp

src/gemm_v2.cpp

@@ -625,8 +625,7 @@ miopenStatus_t CallGemm(const Handle& handle,
         break;
 
         case miopenDouble: {
-            MIOPEN_THROW(miopenStatusBadParm,
-                         "miopenDouble data type not supported by MIOpenGEMM.");
+            MIOPEN_THROW(miopenStatusBadParm, "miopenDouble data type not supported by rocBLAS.");
         };
         break;
         }
@@ -892,8 +891,7 @@ miopenStatus_t CallGemmStridedBatched(const Handle& handle,
         }
 
         case miopenDouble: {
-            MIOPEN_THROW(miopenStatusBadParm,
-                         "miopenDouble data type not supported by MIOpenGEMM.");
+            MIOPEN_THROW(miopenStatusBadParm, "miopenDouble data type not supported by rocBLAS.");
         }
         break;
         }
@@ -1157,8 +1155,7 @@ miopenStatus_t CallGemmStridedBatchedSequential(const Handle& handle,
         }
 
         case miopenDouble: {
-            MIOPEN_THROW(miopenStatusBadParm,
-                         "miopenDouble data type not supported by MIOpenGEMM.");
+            MIOPEN_THROW(miopenStatusBadParm, "miopenDouble data type not supported by rocBLAS.");
         }
         break;
         }

src/hip/handlehip.cpp

@@ -30,7 +30,6 @@
 #include <miopen/binary_cache.hpp>
 #include <miopen/env.hpp>
 #include <miopen/errors.hpp>
-#include <miopen/gemm_geometry.hpp>
 #include <miopen/handle_lock.hpp>
 #include <miopen/invoker.hpp>
 #include <miopen/kernel_cache.hpp>

src/hipoc/hipoc_program.cpp

@@ -328,8 +328,7 @@ void HIPOCProgramImpl::BuildCodeObject(std::string params,
     }
     else if(miopen::EndsWith(filename, ".cl"))
     {
-        params +=
-            " -Werror" + (is_kernel_str ? MiopengemmWarningsString() : OclKernelWarningsString());
+        params += " -Werror" + OclKernelWarningsString();
     }
 #else
     if(miopen::EndsWith(filename, ".cpp") || miopen::EndsWith(filename, ".cl"))

src/include/miopen/handle.hpp

@@ -66,10 +66,6 @@
 namespace miopen {
 
 struct HandleImpl;
-#if MIOPEN_USE_MIOPENGEMM
-struct GemmGeometry;
-using GemmKey = std::pair<std::string, std::string>;
-#endif
 
 #if MIOPEN_USE_ROCBLAS
 using rocblas_handle_ptr = MIOPEN_MANAGE_PTR(rocblas_handle, rocblas_destroy_handle);
@@ -225,9 +221,6 @@ struct Handle : miopenHandle
 
     std::unique_ptr<HandleImpl> impl;
     std::unordered_map<std::string, std::vector<miopenConvSolution_t>> find_map;
-#if MIOPEN_USE_MIOPENGEMM
-    std::unordered_map<GemmKey, std::unique_ptr<GemmGeometry>, SimpleHash> geo_map;
-#endif
 
     Invoker PrepareInvoker(const InvokerFactory& factory,
                            const std::vector<solver::KernelInfo>& kernels) const;

src/include/miopen/kernel_warnings.hpp

@@ -31,7 +31,6 @@
 
 namespace miopen {
 
-const std::string& MiopengemmWarningsString();
 const std::string& OclKernelWarningsString();
 const std::string& HipKernelWarningsString();
 

src/include/miopen/rocm_features.hpp

@@ -43,28 +43,6 @@
 /// To be removed as soon as support for ROCm 3.x is discontinued.
 #define WORKAROUND_MLOPEN_ISSUE_1711 (HIP_PACKAGE_VERSION_FLAT < 4000000000ULL)
 
-/// W/A for MIOpenGEMM issues with ROCm 4.1 and newer ROCm
-/// versions. The issue is highly likely related to the
-/// issues in the OpenCL compiler or in MIOpenGEMM itself.
-/// MIOpenGEMM is used only for OCL BE and deprecated.
-/// Related ticket: http://ontrack-internal.amd.com/browse/SWDEV-276757
-///
-/// Some failing cases:
-/// test_immed_conv2d --float --cmode conv --pmode default --group-count 1
-///  --input 1, 3, 224, 224 --weights 1, 3, 11, 11
-///   --pads_strides_dilations 1 1 1 1 1 1 --trans_output_pads 0 0
-///  --input 1, 3, 224, 224 --weights 1, 3, 7, 7
-///   --pads_strides_dilations 3 3 2 2 1 1 --trans_output_pads 0 0
-/// test_immed_conv3d --float --cmode conv --pmode default --group-count 1
-///  --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) 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))
-
 #define ROCM_FEATURE_TARGETID_OFF (HIP_PACKAGE_VERSION_FLAT < 4001000000ULL)
 
 /// Return type of llvm.amdgcn.buffer.atomic.fadd.f32 can't be detected.

src/include/miopen/solver.hpp

@@ -2830,7 +2830,7 @@ struct ConvDirectNaiveConvFwd final : ConvSolver
     bool IsApplicable(const ExecutionContext&, const ProblemDescription&) const override;
     bool IsDynamic() const override { return true; }
     /// Use very small fixed value enough to backup GEMM for cases when
-    /// GEMM is disabled due to MIOpenGemm or OCL compiler issues.
+    /// GEMM is disabled.
     float GetWti(const ExecutionContext&, const ProblemDescription&) const override
     {
         return 0.01f;
@@ -2848,7 +2848,7 @@ struct ConvDirectNaiveConvBwd final : ConvSolver
     bool IsApplicable(const ExecutionContext&, const ProblemDescription&) const override;
     bool IsDynamic() const override { return true; }
     /// Use very small fixed value enough to backup GEMM for cases when
-    /// GEMM is disabled due to MIOpenGemm or OCL compiler issues.
+    /// GEMM is disabled.
     float GetWti(const ExecutionContext&, const ProblemDescription&) const override
     {
         return 0.01f;
@@ -2866,7 +2866,7 @@ struct ConvDirectNaiveConvWrw final : ConvSolver
     bool IsApplicable(const ExecutionContext&, const ProblemDescription&) const override;
     bool IsDynamic() const override { return true; }
     /// Use very small fixed value enough to backup GEMM for cases when
-    /// GEMM is disabled due to MIOpenGemm or OCL compiler issues.
+    /// GEMM is disabled.
     float GetWti(const ExecutionContext&, const ProblemDescription&) const override
     {
         return 0.01f;
@@ -3114,8 +3114,6 @@ struct GemmBwd1x1_stride1 final : GemmBwdBase
 
 private:
     size_t GetWorkspaceSize(const ExecutionContext&, const conv::ProblemDescription&) const;
-    bool IsApplicableBeforeWorkaround(const ExecutionContext&,
-                                      const conv::ProblemDescription&) const;
     bool IsApplicable(const ExecutionContext&, const conv::ProblemDescription&) const;
     ConvSolution GetSolution(const ExecutionContext&, const conv::ProblemDescription&) const;