Merge branch 'master' into fp16

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
 half,https://github.com/pfultz2/half/archive/1.12.0.tar.gz -X header -H sha256:0a08660b68abb176ebc2a0cdf8de46e3182a7f46c66443bb80dbfaaec98cf969

src/dead_code_elimination.cpp

@@ -3,9 +3,30 @@
 #include <migraph/instruction.hpp>
 #include <migraph/iterator_for.hpp>
 #include <migraph/functional.hpp>
+#include <migraph/ranges.hpp>
 
 namespace migraph {
 
+template <class Range, class Iterator>
+std::ptrdiff_t bidistance(const Range& r, Iterator start, Iterator last)
+{
+    auto start_forward   = start;
+    auto start_backwards = start;
+    std::size_t n        = 0;
+    while(start_forward != last and start_backwards != last)
+    {
+        n++;
+        if(start_forward != r.end())
+            start_forward++;
+        if(start_backwards != r.begin())
+            start_backwards--;
+    }
+    if(start_forward == last)
+        return n;
+    else
+        return -n;
+}
+
 void dead_code_elimination::apply(program& p) const
 {
     auto last = std::prev(p.end());
@@ -16,18 +37,21 @@ void dead_code_elimination::apply(program& p) const
         if(ins == p.begin())
             continue;
         const auto i = std::prev(ins);
-        // Skip instruction with empty shape as output unless its a builtin
-        if(i->get_shape().elements() == 0 and not(i->name().front() == '@'))
-            continue;
         // Skip the last instruction
         if(i == last)
             break;
+        // Skip instruction with empty shape as output unless its a builtin
+        if(i->get_shape().elements() == 0 and not(i->name().front() == '@'))
+            continue;
+        assert(bidistance(p, i, last) > 0);
         fix([&](auto self, auto leaf) {
             assert(p.has_instruction(leaf));
             if(leaf->outputs().empty())
             {
                 auto args = leaf->inputs();
                 leaf->clear_arguments();
+                assert(bidistance(p, last, leaf) < 0);
+                assert(leaf != ins);
                 p.move_instruction(leaf, p.end());
                 for(auto arg : args)
                     self(arg);

src/include/migraph/operation.hpp

@@ -256,7 +256,6 @@ struct operation
 
         shape compute_shape(const std::vector<shape>& input) const override
         {
-
             return private_detail_te_value.compute_shape(input);
         }
 

src/include/migraph/operators.hpp

@@ -306,10 +306,6 @@ struct contiguous
         check_shapes{inputs, *this}.has(1);
         auto lens = inputs.at(0).lens();
         auto t    = inputs.at(0).type();
-        if(lens.size() < 2)
-        {
-            MIGRAPH_THROW("Number of dimensions should exceed 1");
-        }
         return {t, lens};
     }
 };
@@ -761,6 +757,27 @@ struct broadcast
     }
 };
 
+struct scalar
+{
+    shape scalar_bcast;
+
+    std::string name() const { return "scalar"; }
+
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        assert(check_shapes{inputs}.has(1).only_dims(1).size() == 1);
+        auto t     = inputs.at(0).type();
+        auto input = inputs.at(0);
+        std::vector<std::size_t> strides(scalar_bcast.lens().size(), 0);
+        return {t, scalar_bcast.lens(), strides};
+    }
+
+    argument compute(context&, shape output_shape, std::vector<argument> args) const
+    {
+        return {std::move(output_shape), std::move(args.at(0).data)};
+    }
+};
+
 struct binary
 {
     shape compute_shape(std::vector<shape> inputs) const

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/include/migraph/shape.hpp

@@ -92,6 +92,8 @@ struct shape
     /// Returns true if the shape is in its standard format. That is, the shape is both packed and
     /// not transposed.
     bool standard() const;
+    /// Returns true if all strides are equal to 0 (scalar tensor)
+    bool scalar() const;
 
     friend bool operator==(const shape& x, const shape& y);
     friend bool operator!=(const shape& x, const shape& y);

src/onnx/onnx.cpp

@@ -56,6 +56,7 @@ struct onnx_parser
         add_generic_op("Sub", op::sub{});
         add_generic_op("Sum", op::add{});
 
+        add_mem_op("ImageScaler", &onnx_parser::parse_imagescaler);
         add_mem_op("LeakyRelu", &onnx_parser::parse_leaky_relu);
         add_mem_op("Constant", &onnx_parser::parse_constant);
         add_mem_op("Conv", &onnx_parser::parse_conv);
@@ -315,7 +316,7 @@ struct onnx_parser
                                      attribute_map attributes,
                                      std::vector<instruction_ref> args)
     {
-        float alpha = 0.01;
+        float alpha = 0.01; // default alpha val for leaky relu
         if(contains(attributes, "alpha"))
         {
             alpha = parse_value(attributes.at("alpha")).at<float>();
@@ -324,6 +325,34 @@ struct onnx_parser
         return prog.add_instruction(op, args.front());
     }
 
+    instruction_ref parse_imagescaler(const std::string&,
+                                      attribute_map attributes,
+                                      std::vector<instruction_ref> args)
+    {
+        float scale = 1.0;
+        std::vector<float> bias{};
+        if(contains(attributes, "scale"))
+        {
+            scale = parse_value(attributes.at("scale")).at<float>();
+        }
+
+        if(contains(attributes, "bias"))
+        {
+            auto&& bias_floats = attributes["bias"].floats();
+            bias               = std::vector<float>(bias_floats.begin(), bias_floats.end());
+        }
+        auto input_shape = args.front()->get_shape();
+
+        auto scale_val = prog.add_literal(scale);
+        auto bias_vals = prog.add_literal(
+            migraph::literal{migraph::shape{migraph::shape::float_type, {bias.size()}}, bias});
+
+        auto scale_tensor = prog.add_instruction(migraph::op::scalar{input_shape}, scale_val);
+        auto img_scaled   = prog.add_instruction(migraph::op::mul{}, args.front(), scale_tensor);
+        auto bias_bcast   = prog.add_instruction(migraph::op::broadcast{1, input_shape}, bias_vals);
+        return prog.add_instruction(migraph::op::add{}, img_scaled, bias_bcast);
+    }
+
     void parse_from(std::istream& is)
     {
         onnx::ModelProto model;

src/opt/memory_coloring_impl.cpp

@@ -195,6 +195,7 @@ void memory_coloring_impl::register_operand_alias()
     operand_alias["broadcast"]         = 0;
     operand_alias["reshape"]           = 0;
     operand_alias["pass"]              = 0;
+    operand_alias["scalar"]            = 0;
 }
 
 void memory_coloring_impl::rewrite()

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

@@ -29,8 +29,8 @@ struct shape_impl
         : m_type(t), m_lens(std::move(l)), m_strides(std::move(s))
     {
         assert(m_lens.size() == m_strides.size());
-        assert(std::any_of(m_strides.begin(), m_strides.end(), [](auto x) { return x > 0; }) and
-               "At least one stride must be non-zero");
+        // assert(std::any_of(m_strides.begin(), m_strides.end(), [](auto x) { return x > 0; }) and
+        //        "At least one stride must be non-zero");
         m_standard = this->elements() == this->element_space() and
                      std::is_sorted(m_strides.rbegin(), m_strides.rend());
     }
@@ -153,6 +153,13 @@ bool shape::broadcasted() const
                            std::multiplies<std::size_t>()) == 0;
 }
 
