Merge branch 'develop' into trace-eval

src/targets/gpu/lowering.cpp

@@ -45,6 +45,7 @@
 #include <migraphx/gpu/pad.hpp>
 #include <migraphx/gpu/gather.hpp>
 #include <migraphx/gpu/lrn.hpp>
+#include <migraphx/gpu/convert.hpp>
 #include <migraphx/gpu/clip.hpp>
 #include <utility>
 #include <functional>
@@ -102,6 +103,7 @@ struct miopen_apply
         add_extend_op<hip_logsoftmax, op::logsoftmax>("logsoftmax");
         add_extend_op<hip_gather, op::gather>("gather");
         add_extend_op<hip_pad, op::pad>("pad");
+        add_extend_op<hip_convert, op::convert>("convert");
         add_extend_op<hip_clip, op::clip>("clip");
 
         add_lrn_op();

src/targets/gpu/target.cpp

@@ -51,7 +51,7 @@ std::vector<pass> target::get_passes(migraphx::context& gctx) const
         propagate_constant{},
         dead_code_elimination{},
         auto_contiguous{},
-        //simplify_reshapes{},
+        simplify_reshapes{},
         dead_code_elimination{},
         lowering{ctx},
         eliminate_concat{concat_gpu_optimization{}},

src/tf/tf.cpp

@@ -53,15 +53,16 @@ struct tf_parser
     template <class T>
     std::vector<T> parse_axes(std::vector<T> axes) const
     {
-        std::vector<T> new_axes;
         if(is_nhwc)
         {
+            std::vector<T> new_axes;
             std::transform(axes.begin(),
                            axes.end(),
                            std::back_inserter(new_axes),
                            [&](size_t axis) { return parse_axis(axis); });
+            return new_axes;
         }
-        return new_axes;
+        return axes;
     }
 
     // tf stores certain attributes such as strides, dilations, as a 4D input.
@@ -392,7 +393,9 @@ struct tf_parser
         int64_t out_channels = num_channels * multiplier;
         new_weights_shape[0] = out_channels;
         new_weights_shape[1] = 1;
-        auto new_weights     = prog.add_instruction(op::reshape{new_weights_shape}, weights);
+        // Make sure weights are contiguous before doing reshape
+        auto cweights    = prog.add_instruction(op::contiguous{}, weights);
+        auto new_weights = prog.add_instruction(op::reshape{new_weights_shape}, cweights);
 
         return prog.add_instruction(op, {args[0], new_weights});
     }
@@ -426,17 +429,21 @@ struct tf_parser
     instruction_ref
     parse_mean(const std::string&, attribute_map attributes, std::vector<instruction_ref> args)
     {
-
         auto axes      = parse_axes(args[1]->eval().get<int32_t>().to_vector());
         bool keep_dims = attributes.at("keep_dims").b();
         std::vector<int32_t> hw_axes{2, 3};
-        if(axes == hw_axes and keep_dims)
+        // check if conditions for GlobalAvgPool are met
+        auto lens = args[0]->get_shape().lens();
+        if(axes == hw_axes and lens.size() == 4)
         {
             op::pooling op{"average"};
-            std::vector<size_t> input_dims{args[0]->get_shape().lens()};
-            op.lengths[0] = input_dims[2];
-            op.lengths[1] = input_dims[3];
-            return prog.add_instruction(op, args.front());
+            op.lengths[0] = lens[2];
+            op.lengths[1] = lens[3];
+            auto l0       = prog.add_instruction(op, args.front());
+            if(keep_dims)
+                return l0;
+            return prog.add_instruction(
+                op::squeeze{std::vector<int64_t>(hw_axes.begin(), hw_axes.end())}, l0);
         }
         MIGRAPHX_THROW("MIGraphX does not support mean outside of GlobalAvgPool transformation");
     }

test/cpu_ops_test.cpp

@@ -3,6 +3,7 @@
 #include <migraphx/literal.hpp>
 #include <migraphx/operators.hpp>
 #include <migraphx/instruction.hpp>
+#include <migraphx/quantization.hpp>
 #include <migraphx/cpu/target.hpp>
 #include <migraphx/verify.hpp>
 #include <migraphx/onnx.hpp>
@@ -1557,6 +1558,34 @@ TEST_CASE(fp16_test)
     EXPECT(migraphx::verify_range(results_vector, gold));
 }
 
