diff --git a/Dockerfile b/Dockerfile
index 4b2a38de5f..8ebaa17969 100755
--- a/Dockerfile
+++ b/Dockerfile
@@ -111,13 +111,4 @@ RUN pip3 install -r /doc-requirements.txt
 # Composable Kernel requires this version cmake
 RUN pip3 install --upgrade cmake==3.27.5
 
-# Use parallel job to accelerate tensile build
-# Workaround for Tensile with TargetID feature
-ARG USE_TARGETID="OFF"
-RUN if [ "$USE_TARGETID" = "ON" ] ; then export HIPCC_LINK_FLAGS_APPEND='-O3 -parallel-jobs=4' && export HIPCC_COMPILE_FLAGS_APPEND='-O3 -Wno-format-nonliteral -parallel-jobs=4' && rm -f /usr/bin/hipcc; fi
-
-# install last released miopentensile in default (master), install latest commits when MIOTENSILE_VER="latest" (develop)
-ARG MIOTENSILE_VER="default"
-RUN if [ "$USE_TARGETID" = "OFF" ] ; then echo "MIOpenTensile is not installed."; elif [ "$MIOTENSILE_VER" = "latest" ] ; then cget -p $PREFIX install ROCmSoftwarePlatform/MIOpenTensile@94a9047741d16a8eccd290131b78fb1aa69cdcdf; else cget -p $PREFIX install ROCmSoftwarePlatform/MIOpenTensile@94a9047741d16a8eccd290131b78fb1aa69cdcdf; fi
-
 RUN groupadd -f render
diff --git a/Jenkinsfile b/Jenkinsfile
index 7e07d3deba..0b4cfa2205 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -192,10 +192,8 @@ def getDockerImage(Map conf=[:])
     env.DOCKER_BUILDKIT=1
     def prefixpath = conf.get("prefixpath", "/opt/rocm") // one image for each prefix 1: /usr/local 2:/opt/rocm
     def gpu_arch = "gfx900;gfx906;gfx908;gfx90a;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102" // prebuilt dockers should have all the architectures enabled so one image can be used for all stages
-    def miotensile_version = conf.get("miotensile_version", "default") // deprecated
-    def target_id = conf.get("target_id", "OFF") // deprecated
     def mlir_build = conf.get("mlir_build", "ON") // always ON
