Merge branch 'develop' of https://github.com/ROCmSoftwarePlatform/AMDMIGraphX into int8_quantize

src/eliminate_contiguous.cpp

@@ -4,6 +4,8 @@
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/ranges.hpp>
 #include <migraphx/stringutils.hpp>
+#include <migraphx/op/contiguous.hpp>
+#include <migraphx/op/identity.hpp>
 #include <utility>
 
 namespace migraphx {
@@ -82,6 +84,14 @@ void eliminate_contiguous::apply(program& p) const
                 {
                     instruction::replace_argument(ins, arg, prev);
                 }
+                else if(prev->can_eval())
+                {
+                    auto c = op::contiguous{};
+                    auto r = c.compute(c.compute_shape({prev->get_shape()}), {prev->eval()});
+
+                    auto l = p.add_literal(r.get_shape(), r.data());
+                    p.replace_instruction(arg, l);
+                }
             }
         }
     }

src/onnx/onnx.cpp

@@ -323,9 +323,13 @@ struct onnx_parser
         if(contains(attributes, "pads"))
         {
             if(contains(attributes, "auto_pad"))
+            {
+                auto s = attributes["auto_pad"].s();
+                if(contains(attributes, "pads") and to_upper(s) != "NOTSET")
                 {
                     MIGRAPHX_THROW("auto_pad and padding cannot be specified simultaneously");
                 }
+            }
             std::vector<std::int64_t> padding;
             copy(attributes["pads"].ints(), std::back_inserter(padding));
             if(padding.size() != 4)
@@ -372,7 +376,7 @@ struct onnx_parser
         if(args.size() == 3)
         {
             uint64_t axis = 1;
-            auto l1       = prog.add_instruction(op, args[0], args[1]);
+            auto l1       = prog.add_instruction(op, l0, args[1]);
             auto l2 = prog.add_instruction(op::broadcast{axis, l1->get_shape().lens()}, args[2]);
             return prog.add_instruction(op::add{}, l1, l2);
         }

src/targets/gpu/lowering.cpp

@@ -87,6 +87,7 @@ struct miopen_apply
     void init()
     {
         this->last = instruction::get_output_alias(std::prev(prog->end()));
+
         add_miopen_simple_op<miopen_abs>("abs", make_abs);
 
         add_miopen_extend_op<miopen_leaky_relu, op::leaky_relu>("leaky_relu", make_leaky_relu);

test/eliminate_contiguous_test.cpp

@@ -22,7 +22,7 @@ struct eliminate_contiguous_target
 TEST_CASE(standard_op)
 {
     migraphx::program p;
-    auto l = p.add_literal(get_2x2());
+    auto l = p.add_parameter("x", {migraphx::shape::float_type, {2, 2}});
     auto t = p.add_instruction(migraphx::op::transpose{{1, 0}}, l);
     auto c = p.add_instruction(migraphx::op::contiguous{}, t);
     p.add_instruction(pass_standard_op{}, c);
@@ -31,18 +31,40 @@ TEST_CASE(standard_op)
     EXPECT(std::distance(p.begin(), p.end()) == count);
 }
 
-TEST_CASE(non_standard_op)
+TEST_CASE(standard_op_const)
 {
     migraphx::program p;
     auto l = p.add_literal(get_2x2());
     auto t = p.add_instruction(migraphx::op::transpose{{1, 0}}, l);
     auto c = p.add_instruction(migraphx::op::contiguous{}, t);
+    p.add_instruction(pass_standard_op{}, c);
+    p.compile(eliminate_contiguous_target{});
+    EXPECT(std::distance(p.begin(), p.end()) == 2);
+}
+
+TEST_CASE(non_standard_op)
+{
+    migraphx::program p;
+    auto l = p.add_parameter("x", {migraphx::shape::float_type, {2, 2}});
+    auto t = p.add_instruction(migraphx::op::transpose{{1, 0}}, l);
+    auto c = p.add_instruction(migraphx::op::contiguous{}, t);
     p.add_instruction(pass_op{}, c);
     auto count = std::distance(p.begin(), p.end());
     p.compile(eliminate_contiguous_target{});
     EXPECT(std::distance(p.begin(), p.end()) == count);
 }
 
+TEST_CASE(non_standard_op_const)
+{
+    migraphx::program p;
+    auto l = p.add_literal(get_2x2());
+    auto t = p.add_instruction(migraphx::op::transpose{{1, 0}}, l);
+    auto c = p.add_instruction(migraphx::op::contiguous{}, t);
+    p.add_instruction(pass_op{}, c);
+    p.compile(eliminate_contiguous_target{});
+    EXPECT(std::distance(p.begin(), p.end()) == 2);
+}
+
 TEST_CASE(transpose_gemm)
 {
     migraphx::program p;
@@ -59,7 +81,7 @@ TEST_CASE(transpose_gemm)
 TEST_CASE(transpose_standard_op)
 {
     migraphx::program p;
-    auto l  = p.add_literal(get_2x2());
+    auto l  = p.add_parameter("x", {migraphx::shape::float_type, {2, 2}});
     auto t  = p.add_instruction(migraphx::op::transpose{{1, 0}}, l);
     auto c  = p.add_instruction(migraphx::op::contiguous{}, t);
     auto sn = p.add_instruction(migraphx::op::sin{}, c);
@@ -69,6 +91,18 @@ TEST_CASE(transpose_standard_op)
     EXPECT(std::distance(p.begin(), p.end()) == count);
 }
 
+TEST_CASE(transpose_standard_op_const)
+{
+    migraphx::program p;
+    auto l  = p.add_literal(get_2x2());
+    auto t  = p.add_instruction(migraphx::op::transpose{{1, 0}}, l);
+    auto c  = p.add_instruction(migraphx::op::contiguous{}, t);
+    auto sn = p.add_instruction(migraphx::op::sin{}, c);
+    p.add_instruction(pass_standard_op{}, sn);
+    p.compile(eliminate_contiguous_target{});
+    EXPECT(std::distance(p.begin(), p.end()) == 3);
+}
+
 TEST_CASE(no_packed_unary_op)
 {
     migraphx::program p;

test/onnx/onnx_test.cpp

@@ -8,6 +8,11 @@
 #include <migraphx/onnx.hpp>
 #include "test.hpp"
 
+TEST_CASE(conv_autopad_fail_test)
+{
+    EXPECT(test::throws([&] { migraphx::parse_onnx("conv_autopad_fail_test.onnx"); }));
+}
+
 TEST_CASE(pytorch_conv_bias_test)
 {
     migraphx::program p;