Merge branch 'imagescaler' of github.com:ROCmSoftwarePlatform/MIGraph into imagescaler

dev-requirements.txt

@@ -2,8 +2,5 @@ pfultz2/rocm-recipes
 pcre
 danmar/cppcheck@f965e5873 -DHAVE_RULES=1
 ROCm-Developer-Tools/HIP@3a41f286203968421c557338d6fb39c36f3c717c
-# Needed for clang-ocl
-RadeonOpenCompute/rocm-cmake@6240bb3 --build
-RadeonOpenCompute/clang-ocl@799713643b5591a3b877c586ef2c7fbc012af819
 # python/cpython@v3.6.6 -X autotools -H sha256:92aa914572c695c0aeb01b0a214813f414da4b51a371234df514a74761f2bb36
 -f requirements.txt

requirements.txt

 google/protobuf -DCMAKE_POSITION_INDEPENDENT_CODE=On
 RadeonOpenCompute/rocm-cmake@6240bb3 --build
-ROCmSoftwarePlatform/MIOpen,http://gitlab1.amd.com/pfultz/miopen/-/archive/f9baaabdf7c26b70648966306b2e1adb3ba3cab5/miopen-f9baaabdf7c26b70648966306b2e1adb3ba3cab5.tar.bz2
-ROCmSoftwarePlatform/rocBLAS@v14.0.1
+ROCmSoftwarePlatform/MIOpen,http://gitlab1.amd.com/pfultz/miopen/-/archive/04a9a6a6cb3a3e9457fdc207669992c08b7df5d4/miopen-04a9a6a6cb3a3e9457fdc207669992c08b7df5d4.tar.bz2
+ROCmSoftwarePlatform/rocBLAS@v14.3.0
 blaze,https://bitbucket.org/blaze-lib/blaze/get/f0755dea0e03.tar.gz -X header -DHEADER_DIR=blaze

src/include/migraph/ranges.hpp

@@ -45,6 +45,12 @@ bool contains(const C& c, const T& x)
     return generic_find(c, x) != c.end();
 }
 
+template <class T>
+bool contains(const std::initializer_list<T>& c, const T& x)
+{
+    return generic_find(c, x) != c.end();
+}
+
 template <class T, class U>
 bool contains(const std::initializer_list<T>& c, const U& x)
 {

src/program.cpp

@@ -427,7 +427,7 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
     }
 
     double total_time             = common_average(total_vec);
-    double rate                   = std::ceil(1000.0 / total_time);
+    double rate                   = 1000.0 / total_time;
     double overhead_time          = common_average(overhead_vec);
     double overhead_percent       = overhead_time * 100.0 / total_time;
     double total_instruction_time = 0.0;

src/shape.cpp

@@ -157,9 +157,7 @@ bool shape::scalar() const
 {
     assert(this->lens().size() == this->strides().size());
     // if any stride > 0, then accumulate will return false
-    return std::accumulate(this->strides().begin(),
-                           this->strides().end(),
-                           std::size_t(0)) == 0;
+    return std::accumulate(this->strides().begin(), this->strides().end(), std::size_t(0)) == 0;
 }
 
 bool shape::standard() const { return impl->m_standard; }

src/targets/gpu/device/include/migraph/gpu/device/nary.hpp

@@ -333,7 +333,8 @@ nary(hipStream_t stream, const argument& result, const argument& arg1, const arg
 {
     return [=](auto f) {
         // TODO: Check result and arg1 shape is the same
-        if(arg1.get_shape().standard() and arg2.get_shape().broadcasted() and not arg2.get_shape().scalar())
+        if(arg1.get_shape().standard() and arg2.get_shape().broadcasted() and
+           not arg2.get_shape().scalar())
         {
             auto not_zero       = [](auto x) { return x != 0; };
             const auto& strides = arg2.get_shape().strides();

src/targets/gpu/fuse_ops.cpp

@@ -134,15 +134,12 @@ MIGRAPH_PRED_MATCHER(fusable_conv, instruction_ref ins)
         return false;
     auto wei = ins->inputs().at(1)->get_shape();
     assert(wei.lens().size() == 4);
-    auto channels = wei.lens()[1] * wei.lens()[0];
-    if(wei.lens()[0] > 64 and channels > 32768)
-        return false;
     auto conv = any_cast<miopen_convolution>(ins->get_operator());
-    if(conv.algo == miopenConvolutionFwdAlgoWinograd)
+    if(wei.lens()[1] > 512 and conv.algo != miopenConvolutionFwdAlgoWinograd)
         return false;
     auto op = conv.op;
-    return op.padding == make_array<size_t>(0, 0) and op.stride == make_array<size_t>(1, 1) and
-           op.dilation == make_array<size_t>(1, 1);
+    return contains({{0, 0}, {1, 1}, {2, 2}}, op.padding) and
+           contains({{0, 0}, {1, 1}}, op.stride) and op.dilation == make_array<size_t>(1, 1);
 }
 
 struct hip_triadd

test/cpu_ops_test.cpp

@@ -543,36 +543,46 @@ void imagescaler_test()
 {
     migraph::program p;
     migraph::shape s{migraph::shape::float_type, {1, 3, 2, 2}};
-    auto img = p.add_literal(migraph::literal{s, {
-                                                  0.2, 0.3,
-                                                  0.5, 0.4,
-                                                  
-                                                  0.7, 0.8,
-                                                  0.1, 0.9,
-                                                
-                                                  0.15, 0.25,
-                                                  0.35, 0.45
-                                                  }});
+    auto img           = p.add_literal(migraph::literal{s,
+                                              {0.2,
+                                               0.3,
+                                               0.5,
+                                               0.4,
+
+                                               0.7,
+                                               0.8,
+                                               0.1,
+                                               0.9,
+
+                                               0.15,
+                                               0.25,
+                                               0.35,
+                                               0.45}});
     auto scale_val     = p.add_literal(2.f);
     auto scaled_tensor = p.add_instruction(migraph::op::scalar{s}, scale_val);
     auto img_scaled    = p.add_instruction(migraph::op::mul{}, img, scaled_tensor);
-    auto bias_vals = p.add_literal(migraph::literal{migraph::shape{migraph::shape::float_type, {3}}, {0.01, 0.02, 0.03}});
+    auto bias_vals     = p.add_literal(
+        migraph::literal{migraph::shape{migraph::shape::float_type, {3}}, {0.01, 0.02, 0.03}});
     auto bias_bcast = p.add_instruction(migraph::op::broadcast{1, s}, bias_vals);
     p.add_instruction(migraph::op::add{}, img_scaled, bias_bcast);
     p.compile(migraph::cpu::cpu_target{});
     auto result = p.eval({});
     std::vector<float> results_vector(12);
     result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
-    std::vector<float> gold = {
-                               0.41, 0.61,
-                               1.01, 0.81,
-                            
-                               1.42, 1.62,
-                               0.22, 1.82,
-                               
-                               0.33, 0.53,
-                               0.73, 0.93
-                               };
+    std::vector<float> gold = {0.41,
+                               0.61,
+                               1.01,
+                               0.81,
+
+                               1.42,
+                               1.62,
+                               0.22,
+                               1.82,
+
+                               0.33,
+                               0.53,
+                               0.73,
+                               0.93};
     EXPECT(migraph::verify_range(results_vector, gold));
 }