-    def dockerArgs = "--build-arg BUILDKIT_INLINE_CACHE=1 --build-arg PREFIX=${prefixpath} --build-arg GPU_ARCH='${gpu_arch}' --build-arg MIOTENSILE_VER='${miotensile_version}' --build-arg USE_TARGETID='${target_id}' --build-arg USE_MLIR='${mlir_build}' "
+    def dockerArgs = "--build-arg BUILDKIT_INLINE_CACHE=1 --build-arg PREFIX=${prefixpath} --build-arg GPU_ARCH='${gpu_arch}' --build-arg USE_MLIR='${mlir_build}' "
     if(env.CCACHE_HOST)
     {
         def check_host = sh(script:"""(printf "PING\r\n";) | nc -N ${env.CCACHE_HOST} 6379 """, returnStdout: true).trim()
diff --git a/README.md b/README.md
index 440cd7bbc4..085ead7ce1 100755
--- a/README.md
+++ b/README.md
@@ -33,9 +33,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/)
@@ -43,7 +40,6 @@ python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
   * Version 1.79 is recommended, older version may need patches to work on newer systems, e.g. boost1{69,70,72} w/glibc-2.34
 * [SQLite3](https://sqlite.org/index.html) - reading and writing performance database
 * lbzip2 - multi-threaded compress or decompress utility
-* [MIOpenTENSILE](https://github.com/ROCmSoftwarePlatform/MIOpenTensile) - users can enable this library using the cmake configuration flag`-DMIOPEN_USE_MIOPENTENSILE=On`. (deprecated after ROCm 5.1.1)
 * [rocBLAS](https://github.com/ROCm/rocBLAS) - AMD library for Basic Linear Algebra Subprograms (BLAS) on the ROCm platform.
   * Minimum version branch for pre-ROCm 3.5 [master-rocm-2.10](https://github.com/ROCm/rocBLAS/tree/master-rocm-2.10)
   * Minimum version branch for post-ROCm 3.5 [master-rocm-3.5](https://github.com/ROCm/rocBLAS/releases/tag/rocm-3.5.0)
@@ -106,8 +102,6 @@ This prefix can used to specify the dependency path during the configuration pha
 
 * MIOpen's HIP backend uses [rocBLAS](https://github.com/ROCm/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
diff --git a/docs/DebugAndLogging.md b/docs/DebugAndLogging.md
index ffa30cbed8..8996580208 100644
--- a/docs/DebugAndLogging.md
+++ b/docs/DebugAndLogging.md
@@ -168,8 +168,7 @@ The `ROCBLAS_LAYER` environmental variable can be set to output GEMM information
 * `ROCBLAS_LAYER=2` - is set to 2, then there is bench logging
 * `ROCBLAS_LAYER=3` - is set to 3, then there is both trace and bench logging
 
-Additionally, using environment variable "MIOPEN_GEMM_ENFORCE_BACKEND", can override the default behavior. The default behavior which is to use
-both MIOpenGEMM and rocBlas depending on the input configuration:
+Additionally, the environment variable "MIOPEN_GEMM_ENFORCE_BACKEND" can be set to override default GEMM backend (Default GEMM backend is rocBLAS):
 
 * `MIOPEN_GEMM_ENFORCE_BACKEND=1`, use rocBLAS if enabled
 * `MIOPEN_GEMM_ENFORCE_BACKEND=2`, reserved
diff --git a/docs/install.md b/docs/install.md
index de8cd2d4e4..bf99c76b8f 100644
--- a/docs/install.md
+++ b/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/) 
@@ -72,5 +69,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/ROCm/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/ROCm/MIOpenGEMM) by default. Users can install MIOpenGEMM minimum release by using `apt-get install miopengemm`.
diff --git a/include/miopen/config.h.in b/include/miopen/config.h.in
index d87f5e105d..b95942e76e 100644
--- a/include/miopen/config.h.in
+++ b/include/miopen/config.h.in
@@ -29,8 +29,6 @@
 #cmakedefine01 MIOPEN_BACKEND_OPENCL
 #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 +84,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_ROCBLAS)
 
 // Usage of "defined" operator within macro expansion is undefined behavior,
 // so "defined(NDEBUG)" cannot be used there... unlike the following macro:
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 1d5548db7e..364b92d7ef 100644
--- a/src/CMakeLists.txt
+++ b/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/norm/forward_layernorm.cpp
     solver/norm/forward_layernorm2d_ck.cpp
diff --git a/src/anyramdb.cpp b/src/anyramdb.cpp
index 9095ed36b9..7985a4fdfe 100644
--- a/src/anyramdb.cpp
+++ b/src/anyramdb.cpp
@@ -28,7 +28,6 @@
 
 #include <miopen/errors.hpp>
 #include <miopen/logger.hpp>
-#include <miopen/md5.hpp>
 
 #include <boost/filesystem/operations.hpp>
 #include <boost/filesystem.hpp>
diff --git a/src/binary_cache.cpp b/src/binary_cache.cpp
index 37ef2c18df..a14b6c56fb 100644
--- a/src/binary_cache.cpp
+++ b/src/binary_cache.cpp
@@ -145,52 +145,37 @@ KDb GetDb(const TargetProperties& target, size_t num_cu)
 }
 #endif
 
-boost::filesystem::path GetCacheFile(const std::string& device,
-                                     const std::string& name,
-                                     const std::string& args,
-                                     bool is_kernel_str)
+boost::filesystem::path
+GetCacheFile(const std::string& device, const std::string& name, const std::string& args)
 {
-    const std::string filename = (is_kernel_str ? miopen::md5(name) : name) + ".o";
+    const std::string filename = name + ".o";
     return GetCachePath(false) / miopen::md5(device + ":" + args) / filename;
 }
 
 #if MIOPEN_ENABLE_SQLITE_KERN_CACHE
-static inline std::string GetFilenameForInfo2Logging(const bool is_kernel_str,
-                                                     const std::string& filename,
-                                                     const std::string& name)
-{
-    if(!miopen::IsLogging(miopen::LoggingLevel::Info2))
-        return {}; // Used only in MIOPEN_LOG_I2 -- optimize for speed.
-    if(is_kernel_str)
-        return filename + " size=" + std::to_string(name.size());
-    return filename;
-}
-
 std::string LoadBinary(const TargetProperties& target,
                        const size_t num_cu,
                        const std::string& name,
-                       const std::string& args,
-                       bool is_kernel_str)
+                       const std::string& args)
 {
     if(miopen::IsCacheDisabled())
         return {};
 
     auto db = GetDb(target, num_cu);
 
-    const std::string filename = (is_kernel_str ? miopen::md5(name) : name) + ".o";
+    const std::string filename = name + ".o";
     const KernelConfig cfg{filename, args, ""};
 
-    const auto verbose_name = GetFilenameForInfo2Logging(is_kernel_str, filename, name);
-    MIOPEN_LOG_I2("Loading binary for: " << verbose_name << "; args: " << args);
+    MIOPEN_LOG_I2("Loading binary for: " << filename << "; args: " << args);
     auto record = db.FindRecord(cfg);
     if(record)
     {
-        MIOPEN_LOG_I2("Successfully loaded binary for: " << verbose_name << "; args: " << args);
+        MIOPEN_LOG_I2("Successfully loaded binary for: " << filename << "; args: " << args);
         return record.get();
     }
     else
     {
-        MIOPEN_LOG_I2("Unable to load binary for: " << verbose_name << "; args: " << args);
+        MIOPEN_LOG_I2("Unable to load binary for: " << filename << "; args: " << args);
         return {};
     }
 }
@@ -199,33 +184,30 @@ void SaveBinary(const std::string& hsaco,
                 const TargetProperties& target,
                 const std::size_t num_cu,
                 const std::string& name,
-                const std::string& args,
-                bool is_kernel_str)
+                const std::string& args)
 {
     if(miopen::IsCacheDisabled())
         return;
 
     auto db = GetDb(target, num_cu);
 
-    const std::string filename = (is_kernel_str ? miopen::md5(name) : name) + ".o";
+    const std::string filename = name + ".o";
     KernelConfig cfg{filename, args, hsaco};
 
-    const auto verbose_name = GetFilenameForInfo2Logging(is_kernel_str, filename, name);
-    MIOPEN_LOG_I2("Saving binary for: " << verbose_name << "; args: " << args);
+    MIOPEN_LOG_I2("Saving binary for: " << filename << "; args: " << args);
     db.StoreRecord(cfg);
 }
 #else
 boost::filesystem::path LoadBinary(const TargetProperties& target,
                                    const size_t num_cu,
                                    const std::string& name,
-                                   const std::string& args,
-                                   bool is_kernel_str)
+                                   const std::string& args)
 {
     if(miopen::IsCacheDisabled())
         return {};
 
     (void)num_cu;
-    auto f = GetCacheFile(target.DbId(), name, args, is_kernel_str);
+    auto f = GetCacheFile(target.DbId(), name, args);
     if(boost::filesystem::exists(f))
     {
         return f.string();
@@ -239,8 +221,7 @@ boost::filesystem::path LoadBinary(const TargetProperties& target,
 void SaveBinary(const boost::filesystem::path& binary_path,
                 const TargetProperties& target,
                 const std::string& name,
-                const std::string& args,
-                bool is_kernel_str)
+                const std::string& args)
 {
     if(miopen::IsCacheDisabled())
     {
@@ -248,7 +229,7 @@ void SaveBinary(const boost::filesystem::path& binary_path,
     }
     else
     {
-        auto p = GetCacheFile(target.DbId(), name, args, is_kernel_str);
+        auto p = GetCacheFile(target.DbId(), name, args);
         boost::filesystem::create_directories(p.parent_path());
         boost::filesystem::rename(binary_path, p);
     }
diff --git a/src/db.cpp b/src/db.cpp
index a457178a70..e1ee671fa6 100644
--- a/src/db.cpp
+++ b/src/db.cpp
@@ -28,7 +28,6 @@
 #include <miopen/errors.hpp>
 #include <miopen/lock_file.hpp>
 #include <miopen/logger.hpp>
-#include <miopen/md5.hpp>
 
 #include <boost/date_time/posix_time/posix_time_types.hpp>
 #include <boost/filesystem.hpp>
diff --git a/src/gemm_v2.cpp b/src/gemm_v2.cpp
index fa1969bfb3..57b89c90c3 100644
--- a/src/gemm_v2.cpp
+++ b/src/gemm_v2.cpp
@@ -643,8 +643,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;
         }
@@ -918,8 +917,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;
         }
@@ -1191,8 +1189,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;
         }
diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp
index e071782463..70ddd93453 100644
--- a/src/hip/handlehip.cpp
+++ b/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>
@@ -441,10 +440,8 @@ KernelInvoke Handle::AddKernel(const std::string& algorithm,
                                const std::vector<size_t>& vgd,
                                const std::string& params,
                                std::size_t cache_index,
-                               bool is_kernel_str,
                                const std::string& kernel_src) const
 {
-
     auto obj = this->impl->cache.AddKernel(*this,
                                            algorithm,
                                            network_config,
@@ -454,7 +451,6 @@ KernelInvoke Handle::AddKernel(const std::string& algorithm,
                                            vgd,
                                            params,
                                            cache_index,
-                                           is_kernel_str,
                                            kernel_src);
     return this->Run(obj);
 }
@@ -502,7 +498,6 @@ KernelInvoke Handle::Run(Kernel k) const
 
 Program Handle::LoadProgram(const std::string& program_name,
                             std::string params,
-                            bool is_kernel_str,
                             const std::string& kernel_src) const
 {
     this->impl->set_ctx();
@@ -513,11 +508,8 @@ Program Handle::LoadProgram(const std::string& program_name,
     if(!miopen::EndsWith(program_name, ".mlir"))
         params = params + " -mcpu=" + this->GetTargetProperties().Name();
 
-    auto hsaco = miopen::LoadBinary(this->GetTargetProperties(),
-                                    this->GetMaxComputeUnits(),
-                                    program_name,
-                                    params,
-                                    is_kernel_str);
+    auto hsaco = miopen::LoadBinary(
+        this->GetTargetProperties(), this->GetMaxComputeUnits(), program_name, params);
     if(hsaco.empty())
     {
         const auto arch_target_id = miopen::SplitDelim(arch_name, ':');
@@ -528,8 +520,7 @@ Program Handle::LoadProgram(const std::string& program_name,
             hsaco                = miopen::LoadBinary(this->GetTargetProperties(),
                                        this->GetMaxComputeUnits(),
                                        program_name,
-                                       orig_params + " -mcpu=" + base_arch,
-                                       is_kernel_str);
+                                       orig_params + " -mcpu=" + base_arch);
         }
     }
 
@@ -538,9 +529,8 @@ Program Handle::LoadProgram(const std::string& program_name,
     if(hsaco.empty())
     {
         CompileTimer ct;
-        auto p = HIPOCProgram{
-            program_name, params, is_kernel_str, this->GetTargetProperties(), kernel_src};
-        ct.Log("Kernel", is_kernel_str ? std::string() : program_name);
+        auto p = HIPOCProgram{program_name, params, this->GetTargetProperties(), kernel_src};
+        ct.Log("Kernel", program_name);
 
 // Save to cache
 #if MIOPEN_ENABLE_SQLITE_KERN_CACHE
@@ -550,15 +540,14 @@ Program Handle::LoadProgram(const std::string& program_name,
                            this->GetTargetProperties(),
                            this->GetMaxComputeUnits(),
                            program_name,
-                           params,
-                           is_kernel_str);
+                           params);
 #else
         auto path = miopen::GetCachePath(false) / boost::filesystem::unique_path();
         if(p.IsCodeObjectInMemory())
             miopen::WriteFile(p.GetCodeObjectBlob(), path);
         else
             boost::filesystem::copy_file(p.GetCodeObjectPathname(), path);
-        miopen::SaveBinary(path, this->GetTargetProperties(), program_name, params, is_kernel_str);
+        miopen::SaveBinary(path, this->GetTargetProperties(), program_name, params);
 #endif
         p.FreeCodeObjectFileStorage();
         return p;
diff --git a/src/hipoc/hipoc_program.cpp b/src/hipoc/hipoc_program.cpp
index 4de0c6915b..d7892c36f5 100644
--- a/src/hipoc/hipoc_program.cpp
+++ b/src/hipoc/hipoc_program.cpp
@@ -200,12 +200,11 @@ HIPOCProgramImpl::HIPOCProgramImpl(const std::string& program_name, const std::s
 
 HIPOCProgramImpl::HIPOCProgramImpl(const std::string& program_name,
                                    std::string params,
-                                   bool is_kernel_str,
                                    const TargetProperties& target_,
                                    const std::string& kernel_src)
     : program(program_name), target(target_)
 {
-    BuildCodeObject(params, is_kernel_str, kernel_src);
+    BuildCodeObject(params, kernel_src);
     if(!binary.empty())
     {
         module = CreateModuleInMem(binary);
@@ -313,19 +312,14 @@ void HIPOCProgramImpl::BuildCodeObjectInMemory(const std::string& params,
 }
 #endif // MIOPEN_USE_COMGR
 
-void HIPOCProgramImpl::BuildCodeObject(std::string params,
-                                       bool is_kernel_str,
-                                       const std::string& kernel_src)
+void HIPOCProgramImpl::BuildCodeObject(std::string params, const std::string& kernel_src)
 {
-    std::string filename = is_kernel_str ? "tinygemm.cl" // Fixed name for miopengemm.
-                                         : program;
+    std::string filename = program;
     const auto src       = [&]() -> std::string {
         if(miopen::EndsWith(filename, ".mlir"))
             return {}; // MLIR solutions do not use source code.
         if(!kernel_src.empty())
             return kernel_src;
-        if(is_kernel_str)
-            return program;
         return GetKernelSrc(program);
     }();
 
@@ -336,8 +330,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"))
@@ -354,11 +347,9 @@ void HIPOCProgramImpl::BuildCodeObject(std::string params,
 HIPOCProgram::HIPOCProgram() {}
 HIPOCProgram::HIPOCProgram(const std::string& program_name,
                            std::string params,
-                           bool is_kernel_str,
                            const TargetProperties& target,
                            const std::string& kernel_src)
-    : impl(std::make_shared<HIPOCProgramImpl>(
-          program_name, params, is_kernel_str, target, kernel_src))
+    : impl(std::make_shared<HIPOCProgramImpl>(program_name, params, target, kernel_src))
 {
 }
 
diff --git a/src/include/miopen/binary_cache.hpp b/src/include/miopen/binary_cache.hpp
index 1f00c19502..cbca27ba12 100644
--- a/src/include/miopen/binary_cache.hpp
+++ b/src/include/miopen/binary_cache.hpp
@@ -32,14 +32,14 @@
 #include <boost/filesystem/path.hpp>
 #include <string>
 
+#define FIN_OLD_BINARY_CACHE_COMPAT 1
+
 namespace miopen {
 
 bool IsCacheDisabled();
 
-boost::filesystem::path GetCacheFile(const std::string& device,
-                                     const std::string& name,
-                                     const std::string& args,
-                                     bool is_kernel_str);
+boost::filesystem::path
+GetCacheFile(const std::string& device, const std::string& name, const std::string& args);
 
 boost::filesystem::path GetCachePath(bool is_system);
 
@@ -47,26 +47,44 @@ boost::filesystem::path GetCachePath(bool is_system);
 boost::filesystem::path LoadBinary(const TargetProperties& target,
                                    std::size_t num_cu,
                                    const std::string& name,
-                                   const std::string& args,
-                                   bool is_kernel_str = false);
+                                   const std::string& args);
 void SaveBinary(const boost::filesystem::path& binary_path,
                 const TargetProperties& target,
                 const std::string& name,
-                const std::string& args,
-                bool is_kernel_str = false);
+                const std::string& args);
 #else
 std::string LoadBinary(const TargetProperties& target,
                        std::size_t num_cu,
                        const std::string& name,
-                       const std::string& args,
-                       bool is_kernel_str = false);
+                       const std::string& args);
 
 void SaveBinary(const std::string& hsaco,
                 const TargetProperties& target,
                 std::size_t num_cu,
                 const std::string& name,
-                const std::string& args,
-                bool is_kernel_str = false);
+                const std::string& args);
+
+#if FIN_OLD_BINARY_CACHE_COMPAT
+inline std::string LoadBinary(const TargetProperties& target,
+                              std::size_t num_cu,
+                              const std::string& name,
+                              const std::string& args,
+                              bool)
+{
+    return LoadBinary(target, num_cu, name, args);
+}
+
+inline void SaveBinary(const std::string& hsaco,
+                       const TargetProperties& target,
+                       std::size_t num_cu,
+                       const std::string& name,
+                       const std::string& args,
+                       bool)
+{
+    SaveBinary(hsaco, target, num_cu, name, args);
+}
+#endif
+
 #endif
 
 } // namespace miopen
diff --git a/src/include/miopen/clhelper.hpp b/src/include/miopen/clhelper.hpp
index f15caf9f5b..648bb383af 100644
--- a/src/include/miopen/clhelper.hpp
+++ b/src/include/miopen/clhelper.hpp
@@ -45,7 +45,6 @@ ClProgramPtr LoadProgram(cl_context ctx,
                          const TargetProperties& target,
                          const std::string& program,
                          std::string params,
-                         bool is_kernel_str,
                          const std::string& kernel_src);
 void GetProgramBinary(const ClProgramPtr& program, std::string& binary);
 void SaveProgramBinary(const ClProgramPtr& program, const std::string& name);
diff --git a/src/include/miopen/gemm_geometry.hpp b/src/include/miopen/gemm_geometry.hpp
deleted file mode 100644
index f95aef2ac2..0000000000
--- a/src/include/miopen/gemm_geometry.hpp
+++ /dev/null
@@ -1,88 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2017 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_MIOPEN_GEMM_GEOMETRY_HPP_
-#define GUARD_MIOPEN_GEMM_GEOMETRY_HPP_
-
-#include <miopen/config.h>
-#include <miopen/kernel_cache.hpp>
-#include <miopen/tensor.hpp>
-
-#if MIOPEN_USE_MIOPENGEMM
-#include <miopengemm/miogemm.hpp>
-
-#include <mutex>
-
-namespace miopen {
-
-struct GemmGeometry
-{
-    std::string algorithm_name;
-    float alpha{};
-    float beta{};
-    MIOpenGEMM::Geometry tgg{};
-    bool beta_kern_req{};
-
-    /* jn : if miopengemm returned a beta kernel.
-     * not the same as beta_kern_req(uired), as
-     * if beta == 1, beta kernel is returned but
-     * not required.
-     * we still need to know if it was returned,
-     * as the function signature of the main kernel
-     * is then different.
-     * */
-    bool beta_kern_returned{};
-    std::array<int, 2> beta_kern_args = {{0, 0}};
-
-    GemmGeometry() {}
-    GemmGeometry(std::string algo_name, float palpha, float pbeta, MIOpenGEMM::Geometry ptgg)
-        : algorithm_name(algo_name), alpha(palpha), beta(pbeta), tgg(ptgg)
-    {
-        beta_kern_req      = false;
-        beta_kern_returned = false;
-    }
-
-    void EnableBetaKernel(bool enable);
-
-    void FindSolution(float time,
-                      Handle& handle,
-                      ConstData_t a,
-                      ConstData_t b,
-                      Data_t c,
-                      bool enforce_determinism);
-
-    void RunGemm(const Handle& handle,
-                 ConstData_t a,
-                 ConstData_t b,
-                 Data_t c,
-                 int a_offset,
-                 int b_offset,
-                 int c_offset);
-};
-
-} // namespace miopen
-#endif // MIOPEN_USE_MIOPENGEMM
-
-#endif // GUARD_MIOPEN_GEMM_GEOMETRY_HPP_
diff --git a/src/include/miopen/generic_search.hpp b/src/include/miopen/generic_search.hpp
index 00ccd95866..c3aeea319f 100644
--- a/src/include/miopen/generic_search.hpp
+++ b/src/include/miopen/generic_search.hpp
@@ -351,7 +351,7 @@ void CompileAgent(size_t thread_index,
         {
             if(profile_h.HasProgram(kernel.kernel_file, kernel.comp_options))
                 continue;
-            std::ignore = profile_h.LoadProgram(kernel.kernel_file, kernel.comp_options, false, "");
+            std::ignore = profile_h.LoadProgram(kernel.kernel_file, kernel.comp_options, "");
         }
         auto tup = std::make_tuple<PerformanceConfig, ConvSolution, bool>(
             std::move(current_config), std::move(current_solution), false);
diff --git a/src/include/miopen/handle.hpp b/src/include/miopen/handle.hpp
index 7d1bb79a37..5e8cdfac86 100644
--- a/src/include/miopen/handle.hpp
+++ b/src/include/miopen/handle.hpp
@@ -63,13 +63,11 @@
 #endif
 #endif
 
+#define FIN_OLD_HANDLE_COMPAT 1
+
 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);
@@ -109,7 +107,6 @@ struct Handle : miopenHandle
                            const std::vector<size_t>& vgd,
                            const std::string& params,
                            std::size_t cache_index       = 0,
-                           bool is_kernel_str            = false,
                            const std::string& kernel_src = "") const;
 
     void ClearKernels(const std::string& algorithm, const std::string& network_config) const;
@@ -136,9 +133,18 @@ struct Handle : miopenHandle
 
     Program LoadProgram(const std::string& program_name,
                         std::string params,
-                        bool is_kernel_str,
                         const std::string& kernel_src) const;
 
+#if FIN_OLD_HANDLE_COMPAT
+    Program LoadProgram(const std::string& program_name,
+                        std::string params,
+                        bool,
+                        const std::string& kernel_src) const
+    {
+        return LoadProgram(program_name, params, kernel_src);
+    }
+#endif
+
     bool HasProgram(const std::string& program_name, const std::string& params) const;
     void ClearProgram(const std::string& program_name, const std::string& params) const;
     void AddProgram(Program prog, const std::string& program_name, const std::string& params) const;
@@ -225,9 +231,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;
diff --git a/src/include/miopen/hipoc_program.hpp b/src/include/miopen/hipoc_program.hpp
index 0e569fc80b..c6be6544f0 100644
--- a/src/include/miopen/hipoc_program.hpp
+++ b/src/include/miopen/hipoc_program.hpp
@@ -47,7 +47,6 @@ struct HIPOCProgram
     /// Other ctors only guarantee to initialize module.
     HIPOCProgram(const std::string& program_name,
                  std::string params,
-                 bool is_kernel_str,
                  const TargetProperties& target,
                  const std::string& kernel_src);
     HIPOCProgram(const std::string& program_name, const boost::filesystem::path& hsaco);
diff --git a/src/include/miopen/hipoc_program_impl.hpp b/src/include/miopen/hipoc_program_impl.hpp
index d49c7bf306..f9067dc0ce 100644
--- a/src/include/miopen/hipoc_program_impl.hpp
+++ b/src/include/miopen/hipoc_program_impl.hpp
@@ -45,7 +45,6 @@ struct HIPOCProgramImpl
 
     HIPOCProgramImpl(const std::string& program_name,
                      std::string params,
-                     bool is_kernel_str,
                      const TargetProperties& target_,
                      const std::string& kernel_src);
 
@@ -65,7 +64,7 @@ struct HIPOCProgramImpl
                                  const std::string& filename);
 #endif
 
-    void BuildCodeObject(std::string params, bool is_kernel_str, const std::string& kernel_src);
+    void BuildCodeObject(std::string params, const std::string& kernel_src);
 };
 } // namespace miopen
 #endif // GUARD_MIOPEN_HIPOC_PROGRAM_IMPL_HPP
diff --git a/src/include/miopen/kernel_cache.hpp b/src/include/miopen/kernel_cache.hpp
index 1f9e631809..3b2909b27c 100644
--- a/src/include/miopen/kernel_cache.hpp
+++ b/src/include/miopen/kernel_cache.hpp
@@ -73,7 +73,6 @@ class KernelCache
                      const std::vector<size_t>& vgd,
                      std::string params            = "",
                      std::size_t cache_index       = 0,
-                     bool is_kernel_miopengemm_str = false,
                      const std::string& kernel_src = "");
 
     void AddKernel(Key key, Kernel k, std::size_t cache_index);
