Set find mode in miopen to normal (#711)

src/targets/gpu/CMakeLists.txt

@@ -256,6 +256,17 @@ target_compile_definitions(migraphx_gpu PRIVATE
     "-DMIGRAPHX_EXTRACT_KERNEL=${MIGRAPHX_EXTRACT_KERNEL}"
 )
 
+# Check miopen find mode api
+include(CheckLibraryExists)
+get_target_property(MIOPEN_LOCATION MIOpen LOCATION)
+check_library_exists(MIOpen "miopenHiddenSetConvolutionFindMode" "${MIOPEN_LOCATION}" HAS_FIND_MODE_API)
+if(HAS_FIND_MODE_API)
+    target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_MODE_API)
+    message(STATUS "MIOpen has find mode api")
+else()
+    message(STATUS "MIOpen does not have find mode api")
+endif()
+
 # Workaround broken rocblas headers
 target_compile_definitions(migraphx_gpu PUBLIC -D__HIP_PLATFORM_HCC__=1)
 target_compile_options(migraphx_gpu PRIVATE -std=c++17)

src/targets/gpu/include/migraphx/gpu/miopen.hpp

@@ -9,6 +9,11 @@
 #include <miopen/miopen.h>
 #include <migraphx/config.hpp>
 
+#ifdef HAS_FIND_MODE_API
+extern "C" miopenStatus_t miopenHiddenSetConvolutionFindMode(miopenConvolutionDescriptor_t convDesc,
+                                                             int findMode);
+#endif
+
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
@@ -97,6 +102,9 @@ inline convolution_descriptor make_conv(const T& op)
         c.get(), padding.size(), padding.data(), stride.data(), dilation.data(), c_mode);
     if(op.group > 1)
         miopenSetConvolutionGroupCount(c.get(), op.group);
+#ifdef HAS_FIND_MODE_API
+    miopenHiddenSetConvolutionFindMode(c.get(), 1); // Normal mode
+#endif
     return c;
 }