CP for Find-2.0 and 3D-Unet fix (#1448)

* Find2.0 changes for the Quant  and De-Convolution (#1408)
* fix for 3d-UNet

src/include/migraphx/reflect.hpp

@@ -56,11 +56,11 @@ auto reflect_impl(rank<0>, T&, Selector)
 }
 
 template <class T>
-auto reflectable_impl(rank<1>, T&& x)
+auto reflectable_impl(rank<1>, const T& x)
     -> decltype(T::reflect(x, reflect_placeholder{}), std::true_type{});
 
 template <class T>
-auto reflectable_impl(rank<0>, T &&) -> decltype(std::false_type{});
+auto reflectable_impl(rank<0>, const T&) -> decltype(std::false_type{});
 
 template <class T>
 struct remove_rvalue_reference
@@ -111,8 +111,18 @@ auto reflect(T& x, Selector f)
 template <class T>
 auto reflect_tie(T& x)
 {
-    return reflect(x, [](auto&& y, auto&&...) { return detail::wrap<decltype(y)>(y); })(
-        [](auto&&... xs) { return detail::auto_tuple(xs.get()...); });
+    return reflect(x, [](auto&& y, auto&&...) {
+        // cppcheck-suppress UnnecessaryElseStatement
+        if constexpr(is_reflectable<decltype(y)>{})
+        {
+            auto t = reflect_tie(y);
+            return detail::wrap<decltype(t)>(t);
+        }
+        else
+        {
+            return detail::wrap<decltype(y)>(y);
+        }
+    })([](auto&&... xs) { return detail::auto_tuple(xs.get()...); });
 }
 
 template <class T, class F>

src/include/migraphx/streamutils.hpp

@@ -26,7 +26,9 @@
 
 #include <ostream>
 #include <algorithm>
+#include <migraphx/reflect.hpp>
 #include <migraphx/rank.hpp>
+#include <migraphx/requires.hpp>
 #include <migraphx/config.hpp>
 #include <vector>
 
@@ -83,6 +85,20 @@ auto stream_write_value_impl(rank<0>, std::ostream& os, const Range& r)
     os << "}";
 }
 
+template <class T, MIGRAPHX_REQUIRES(is_reflectable<T>{})>
+void stream_write_value_impl(rank<0>, std::ostream& os, const T& x)
+{
+    char delim = '{';
+    reflect_each(x, [&](auto&& y, auto name) {
+        os << delim;
+        os << name << "=";
+        stream_write_value_impl(rank<2>{}, os, y);
+        delim = ',';
+    });
+    if(delim == ',')
+        os << "}";
+}
+
 } // namespace detail
 
 template <class T>

src/targets/gpu/CMakeLists.txt

@@ -84,8 +84,6 @@ add_library(migraphx_gpu
     compile_hip.cpp
     compile_hip_code_object.cpp
     compiler.cpp
-    convolution.cpp
-    deconvolution.cpp
     device_name.cpp
     elu.cpp
     fuse_mlir.cpp
@@ -110,7 +108,6 @@ add_library(migraphx_gpu
     pad.cpp
     perfdb.cpp
     pooling.cpp
-    quant_convolution.cpp
     reverse.cpp
     rnn_variable_seq_lens.cpp
     rocblas.cpp
@@ -146,14 +143,11 @@ register_migraphx_gpu_ops(hip_
 register_migraphx_gpu_ops(miopen_
     abs
     contiguous
-    convolution
-    deconvolution
     elu
     int8_conv_pack
     leaky_relu
     lrn
     pooling
-    quant_convolution
 )
 register_op(migraphx_gpu 
     HEADER migraphx/gpu/rnn_variable_seq_lens.hpp 
@@ -167,6 +161,9 @@ register_op(migraphx_gpu
     HEADER migraphx/gpu/gemm.hpp 
     OPERATORS gpu::rocblas_gemm<op::dot> gpu::rocblas_gemm<op::quant_dot>
     INCLUDES migraphx/gpu/context.hpp)
+register_op(migraphx_gpu HEADER migraphx/gpu/convolution.hpp 
+    OPERATORS gpu::miopen_convolution<op::convolution> gpu::miopen_convolution<op::deconvolution> gpu::miopen_convolution<op::quant_convolution>
+    INCLUDES migraphx/gpu/context.hpp)
 rocm_set_soversion(migraphx_gpu ${MIGRAPHX_SO_VERSION})
 rocm_clang_tidy_check(migraphx_gpu)
 
@@ -239,11 +236,13 @@ get_target_property(MIOPEN_LOCATION MIOpen LOCATION)
 check_library_exists(MIOpen "miopenHiddenSetConvolutionFindMode" "${MIOPEN_LOCATION}" HAS_FIND_MODE_API)
 check_library_exists(MIOpen "miopenFindSolutions" "${MIOPEN_LOCATION}" HAS_FIND_2_API)
 
-if(HAS_FIND_2_API) 
+set(MIGRAPHX_USE_FIND_2_API "${HAS_FIND_2_API}" CACHE BOOL "")
+
+if(MIGRAPHX_USE_FIND_2_API) 
     target_compile_definitions(migraphx_gpu PUBLIC -DMIGRAPHX_HAS_FIND_2_API)
     message(STATUS "MIGraphx is using Find-2.0 API of MIOpen")
 else()
-    message(STATUS "MIOpen does not have Find-2.0 API")
+    message(STATUS "MIGraphx is using legacy Find API in MIOpen")
 endif()
 
 if(HAS_FIND_MODE_API)