diff --git a/src/include/miopen/kernel_warnings.hpp b/src/include/miopen/kernel_warnings.hpp
index e913894470..399b004cda 100644
--- a/src/include/miopen/kernel_warnings.hpp
+++ b/src/include/miopen/kernel_warnings.hpp
@@ -31,7 +31,6 @@
 
 namespace miopen {
 
-const std::string& MiopengemmWarningsString();
 const std::string& OclKernelWarningsString();
 const std::string& HipKernelWarningsString();
 
diff --git a/src/include/miopen/miopengemm.hpp b/src/include/miopen/miopengemm.hpp
deleted file mode 100644
index 33e3042bb0..0000000000
--- a/src/include/miopen/miopengemm.hpp
+++ /dev/null
@@ -1,62 +0,0 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2017 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_MIOPEN_MIOPENGEMM_HPP_
-#define GUARD_MIOPEN_MIOPENGEMM_HPP_
-
-#include <miopen/config.h>
-
-#if MIOPEN_USE_MIOPENGEMM
-#include <miopengemm/miogemm.hpp>
-
-namespace miopen {
-
-struct Handle;
-
-void AddMiopengemmSolution(const Handle& handle,
-                           const std::string& algorithm_name,
-                           const std::string& network_config,
-                           const MIOpenGEMM::Geometry& mgg,
-                           ConstData_t A,
-                           ConstData_t B,
-                           Data_t C,
-                           float time,
-                           bool enforce_determinism);
-
-void RunMiopengemmSolution(const Handle& handle,
-                           const decltype(handle.GetKernels("_", "_"))& kernels,
-                           float alpha,
-                           ConstData_t A,
-                           int a_offset,
-                           ConstData_t B,
-                           int b_offset,
-                           float beta,
-                           Data_t C,
-                           int c_offset);
-
-} // namespace miopen
-#endif // MIOPEN_USE_MIOPENGEMM
-
-#endif // GUARD_MIOPEN_MIOPENGEMM_HPP_
diff --git a/src/include/miopen/rocm_features.hpp b/src/include/miopen/rocm_features.hpp
index f8f2b754da..ab9d5033a1 100644
--- a/src/include/miopen/rocm_features.hpp
+++ b/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.
diff --git a/src/include/miopen/solver.hpp b/src/include/miopen/solver.hpp
index 972cd2906a..2b334499ac 100644
--- a/src/include/miopen/solver.hpp
+++ b/src/include/miopen/solver.hpp
@@ -2910,7 +2910,7 @@ struct ConvDirectNaiveConvFwd final : ConvSolver
                       const miopen::conv::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 miopen::conv::ProblemDescription&) const override
     {
         return 0.01f;
@@ -2930,7 +2930,7 @@ struct ConvDirectNaiveConvBwd final : ConvSolver
                       const miopen::conv::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 miopen::conv::ProblemDescription&) const override
     {
         return 0.01f;
@@ -2950,7 +2950,7 @@ struct ConvDirectNaiveConvWrw final : ConvSolver
                       const miopen::conv::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 miopen::conv::ProblemDescription&) const override
     {
         return 0.01f;
@@ -3091,10 +3091,6 @@ struct GemmBwd1x1_stride1 final : GemmBwdBase
     ConvSolution GetSolution(const ExecutionContext&,
                              const miopen::conv::ProblemDescription& problem) const override;
 
-private:
-    bool IsApplicableBeforeWorkaround(const ExecutionContext&,
-                                      const miopen::conv::ProblemDescription&) const;
-
     friend struct GemmBwdRest;
 };
 