+TEST_CASE(fp32_fp16_test)
+{
+    auto create_program = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        std::vector<float> data(2 * 3);
+        std::iota(data.begin(), data.end(), 1.0f);
+        auto l1 = p.add_literal(migraphx::literal(s, data));
+        auto l2 = p.add_literal(migraphx::literal(s, data));
+        p.add_instruction(migraphx::op::add{}, l1, l2);
+        return p;
+    };
+
+    auto test_case = [&](std::vector<std::string>&& op_names) {
+        std::vector<float> gold_res = {2.0, 4.0, 6.0, 8.0, 10.0, 12.0};
+        auto p                      = create_program();
+        migraphx::quantize(p, op_names);
+        p.compile(migraphx::cpu::target{});
+        auto result = p.eval({});
+        std::vector<float> res;
+        result.visit([&](auto output) { res.assign(output.begin(), output.end()); });
+        EXPECT(migraphx::verify_range(res, gold_res));
+    };
+
+    test_case({"all"});
+    test_case({"add"});
+}
+
 TEST_CASE(clip_test)
 {
     migraphx::program p;

test/gpu/miopen.cpp

@@ -10,6 +10,7 @@
 #include <migraphx/type_name.hpp>
 #include <migraphx/verify_args.hpp>
 #include <migraphx/instruction.hpp>
+#include <migraphx/quantization.hpp>
 
 #include <miopen/miopen.h>
 
@@ -1250,22 +1251,6 @@ struct test_contiguous : verify_program<test_contiguous>
     }
 };
 
-struct test_eliminate_contiguous : verify_program<test_eliminate_contiguous>
-{
-    migraphx::program create_program() const
-    {
-        migraphx::program p;
-        migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 5}};
-        auto seq = p.add_parameter("seq", s);
-        std::vector<int64_t> perm{0, 2, 1, 3};
-        auto tran_seq = p.add_instruction(migraphx::op::transpose{perm}, seq);
-        std::vector<int64_t> out_shape{0, 0, -1};
-        p.add_instruction(migraphx::op::reshape{out_shape}, tran_seq);
-
-        return p;
-    }
-};
-
 struct test_transpose : verify_program<test_transpose>
 {
     migraphx::program create_program() const
@@ -3360,4 +3345,70 @@ struct test_logsoftmax_1 : verify_program<test_logsoftmax_1<Axis>>
 template struct test_logsoftmax_1<0>;
 template struct test_logsoftmax_1<1>;
 
+struct test_fp32_fp16_lall : verify_program<test_fp32_fp16_lall>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        std::vector<float> data(2 * 3);
+        std::iota(data.begin(), data.end(), 1.0f);
+        auto l1 = p.add_literal(migraphx::literal(s, data));
+        auto l2 = p.add_parameter("p2", s);
+        p.add_instruction(migraphx::op::add{}, l1, l2);
+        migraphx::quantize(p, {"all"});
+        return p;
+    };
+};
+
+struct test_fp32_fp16_ladd : verify_program<test_fp32_fp16_ladd>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        std::vector<float> data(2 * 3);
+        std::iota(data.begin(), data.end(), 1.0f);
+        auto l1 = p.add_literal(migraphx::literal(s, data));
+        auto l2 = p.add_parameter("p2", s);
+        p.add_instruction(migraphx::op::add{}, l1, l2);
+        migraphx::quantize(p, {"add"});
+        return p;
+    };
+};
+
+struct test_fp32_fp16_add : verify_program<test_fp32_fp16_add>
+{
+    migraphx::program create_program()
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto p1   = p.add_parameter("x", s);
+        auto p2   = p.add_parameter("y", s);
+        auto sum  = p.add_instruction(migraphx::op::add{}, p1, p2);
+        auto diff = p.add_instruction(migraphx::op::sub{}, sum, p2);
+        p.add_instruction(migraphx::op::add{}, diff, p1);
+        migraphx::quantize(p, {"add"});
+
+        return p;
+    };
+};
+
+struct test_fp32_fp16_sub : verify_program<test_fp32_fp16_sub>
+{
+    migraphx::program create_program()
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto p1   = p.add_parameter("x", s);
+        auto p2   = p.add_parameter("y", s);
+        auto sum  = p.add_instruction(migraphx::op::add{}, p1, p2);
+        auto diff = p.add_instruction(migraphx::op::sub{}, sum, p2);
+        p.add_instruction(migraphx::op::add{}, diff, p1);
+        migraphx::quantize(p, {"sub"});
+
+        return p;
+    };
+};
+
 int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/tf/tf_test.cpp

@@ -136,8 +136,9 @@ TEST_CASE(depthwiseconv_test)
     op.group        = 3;
     auto l2         = p.add_instruction(migraphx::op::transpose{{0, 3, 1, 2}}, l1);
     auto l3         = p.add_instruction(migraphx::op::transpose{{1, 3, 0, 2}}, l2);