+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;
+}
+
 bool shape::standard() const { return impl->m_standard; }
 
 std::size_t shape::element_space() const { return impl->element_space(); }

src/targets/gpu/CMakeLists.txt

@@ -10,10 +10,11 @@ if(NOT TARGET MIOpen)
     message(SEND_ERROR "Cant find miopen")
 endif()
 
-add_library(migraph_device 
+add_library(migraph_device
     device/add.cpp
     device/add_relu.cpp
     device/contiguous.cpp
+    device/mul.cpp
     device/concat.cpp
 )
 rocm_clang_tidy_check(migraph_device)
@@ -36,6 +37,7 @@ add_library(migraph_gpu
     relu.cpp
     leaky_relu.cpp
     add.cpp
+    mul.cpp
     batchnorm.cpp
     write_literals.cpp
     rocblas.cpp

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())
+        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/device/mul.cpp

+#include <migraph/gpu/device/mul.hpp>
+#include <migraph/gpu/device/nary.hpp>
+
+namespace migraph {
+namespace gpu {
+namespace device {
+
+void mul(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2)
+{
+    nary(stream, result, arg1, arg2)([](auto x, auto y) { return x * y; });
+}
+
+void mul(hipStream_t stream,
+         const argument& result,
+         const argument& arg1,
+         const argument& arg2,
+         const argument& arg3)
+{
+    nary(stream, result, arg1, arg2, arg3)([](auto x, auto y, auto z) { return x * y * z; });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace migraph

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

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

+
+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_MUL_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_MUL_HPP
+
+#include <migraph/argument.hpp>
+#include <hip/hip_runtime_api.h>
+
+namespace migraph {
+namespace gpu {
+namespace device {
+
+void mul(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2);
+
+void mul(hipStream_t stream,
+         const argument& result,
+         const argument& arg1,
+         const argument& arg2,
+         const argument& arg3);
+
+} // namespace device
+} // namespace gpu
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/mul.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_MUL_HPP
+#define MIGRAPH_GUARD_RTGLIB_MUL_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/mul.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct hip_mul
+{
+    std::string name() const { return "gpu::mul"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument compute(context&, const shape&, const std::vector<argument>& args) const;
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/lowering.cpp

@@ -19,6 +19,7 @@
 #include <migraph/gpu/leaky_relu.hpp>
 #include <migraph/gpu/softmax.hpp>
 #include <migraph/gpu/add.hpp>
+#include <migraph/gpu/mul.hpp>
 #include <migraph/gpu/batchnorm.hpp>
 #include <migraph/gpu/pooling.hpp>
 #include <migraph/gpu/gemm.hpp>
@@ -65,6 +66,10 @@ struct miopen_apply
             {
                 check_shape(s, apply_add(it));
             }
+            else if(it->name() == "mul")
+            {
+                check_shape(s, apply_mul(it));
+            }
             else if(it->name() == "dot")
             {
                 check_shape(s, apply_gemm(it));
@@ -163,6 +168,13 @@ struct miopen_apply
             ins, hip_add{}, ins->inputs().at(0), ins->inputs().at(1), output);
     }
 
+    instruction_ref apply_mul(instruction_ref ins)
+    {
+        auto output = insert_allocation(ins, ins->get_shape());
+        return prog->replace_instruction(
+            ins, hip_mul{}, ins->inputs().at(0), ins->inputs().at(1), output);
+    }
+
     instruction_ref apply_gemm(instruction_ref ins)
     {
         auto&& op   = any_cast<op::dot>(ins->get_operator());

src/targets/gpu/mul.cpp

+#include <migraph/gpu/mul.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+shape hip_mul::compute_shape(const std::vector<shape>& inputs) const
+{
+    // check_shapes{inputs, *this}.has(3).standard();
+    check_shapes{inputs, *this}.has(3);
+    return inputs.at(0);
+}
+
+argument hip_mul::compute(context& ctx, const shape&, const std::vector<argument>& args) const
+{
+    device::mul(ctx.get_stream().get(), args[2], args[0], args[1]);
+    return args[2];
+}
+
+} // namespace gpu
+
+} // namespace migraph

test/cpu_ops_test.cpp

@@ -539,6 +539,53 @@ void leaky_relu_test()
     EXPECT(migraph::verify_range(results_vector, gold));
 }
 
+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 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_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};
+    EXPECT(migraph::verify_range(results_vector, gold));
+}
+
 void reshape_test()
 {
     migraph::shape a_shape{migraph::shape::float_type, {24, 1, 1, 1}};
@@ -994,6 +1041,7 @@ int main()
     add_test();
     broadcast_test();
     add_broadcast_test();
+    imagescaler_test();
     sub_test();
     mul_test();
     div_test();