diff --git a/src/include/miopen/solver/gemm_common.hpp b/src/include/miopen/solver/gemm_common.hpp
deleted file mode 100644
index 14ab5b8444..0000000000
--- a/src/include/miopen/solver/gemm_common.hpp
+++ /dev/null
@@ -1,44 +0,0 @@
-/*******************************************************************************
- *
- * 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 solver {
-namespace conv {
-namespace gemm {
-
-bool IsWorkaroundIssue1315(const miopen::ExecutionContext& ctx);
-
-} // namespace gemm
-} // namespace conv
-} // namespace solver
-} // namespace miopen
-
-#endif
diff --git a/src/kernel_cache.cpp b/src/kernel_cache.cpp
index 88e119bdcd..1c02c7d802 100644
--- a/src/kernel_cache.cpp
+++ b/src/kernel_cache.cpp
@@ -100,7 +100,6 @@ Kernel KernelCache::AddKernel(const Handle& h,
                               const std::vector<size_t>& vgd,
                               std::string params,
                               std::size_t cache_index,
-                              bool is_kernel_miopengemm_str,
                               const std::string& kernel_src)
 {
     const std::pair<std::string, std::string> key = std::make_pair(algorithm, network_config);
@@ -116,12 +115,7 @@ Kernel KernelCache::AddKernel(const Handle& h,
     }
     else
     {
-        if(!is_kernel_miopengemm_str) // default value
-        {
-            is_kernel_miopengemm_str = algorithm.find("ImplicitGEMM") == std::string::npos &&
-                                       algorithm.find("GEMM") != std::string::npos;
-        }
-        program = h.LoadProgram(program_name, params, is_kernel_miopengemm_str, kernel_src);
+        program = h.LoadProgram(program_name, params, kernel_src);
         program_map[std::make_pair(program_name, params)] = program;
     }
 
diff --git a/src/kernel_warnings.cpp b/src/kernel_warnings.cpp
index 253cdb9c17..5a9c67b37b 100644
--- a/src/kernel_warnings.cpp
+++ b/src/kernel_warnings.cpp
@@ -32,7 +32,7 @@
 
 namespace miopen {
 
-static std::vector<std::string> OclKernelWarnings(const bool is_miopengemm)
+static std::vector<std::string> OclKernelWarnings()
 {
     std::vector<std::string> rv = {
         "-Weverything",
@@ -52,9 +52,7 @@ static std::vector<std::string> OclKernelWarnings(const bool is_miopengemm)
         "-Wno-unused-macros",
         "-Wno-declaration-after-statement", // W/A for SWDEV-337356
     };
-    // W/A for SWDEV-270602. We'll remove this when we stop using MIOpenGEMM (deprecated).
-    if(is_miopengemm)
-        rv.emplace_back("-Wno-tautological-unsigned-zero-compare");
+
     return rv;
 }
 
@@ -94,18 +92,6 @@ static std::string MakeKernelWarningsString(const std::vector<std::string>& kern
     return prefix + JoinStrings(kernel_warnings, prefix);
 }
 
-const std::string& MiopengemmWarningsString()
-{
-#if MIOPEN_BACKEND_OPENCL
-    const std::string prefix = " -Wf,";
-#else
-    const std::string prefix = " ";
-#endif
-
-    static const std::string result = MakeKernelWarningsString(OclKernelWarnings(true), prefix);
-    return result;
-}
-
 const std::string& OclKernelWarningsString()
 {
 #if MIOPEN_BACKEND_OPENCL
@@ -114,7 +100,7 @@ const std::string& OclKernelWarningsString()
     const std::string prefix = " ";
 #endif
 
-    static const std::string result = MakeKernelWarningsString(OclKernelWarnings(false), prefix);
+    static const std::string result = MakeKernelWarningsString(OclKernelWarnings(), prefix);
     return result;
 }
 
diff --git a/src/nogpu/handle.cpp b/src/nogpu/handle.cpp
index 9536b97e09..d601347fbd 100644
--- a/src/nogpu/handle.cpp
+++ b/src/nogpu/handle.cpp
@@ -30,7 +30,6 @@
 #include <miopen/binary_cache.hpp>
 #include <miopen/target_properties.hpp>
 #include <miopen/errors.hpp>
-#include <miopen/gemm_geometry.hpp>
 #include <miopen/handle_lock.hpp>
 #include <miopen/invoker.hpp>
 #include <miopen/kernel_cache.hpp>
@@ -43,9 +42,7 @@
 #endif
 
 #include <boost/filesystem.hpp>
-#include <miopen/handle_lock.hpp>
 #include <miopen/load_file.hpp>
-#include <miopen/gemm_geometry.hpp>
 
 #ifndef _WIN32
 #include <unistd.h>
@@ -56,6 +53,7 @@
 #include <chrono>
 #include <thread>
 #include <miopen/nogpu/handle_impl.hpp>
+
 namespace miopen {
 
 Handle::Handle(miopenAcceleratorQueue_t /* stream */) : Handle::Handle() {}