-    auto l4         = p.add_instruction(migraphx::op::reshape{{3, 1, 3, 3}}, l3);
-    p.add_instruction(op, l0, l4);
+    auto l4         = p.add_instruction(migraphx::op::contiguous{}, l3);
+    auto l5         = p.add_instruction(migraphx::op::reshape{{3, 1, 3, 3}}, l4);
+    p.add_instruction(op, l0, l5);
     auto prog = migraphx::parse_tf("depthwise_conv_test.pb", true);
 
     EXPECT(p == prog);
@@ -168,6 +169,40 @@ TEST_CASE(matmul_test)
     EXPECT(p == prog);
 }
 
+TEST_CASE(mean_test)
+{
+    migraphx::program p;
+    migraphx::literal l{migraphx::shape{migraphx::shape::int32_type, {2}}, {2, 3}};
+    auto l0 = p.add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
+    p.add_literal(l);
+    p.add_literal(l);
+    migraphx::op::pooling op;
+    op.lengths = {16, 16};
+    auto l3    = p.add_instruction(op, l0);
+    p.add_instruction(migraphx::op::squeeze{{2, 3}}, l3);
+    p.add_instruction(op, l0);
+    auto prog = migraphx::parse_tf("mean_test.pb", false);
+
+    EXPECT(p == prog);
+}
+
+TEST_CASE(mean_test_nhwc)
+{
+    migraphx::program p;
+    migraphx::literal l{migraphx::shape{migraphx::shape::int32_type, {2}}, {1, 2}};
+    auto l0 = p.add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
+    p.add_literal(l);
+    p.add_literal(l);
+    migraphx::op::pooling op;
+    op.lengths = {16, 16};
+    auto l3    = p.add_instruction(op, l0);
+    p.add_instruction(migraphx::op::squeeze{{2, 3}}, l3);
+    p.add_instruction(op, l0);
+    auto prog = migraphx::parse_tf("mean_test_nhwc.pb", true);
+
+    EXPECT(p == prog);
+}
+
 TEST_CASE(mul_test)
 {
     migraphx::program p;

test/type_conversion.cpp

+#include <iostream>
+#include <vector>
+#include <migraphx/literal.hpp>
+#include <migraphx/operators.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/cpu/target.hpp>
+#include <migraphx/verify.hpp>
+#include <migraphx/quantization.hpp>
+#include <migraphx/dead_code_elimination.hpp>
+#include <migraphx/propagate_constant.hpp>
+#include <migraphx/pass_manager.hpp>
+#include <migraphx/onnx.hpp>
+#include "test.hpp"
+#include <migraphx/half.hpp>
+
+TEST_CASE(param_add)
+{
+    auto create_program_float = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto p1 = p.add_parameter("x", s);
+        auto p2 = p.add_parameter("y", s);
+        p.add_instruction(migraphx::op::add{}, p1, p2);
+
+        return p;
+    };
+
+    auto create_program_half = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto p1  = p.add_parameter("x", s);
+        auto hp1 = p.insert_instruction(std::next(p1), migraphx::op::convert{}, p1);
+        auto p2  = p.add_parameter("y", s);
+        auto hp2 = p.insert_instruction(std::next(p2), migraphx::op::convert{}, p2);
+        auto hs  = p.add_instruction(migraphx::op::add{}, hp1, hp2);
+        p.add_instruction(migraphx::op::convert{migraphx::shape::float_type}, hs);
+
+        return p;
+    };
+
+    {
+        auto p1 = create_program_float();
+        auto p2 = create_program_half();
+
+        migraphx::quantize(p1);
+        EXPECT(p1 == p2);
+    }
+
+    {
+        auto p1 = create_program_float();
+        auto p2 = create_program_half();
+
+        migraphx::quantize(p1, {"add"});
+        EXPECT(p1 == p2);
+    }
+}
+
+TEST_CASE(param_add_sub)
+{
+    auto create_program_float = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto p1   = p.add_parameter("x", s);
+        auto p2   = p.add_parameter("y", s);
+        auto sum  = p.add_instruction(migraphx::op::add{}, p1, p2);
+        auto diff = p.add_instruction(migraphx::op::sub{}, sum, p2);
+        p.add_instruction(migraphx::op::add{}, diff, p1);
+
+        return p;
+    };
+
+    auto create_program_half_add = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto p1  = p.add_parameter("x", s);
+        auto hp1 = p.insert_instruction(
+            std::next(p1), migraphx::op::convert{migraphx::shape::half_type}, p1);
+        auto p2  = p.add_parameter("y", s);
+        auto hp2 = p.insert_instruction(
+            std::next(p2), migraphx::op::convert{migraphx::shape::half_type}, p2);
+        auto hsum  = p.add_instruction(migraphx::op::add{}, hp1, hp2);
+        auto sum   = p.add_instruction(migraphx::op::convert{migraphx::shape::float_type}, hsum);
+        auto diff  = p.add_instruction(migraphx::op::sub{}, sum, p2);
+        auto hdiff = p.add_instruction(
+            migraphx::op::convert{migraphx::op::convert{migraphx::shape::half_type}}, diff);
+        auto res = p.add_instruction(migraphx::op::add{}, hdiff, hp1);
+        p.add_instruction(migraphx::op::convert{migraphx::shape::float_type}, res);
+
+        return p;
+    };
+
+    auto create_program_half_sub = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto p1  = p.add_parameter("x", s);
+        auto p2  = p.add_parameter("y", s);
+        auto hp2 = p.insert_instruction(
+            std::next(p2), migraphx::op::convert{migraphx::shape::half_type}, p2);
+        auto sum   = p.add_instruction(migraphx::op::add{}, p1, p2);
+        auto hsum  = p.add_instruction(migraphx::op::convert{migraphx::shape::half_type}, sum);
+        auto hdiff = p.add_instruction(migraphx::op::sub{}, hsum, hp2);
+        auto diff  = p.add_instruction(migraphx::op::convert{migraphx::shape::float_type}, hdiff);
+        p.add_instruction(migraphx::op::add{}, diff, p1);
+
+        return p;
+    };
+
+    auto create_program_half_all = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto p1  = p.add_parameter("x", s);
+        auto hp1 = p.insert_instruction(
+            std::next(p1), migraphx::op::convert{migraphx::shape::half_type}, p1);
+        auto p2  = p.add_parameter("y", s);
+        auto hp2 = p.insert_instruction(
+            std::next(p2), migraphx::op::convert{migraphx::shape::half_type}, p2);
+        auto hsum  = p.add_instruction(migraphx::op::add{}, hp1, hp2);
+        auto hdiff = p.add_instruction(migraphx::op::sub{}, hsum, hp2);
+        auto hres  = p.add_instruction(migraphx::op::add{}, hdiff, hp1);
+        p.add_instruction(migraphx::op::convert{migraphx::shape::float_type}, hres);
+
+        return p;
+    };
+
+    {
+        auto p1 = create_program_float();
+        auto p2 = create_program_half_add();
+
+        migraphx::quantize(p1, {"add"});
+        EXPECT(p1 == p2);
+    }
+
+    {
+        auto p1 = create_program_float();
+        auto p2 = create_program_half_sub();
+
+        migraphx::quantize(p1, {"sub"});
+        EXPECT(p1 == p2);
+    }
+
+    {
+        auto p1 = create_program_float();
+        auto p2 = create_program_half_all();
+
+        migraphx::quantize(p1);
+        migraphx::run_passes(p1, {migraphx::dead_code_elimination{}});
+
+        EXPECT(p1 == p2);
+    }
+}
+
+TEST_CASE(literal_add)
+{
+    auto create_program_float = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        std::vector<float> data(2 * 3);
+        std::iota(data.begin(), data.end(), 1.0f);
+        auto l1 = p.add_literal(migraphx::literal(s, data));
+        auto l2 = p.add_literal(migraphx::literal(s, data));
+        p.add_instruction(migraphx::op::add{}, l1, l2);
+
+        return p;
+    };
+
+    auto create_program_half = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::half_type, {2, 3}};
+        std::vector<migraphx::half> data(2 * 3);
+        std::iota(data.begin(), data.end(), 1.0f);
+        auto l1 = p.add_literal(migraphx::literal(s, data));
+        auto l2 = p.add_literal(migraphx::literal(s, data));
+        auto hs = p.add_instruction(migraphx::op::add{}, l1, l2);
+        p.add_instruction(migraphx::op::convert{migraphx::shape::float_type}, hs);
+
+        return p;
+    };
+
+    {
+        auto p1 = create_program_float();
+        auto p2 = create_program_half();
+
+        migraphx::quantize(p1, {"all"});
+        migraphx::run_passes(p1,
+                             {migraphx::propagate_constant{}, migraphx::dead_code_elimination{}});
+        migraphx::run_passes(p2,
+                             {migraphx::propagate_constant{}, migraphx::dead_code_elimination{}});
+
+        EXPECT(p1 == p2);
+    }
+
+    {
+        auto p1 = create_program_float();
+        auto p2 = create_program_half();
+
+        migraphx::quantize(p1, {"add"});
+        migraphx::run_passes(p1,
+                             {migraphx::propagate_constant{}, migraphx::dead_code_elimination{}});
+        migraphx::run_passes(p2,
+                             {migraphx::propagate_constant{}, migraphx::dead_code_elimination{}});
+        EXPECT(p1 == p2);
+    }
+}
+
+int main(int argc, const char* argv[]) { test::run(argc, argv); }