@@ -111,7 +109,6 @@ KernelInvoke Handle::AddKernel(const std::string& algorithm,
                                const std::vector<size_t>& vgd,
                                const std::string& params,
                                std::size_t cache_index,
-                               bool is_kernel_str,
                                const std::string& kernel_src) const
 {
     auto obj = this->impl->cache.AddKernel(*this,
@@ -123,7 +120,6 @@ KernelInvoke Handle::AddKernel(const std::string& algorithm,
                                            vgd,
                                            params,
                                            cache_index,
-                                           is_kernel_str,
                                            kernel_src);
     return this->Run(obj);
 }
@@ -168,7 +164,6 @@ KernelInvoke Handle::Run(Kernel /* k */) const { return {}; }
 
 Program Handle::LoadProgram(const std::string& program_name,
                             std::string params,
-                            bool is_kernel_str,
                             const std::string& kernel_src) const
 {
     if(!miopen::EndsWith(program_name, ".mlir"))
@@ -176,11 +171,8 @@ Program Handle::LoadProgram(const std::string& program_name,
         params += " -mcpu=" + this->GetTargetProperties().Name();
     }
 
-    auto hsaco       = miopen::LoadBinary(this->GetTargetProperties(),
-                                    this->GetMaxComputeUnits(),
-                                    program_name,
-                                    params,
-                                    is_kernel_str);
+    auto hsaco = miopen::LoadBinary(
+        this->GetTargetProperties(), this->GetMaxComputeUnits(), program_name, params);
     auto pgmImpl     = std::make_shared<HIPOCProgramImpl>();
     pgmImpl->program = program_name;
     pgmImpl->target  = this->GetTargetProperties();
@@ -189,9 +181,8 @@ Program Handle::LoadProgram(const std::string& program_name,
     if(hsaco.empty())
     {
         // avoid the constructor since it implicitly calls the HIP API
-        pgmImpl->BuildCodeObject(params, is_kernel_str, kernel_src);
-// auto p = HIPOCProgram{
-//     program_name, params, is_kernel_str, this->GetTargetProperties(), kernel_src};
+        pgmImpl->BuildCodeObject(params, kernel_src);
+// auto p = HIPOCProgram{program_name, params, this->GetTargetProperties(), kernel_src};
 
 // Save to cache
 #if MIOPEN_ENABLE_SQLITE_KERN_CACHE
@@ -201,15 +192,14 @@ Program Handle::LoadProgram(const std::string& program_name,
                            this->GetTargetProperties(),
                            this->GetMaxComputeUnits(),
                            program_name,
-                           params,
-                           is_kernel_str);
+                           params);
 #else
         auto path = miopen::GetCachePath(false) / boost::filesystem::unique_path();
         if(p.IsCodeObjectInMemory())
             miopen::WriteFile(p.GetCodeObjectBlob(), path);
         else
             boost::filesystem::copy_file(p.GetCodeObjectPathname(), path);
-        miopen::SaveBinary(path, this->GetTargetProperties(), program_name, params, is_kernel_str);
+        miopen::SaveBinary(path, this->GetTargetProperties(), program_name, params);
 #endif
     }
     else
diff --git a/src/ocl/clhelper.cpp b/src/ocl/clhelper.cpp
index 40fea57b15..f25cc9fd3d 100644
--- a/src/ocl/clhelper.cpp
+++ b/src/ocl/clhelper.cpp
@@ -154,26 +154,17 @@ ClProgramPtr LoadProgram(cl_context ctx,
                          const TargetProperties& target,
                          const std::string& program,
                          std::string params,
-                         bool is_kernel_str,
                          const std::string& kernel_src)
 {
     std::string source;
     std::string program_name;
 
-    if(is_kernel_str)
-    {
-        source       = program;
-        program_name = "(unknown)";
-    }
+    program_name = program;
+    // For mlir build, leave both source and kernel_src to be empty
+    if((kernel_src.empty()) && !(miopen::EndsWith(program_name, ".mlir")))
+        source = miopen::GetKernelSrc(program_name);
     else
-    {
-        program_name = program;
-        // For mlir build, leave both source and kernel_src to be empty
-        if((kernel_src.empty()) && !(miopen::EndsWith(program_name, ".mlir")))
-            source = miopen::GetKernelSrc(program_name);
-        else
-            source = kernel_src;
-    }
+        source = kernel_src;
 
     bool load_binary = false;
     if(miopen::EndsWith(program_name, ".s"))
@@ -216,7 +207,7 @@ ClProgramPtr LoadProgram(cl_context ctx,
 #if MIOPEN_BUILD_DEV
         params += " -Werror";
 #ifdef __linux__
-        params += is_kernel_str ? MiopengemmWarningsString() : OclKernelWarningsString();
+        params += OclKernelWarningsString();
 #endif
 #endif
         params += " -cl-std=CL1.2";
diff --git a/src/ocl/handleocl.cpp b/src/ocl/handleocl.cpp
index ce5c680c26..8829fc790f 100644
--- a/src/ocl/handleocl.cpp
+++ b/src/ocl/handleocl.cpp
@@ -39,10 +39,6 @@
 #include <miopen/ocldeviceinfo.hpp>
 #include <miopen/timer.hpp>
 
-#if MIOPEN_USE_MIOPENGEMM
-#include <miopen/gemm_geometry.hpp>
-#endif
-
 #include <boost/filesystem.hpp>
 
 #include <string>
@@ -324,10 +320,8 @@ KernelInvoke Handle::AddKernel(const std::string& algorithm,
                                const std::vector<size_t>& vgd,
                                const std::string& params,
                                std::size_t cache_index,
-                               bool is_kernel_str,
                                const std::string& kernel_src) const
 {
-
     auto obj = this->impl->cache.AddKernel(*this,
                                            algorithm,
                                            network_config,
@@ -337,7 +331,6 @@ KernelInvoke Handle::AddKernel(const std::string& algorithm,
                                            vgd,
                                            params,
                                            cache_index,
-                                           is_kernel_str,
                                            kernel_src);
     return this->Run(obj);
 }
@@ -393,14 +386,10 @@ KernelInvoke Handle::Run(Kernel k) const
 
 Program Handle::LoadProgram(const std::string& program_name,
                             std::string params,
-                            bool is_kernel_str,
                             const std::string& kernel_src) const
 {
-    auto hsaco = miopen::LoadBinary(this->GetTargetProperties(),
-                                    this->GetMaxComputeUnits(),
-                                    program_name,
-                                    params,
-                                    is_kernel_str);
+    auto hsaco = miopen::LoadBinary(
+        this->GetTargetProperties(), this->GetMaxComputeUnits(), program_name, params);
     if(hsaco.empty())
     {
         CompileTimer ct;
@@ -409,25 +398,19 @@ Program Handle::LoadProgram(const std::string& program_name,
                                      this->GetTargetProperties(),
                                      program_name,
                                      params,
-                                     is_kernel_str,
                                      kernel_src);
-        ct.Log("Kernel", is_kernel_str ? std::string() : program_name);
+        ct.Log("Kernel", program_name);
 
 // Save to cache
 #if MIOPEN_ENABLE_SQLITE_KERN_CACHE
         std::string binary;
         miopen::GetProgramBinary(p, binary);
-        miopen::SaveBinary(binary,
-                           this->GetTargetProperties(),
-                           this->GetMaxComputeUnits(),
-                           program_name,
-                           params,
-                           is_kernel_str);
+        miopen::SaveBinary(
+            binary, this->GetTargetProperties(), this->GetMaxComputeUnits(), program_name, params);
 #else
         auto path = miopen::GetCachePath(false) / boost::filesystem::unique_path();
         miopen::SaveProgramBinary(p, path.string());
-        miopen::SaveBinary(
-            path.string(), this->GetTargetProperties(), program_name, params, is_kernel_str);
+        miopen::SaveBinary(path.string(), this->GetTargetProperties(), program_name, params);
 #endif
         return p;
     }
diff --git a/src/ramdb.cpp b/src/ramdb.cpp
index e24abaca47..7c72d77fef 100644
--- a/src/ramdb.cpp
+++ b/src/ramdb.cpp
@@ -29,7 +29,6 @@
 #include <miopen/errors.hpp>
 #include <miopen/lock_file.hpp>
 #include <miopen/logger.hpp>
-#include <miopen/md5.hpp>
 
 #include <boost/filesystem/operations.hpp>
 #include <boost/filesystem.hpp>
diff --git a/src/solver.cpp b/src/solver.cpp
index b4d25652ba..65a9bb4650 100644
--- a/src/solver.cpp
+++ b/src/solver.cpp
@@ -71,7 +71,7 @@ std::vector<Program> PrecompileKernels(const Handle& h, const std::vector<Kernel
                     max_threads{GetTuningThreadsMax()},
                     [&](auto i) {
                         const KernelInfo& k = kernels[i];
-                        programs[i]         = h.LoadProgram(k.kernel_file, k.comp_options, false, "");
+                        programs[i]         = h.LoadProgram(k.kernel_file, k.comp_options, "");
                     });
     // clang-format on
     ct.Log("PrecompileKernels");
diff --git a/src/solver/gemm.cpp b/src/solver/gemm.cpp
index d44ae46987..9f8fecfb7d 100644
--- a/src/solver/gemm.cpp
+++ b/src/solver/gemm.cpp
@@ -31,8 +31,6 @@
 #include <miopen/gemm_v2.hpp>
 #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>
@@ -43,11 +41,6 @@
 
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_CONV_PRECISE_ROCBLAS_TIMING)
 
-/// MIOpenGEMM issues with ROCm 3.7, most likely related to the
-/// issues in the OpenCL compiler. Not reproducible in ROCm 4.0.
-#define WORKAROUND_MIOPENGEMM_ROCM37 \
-    (MIOPEN_USE_MIOPENGEMM && HIP_PACKAGE_VERSION_MAJOR == 3 && HIP_PACKAGE_VERSION_MINOR == 7)
-
 namespace miopen {
 namespace solver {
 namespace conv {
@@ -84,8 +77,6 @@ static inline bool IsAnyBufferFp16(const TensorDescriptor& xDesc,
 bool GemmFwdBase::IsApplicable(const ExecutionContext& ctx, const ProblemDescription& problem) const
 {
 #if MIOPEN_USE_GEMM
-    if(conv::gemm::IsWorkaroundIssue1315(ctx))
-        return false;
     const auto& xDesc = problem.GetIn();
     const auto& wDesc = problem.GetWeights();
     const auto& yDesc = problem.GetOut();
@@ -1025,43 +1016,6 @@ bool GemmFwdRest::IsApplicable(const ExecutionContext& context,
     if(!GemmFwdBase::IsApplicable(context, problem))
         return false;
 
-#if WORKAROUND_MIOPENGEMM_ROCM37
-    {
-        decltype(auto) conv  = problem.GetConv();
-        decltype(auto) xDesc = problem.GetIn();
-        decltype(auto) wDesc = problem.GetWeights();
-
-        const auto spatial_dim  = conv.GetSpatialDimension();
-        const auto& in_spatial  = boost::adaptors::slice(xDesc.GetLengths(), 2, 2 + spatial_dim);
-        const auto& wei_spatial = boost::adaptors::slice(wDesc.GetLengths(), 2, 2 + spatial_dim);
-
-        const auto in_c = xDesc.GetLengths()[1];
-
-        if(conv.GetSpatialDimension() == 2 && conv.group_count == 4 && in_c == 4 &&
-           in_spatial[0] == 161 && in_spatial[1] == 700 && wDesc.GetLengths()[0] == 32 &&
-           wDesc.GetLengths()[1] == 1 && wei_spatial[0] == 5 && wei_spatial[1] == 20 &&
-           miopen::all_of(conv.GetConvPads(), [](auto v) { return v == 0; }) &&
-           miopen::all_of(conv.GetConvStrides(), [](auto v) { return v == 2; }) &&
-           miopen::all_of(conv.GetConvDilations(), [](auto v) { return v == 1; }))
-            return false;
-    }
-#endif
-#if WORKAROUND_MIOPENGEMM_SINCE_ROCM41
-    {
-        decltype(auto) conv  = problem.GetConv();
-        decltype(auto) xDesc = problem.GetIn();
-        decltype(auto) wDesc = problem.GetWeights();
-
-        const std::size_t spatial_dim = conv.GetSpatialDimension();
-        const auto in_spatial  = boost::adaptors::slice(xDesc.GetLengths(), 2, 2 + spatial_dim);
-        const auto wei_spatial = boost::adaptors::slice(wDesc.GetLengths(), 2, 2 + spatial_dim);
-
-        if(miopen::any_of(in_spatial, [](auto v) { return v >= 161; }) &&
-           miopen::any_of(wei_spatial, [](auto v) { return v >= 7; }))
-            return false;
-    }
-#endif
-
     // Todo: This is a rest-of kind of logic. Should be revised later.
     if(GemmFwd1x1_0_1{}.IsApplicable(context, problem))
         return false;
diff --git a/src/solver/gemm_bwd.cpp b/src/solver/gemm_bwd.cpp
index e3bbce86c0..eb196ebae4 100644
--- a/src/solver/gemm_bwd.cpp
+++ b/src/solver/gemm_bwd.cpp
@@ -31,7 +31,6 @@
 #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>
@@ -42,8 +41,6 @@
 
 MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_CONV_PRECISE_ROCBLAS_TIMING)
 
-#define WORKAROUND_MIOPENGEMM_ISSUE_59 1
-
 // copy from convolution.cpp
 // Workaround for issue 1430.
 // Vega20 fails to access GPU memory larger than the return value of GetMaxMemoryAllocSize() of
@@ -100,8 +97,6 @@ SlowdownFactor(int n_oper, const double oper_factor, const double multiple_oper_
 bool GemmBwdBase::IsApplicable(const ExecutionContext& ctx, const ProblemDescription& problem) const
 {
 #if MIOPEN_USE_GEMM
-    if(conv::gemm::IsWorkaroundIssue1315(ctx))
-        return false;
     const auto& dyDesc             = problem.GetIn();
     const auto& wDesc              = problem.GetWeights();
     const auto& dxDesc             = problem.GetOut();
@@ -464,8 +459,8 @@ size_t GemmBwd1x1_stride1::GetWorkspaceSize(const ExecutionContext&,
     return 0;
 }
 
-bool GemmBwd1x1_stride1::IsApplicableBeforeWorkaround(const ExecutionContext& context,
-                                                      const ProblemDescription& problem) const
+bool GemmBwd1x1_stride1::IsApplicable(const ExecutionContext& context,
+                                      const ProblemDescription& problem) const
 {
 #if MIOPEN_USE_GEMM
     if(!GemmBwdBase::IsApplicable(context, problem))
@@ -487,18 +482,6 @@ bool GemmBwd1x1_stride1::IsApplicableBeforeWorkaround(const ExecutionContext& co
 #endif
 }
 
-bool GemmBwd1x1_stride1::IsApplicable(const ExecutionContext& context,
-                                      const ProblemDescription& problem) const
-{
-#if MIOPEN_USE_GEMM && (!MIOPEN_USE_MIOPENGEMM || !WORKAROUND_MIOPENGEMM_ISSUE_59)
-    return IsApplicableBeforeWorkaround(context, problem);
-#else
-    std::ignore = context;
-    std::ignore = problem;
-    return false;
-#endif
-}
-
 ConvSolution GemmBwd1x1_stride1::GetSolution(const ExecutionContext&,
                                              const ProblemDescription& problem) const
 {
@@ -695,7 +678,7 @@ bool GemmBwdRest::IsApplicable(const ExecutionContext& context,
         return false;
 
     return !GemmBwd1x1_stride2{}.IsApplicable(context, problem) &&
-           !GemmBwd1x1_stride1{}.IsApplicableBeforeWorkaround(context, problem) &&
+           !GemmBwd1x1_stride1{}.IsApplicable(context, problem) &&
            GetWorkspaceSize(context, problem) > 0;
 #else
     std::ignore = context;
diff --git a/src/solver/gemm_common.cpp b/src/solver/gemm_common.cpp
deleted file mode 100644
index be0294ad48..0000000000
--- a/src/solver/gemm_common.cpp
+++ /dev/null
@@ -1,59 +0,0 @@
-/*******************************************************************************
- *
- * 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/ROCm/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 solver {
-namespace conv {
-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 conv
-} // namespace solver
-} // namespace miopen
diff --git a/src/solver/gemm_wrw.cpp b/src/solver/gemm_wrw.cpp
index b58d71f104..d204ae0a5f 100644
--- a/src/solver/gemm_wrw.cpp
+++ b/src/solver/gemm_wrw.cpp
@@ -3,7 +3,6 @@
 #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>
 
@@ -67,8 +66,6 @@ SlowdownFactor(int n_oper, const double oper_factor, const double multiple_oper_
 bool GemmWrwBase::IsApplicable(const ExecutionContext& ctx, const ProblemDescription& problem) const
 {
 #if MIOPEN_USE_GEMM
-    if(conv::gemm::IsWorkaroundIssue1315(ctx))
-        return false;
     const auto& dyDesc             = problem.GetIn();
     const auto& dwDesc             = problem.GetWeights();
     const auto& xDesc              = problem.GetOut();
diff --git a/src/sqlite_db.cpp b/src/sqlite_db.cpp
index 89a223ca7d..b697932a95 100644
--- a/src/sqlite_db.cpp
+++ b/src/sqlite_db.cpp
@@ -28,7 +28,6 @@
 #include <miopen/errors.hpp>
 #include <miopen/lock_file.hpp>
 #include <miopen/logger.hpp>
-#include <miopen/md5.hpp>
 #include <miopen/conv/problem_description.hpp>
 #include <miopen/exp_backoff.hpp>
 
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index ffaeb93228..1a15f602ac 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -1118,23 +1118,6 @@ COMMAND $<TARGET_FILE:test_conv2d> ${MIOPEN_TEST_FLOAT_ARG} --input 352 16 7 1 -
 
 if (0) #disabled too many errors
 if(MIOPEN_TEST_LIMIT GREATER 0)
-if(${MIOPEN_USE_MIOPENGEMM} AND (MIOPEN_hip_VERSION VERSION_EQUAL 3.7))
-    add_custom_test(test_conv3d_extra  SKIP_UNLESS_ALL GFX94X_ENABLED GFX103X_ENABLED
-    ENVIRONMENT MIOPEN_LOG_LEVEL=6
-    COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 2 2 2 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 2 2 2 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 1 1 1 2 2 2 ${MIOPEN_TEST_FLAGS_ARGS}
-    #COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 1 1 1 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-
-    #ROCM3.7 compiler problems
-    #COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    #COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 2 2 2 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    #COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 140 602 --weights 16 16 3 11 11 --pads_strides_dilations 1 1 1 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    #COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 140 602 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
-    )
-    message(STATUS "test_conv3d_extra reduced set")
-else()
     add_custom_test(test_conv3d_extra  SKIP_UNLESS_ALL GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED
     COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
     COMMAND $<TARGET_FILE:test_conv3d> ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50  50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 2 2 2 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}
@@ -1148,7 +1131,6 @@ else()
     )
 endif()
 endif()
-endif()
 
 
 add_custom_test(test_conv_3d SKIP_UNLESS_ALL GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED
diff --git a/test/conv_common.hpp b/test/conv_common.hpp
index e387e98b5b..6b3ce738ef 100644
--- a/test/conv_common.hpp
+++ b/test/conv_common.hpp
@@ -60,7 +60,7 @@
 #include "cpu_bias.hpp"
 #include "random.hpp"
 
-#define TEST_DIRECT_SUPPORTED_CONFIG_ONLY (!MIOPEN_USE_ROCBLAS && !MIOPEN_USE_MIOPENTENSILE)
+#define TEST_DIRECT_SUPPORTED_CONFIG_ONLY (!MIOPEN_USE_ROCBLAS)
 
 #define WORKAROUND_MI100_ROM37_HIP_COMPILER_CRASH \
     (HIP_PACKAGE_VERSION_MAJOR == 3 && HIP_PACKAGE_VERSION_MINOR == 7)
diff --git a/test/gtest/cache.cpp b/test/gtest/cache.cpp
index 5c512e6a8a..3163ddafd0 100644
--- a/test/gtest/cache.cpp
+++ b/test/gtest/cache.cpp
@@ -28,7 +28,6 @@
 #include <miopen/kern_db.hpp>
 #include <miopen/temp_file.hpp>
 
-#include <miopen/md5.hpp>
 #include "test.hpp"
 #include "random.hpp"
 
@@ -138,13 +137,6 @@ TEST(TestCache, check_kern_db)
 
 TEST(TestCache, check_cache_file)
 {
-    auto p = miopen::GetCacheFile("gfx", "base", "args", false);
+    auto p = miopen::GetCacheFile("gfx", "base", "args");
     EXPECT_TRUE(p.filename().string() == "base.o");
 }
-
-TEST(TestCache, check_cache_str)
-{
-    auto p    = miopen::GetCacheFile("gfx", "base", "args", true);
-    auto name = miopen::md5("base");
-    EXPECT_TRUE(p.filename().string() == name + ".o");
-}
diff --git a/test/gtest/db_sync.cpp b/test/gtest/db_sync.cpp
index aed0bdaed0..3d06147c69 100644
--- a/test/gtest/db_sync.cpp
+++ b/test/gtest/db_sync.cpp
@@ -470,7 +470,7 @@ void BuildKernel(const std::string& program_file,
 #else
     try
     {
-        auto p = handle.LoadProgram(program_file, program_args, false, "");
+        auto p = handle.LoadProgram(program_file, program_args, "");
     }
     catch(std::exception&)
     {
diff --git a/test/handle_test.cpp b/test/handle_test.cpp
index 2548a7ad4b..409d2864a8 100644
--- a/test/handle_test.cpp
+++ b/test/handle_test.cpp
@@ -94,8 +94,15 @@ void run2s(miopen::Handle& h, std::size_t n, kernel_type_t kern_type)
     auto data_dev = h.Write(data_in);
     if(kern_type == miopenOpenCLKernelType)
     {
-        h.AddKernel("GEMM", "", Write2s(miopenOpenCLKernelType), "write", {n, 1, 1}, {n, 1, 1}, "")(
-            data_dev.get());
+        h.AddKernel("NoAlgo",
+                    "",
+                    "test_ocl.cl",
+                    "write",
+                    {n, 1, 1},
+                    {n, 1, 1},
+                    "",
+                    0,
+                    Write2s(miopenOpenCLKernelType))(data_dev.get());
     }
     else if(kern_type == miopenHIPKernelType)
     {
@@ -107,7 +114,6 @@ void run2s(miopen::Handle& h, std::size_t n, kernel_type_t kern_type)
                     {n, 1, 1},
                     "",
                     0,
-                    false,
                     Write2s(miopenHIPKernelType))(data_dev.get());
     }
     else
@@ -161,11 +167,27 @@ void test_errors(kernel_type_t kern_type)
     if(kern_type == miopenOpenCLKernelType)
     {
         EXPECT(throws([&] {
-            h.AddKernel("GEMM", "", WriteError(kern_type), "write", {1, 1, 1}, {1, 1, 1}, "");
+            h.AddKernel("NoAlgo",
+                        "",
+                        "error_ocl.cl",
+                        "write",
+                        {1, 1, 1},
+                        {1, 1, 1},
+                        "",
+                        0,
+                        WriteError(kern_type));
         }));
         try
         {
-            h.AddKernel("GEMM", "", WriteError(kern_type), "write", {1, 1, 1}, {1, 1, 1}, "");
+            h.AddKernel("NoAlgo",
+                        "",
+                        "error_ocl.cl",
+                        "write",
+                        {1, 1, 1},
+                        {1, 1, 1},
+                        "",
+                        0,
+                        WriteError(kern_type));
         }
         catch(miopen::Exception& e)
         {
@@ -183,7 +205,6 @@ void test_errors(kernel_type_t kern_type)
                         {1, 1, 1},
                         "",
                         0,
-                        false,
                         WriteError(miopenHIPKernelType));
         }));
         try
@@ -196,7 +217,6 @@ void test_errors(kernel_type_t kern_type)
                         {1, 1, 1},
                         "",
                         0,
-                        false,
                         WriteError(miopenHIPKernelType));
         }
         catch(miopen::Exception& e)
@@ -235,7 +255,15 @@ void test_warnings(kernel_type_t kern_type)
     if(kern_type == miopenOpenCLKernelType)
     {
         EXPECT(throws([&] {
-            h.AddKernel("GEMM", "", WriteNop(kern_type), "write", {1, 1, 1}, {1, 1, 1}, "");
+            h.AddKernel("NoAlgo",
+                        "",
+                        "nop_ocl.cl",
+                        "write",
+                        {1, 1, 1},
+                        {1, 1, 1},
+                        "",
+                        0,
+                        WriteNop(kern_type));
             MIOPEN_LOG_E("FAILED: Build of the OpenCL kernel should produce warnings");
         }));
     }
@@ -250,7 +278,6 @@ void test_warnings(kernel_type_t kern_type)
                         {1, 1, 1},
                         "",
                         0,
-                        false,
                         WriteNop(kern_type));
             MIOPEN_LOG_E("FAILED: Build of the HIP kernel 'nop_hip.cpp' should produce warnings");
         }));