Merge branch 'develop' into reshape

src/targets/gpu/include/migraphx/gpu/oper.hpp

@@ -45,7 +45,7 @@ struct unary_device : oper<Derived>
     shape compute_shape(const std::vector<shape>& inputs) const
     {
         check_shapes{inputs, *this}.has(2);
-        return inputs.at(0);
+        return inputs.at(1);
     }
 
     argument compute(context& ctx, const shape&, const std::vector<argument>& args) const
@@ -63,7 +63,7 @@ struct binary_device : oper<Derived>
     shape compute_shape(const std::vector<shape>& inputs) const
     {
         check_shapes{inputs, *this}.has(3);
-        return inputs.at(0);
+        return inputs.at(2);
     }
 
     argument compute(context& ctx, const shape&, const std::vector<argument>& args) const

src/targets/gpu/tanh.cpp

@@ -8,7 +8,7 @@ namespace gpu {
 shape miopen_tanh::compute_shape(const std::vector<shape>& inputs) const
 {
     check_shapes{inputs, *this}.has(2).not_broadcasted();
-    return inputs.at(1);
+    return inputs.at(0);
 }
 
 argument miopen_tanh::compute(context& ctx,

src/targets/gpu/target.cpp

@@ -11,7 +11,7 @@
 #include <migraphx/dead_code_elimination.hpp>
 #include <migraphx/simplify_reshapes.hpp>
 #include <migraphx/simplify_algebra.hpp>
-#include <migraphx/constant_propagate.hpp>
+#include <migraphx/propagate_constant.hpp>
 #include <migraphx/eliminate_contiguous.hpp>
 #include <migraphx/common_subexpression_elimination.hpp>
 #include <migraphx/fwd_conv_batchnorm_rewrite.hpp>
@@ -20,6 +20,7 @@
 #include <migraphx/eliminate_identity.hpp>
 #include <migraphx/gpu/concat_gpu_opt.hpp>
 #include <migraphx/gpu/schedule_model.hpp>
+#include <migraphx/gpu/adjust_allocation.hpp>
 #include <migraphx/eliminate_pad.hpp>
 #include <migraphx/schedule.hpp>
 
@@ -47,7 +48,7 @@ std::vector<pass> target::get_passes(migraphx::context& gctx) const
         //dead_code_elimination{},
         simplify_algebra{},
         dead_code_elimination{},
-        constant_propagate{},
+        propagate_constant{},
         dead_code_elimination{},
         auto_contiguous{},
         simplify_reshapes{},
@@ -57,6 +58,8 @@ std::vector<pass> target::get_passes(migraphx::context& gctx) const
         dead_code_elimination{},
         eliminate_contiguous{},
         dead_code_elimination{},
+        adjust_allocation{},
+        dead_code_elimination{},
         fuse_ops{&ctx},
         dead_code_elimination{},
         write_literals{&ctx},

test/gpu/adjust_allocation.cpp

+#include <migraphx/gpu/adjust_allocation.hpp>
+#include <migraphx/gpu/target.hpp>
+#include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/context.hpp>
+#include <migraphx/dead_code_elimination.hpp>
+#include <migraphx/auto_contiguous.hpp>
+#include <migraphx/eliminate_contiguous.hpp>
+#include <migraphx/iterator_for.hpp>
+#include <migraphx/op/add.hpp>
+#include <migraphx/op/transpose.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/pass_manager.hpp>
+#include <migraphx/op/tanh.hpp>
+#include <basic_ops.hpp>
+#include <test.hpp>
+
+struct lowering_target
+{
+    std::string name() const { return "gpu::lowering"; }
+    std::vector<migraphx::pass> get_passes(migraphx::context& gctx) const
+    {
+        auto& ctx = migraphx::any_cast<migraphx::gpu::context>(gctx);
+        return {migraphx::auto_contiguous{},
+                migraphx::gpu::lowering{ctx},
+                migraphx::dead_code_elimination{},
+                migraphx::eliminate_contiguous{},
+                migraphx::dead_code_elimination{}};
+    }
+    migraphx::gpu::context get_context() const { return migraphx::gpu::context{}; }
+};
+
+TEST_CASE(tanh_shape)
+{
+    auto create_program = [] {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto x   = p.add_parameter("x", s);
+        auto tx  = p.add_instruction(migraphx::op::transpose{{1, 0}}, x);
+        auto txh = p.add_instruction(migraphx::op::tanh{}, tx);
+        p.add_instruction(migraphx::op::add{}, txh, txh);
+
+        return p;
+    };
+
+    auto p1 = create_program();
+    auto p2 = create_program();
+    EXPECT(p1 == p2);
+
+    p1.compile(lowering_target{});
+    p2.compile(lowering_target());
+
+    EXPECT(p1 == p2);
+
+    for(auto ins : iterator_for(p1))
+    {
+        if(ins->name() == "hip::allocate")
+        {
+            migraphx::shape wrong_s{migraphx::shape::float_type, {3, 2}, {1, 3}};
+            ins->replace(wrong_s);
+        }
+    }
+    EXPECT(p1 != p2);
+
+    migraphx::run_passes(p2,
+                         {migraphx::gpu::adjust_allocation{}, migraphx::dead_code_elimination{}});
+    EXPECT(p1 == p2);
+}
+
+int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/gpu/miopen.cpp

@@ -236,8 +236,7 @@ struct test_exp : verify_program<test_exp>
     {
         migraphx::program p;
         migraphx::shape s{migraphx::shape::float_type, {6}};
-        std::vector<float> data{0.1f, 0.2f, 1.f, 2.f, 0.6f, 10.f};
-        auto x = p.add_literal(s, data);
+        auto x = p.add_instruction(migraphx::op::abs{}, p.add_parameter("x", s));
         p.add_instruction(migraphx::op::exp{}, x);
         return p;
     }
@@ -249,8 +248,7 @@ struct test_log : verify_program<test_log>
     {
         migraphx::program p;
         migraphx::shape s{migraphx::shape::float_type, {6}};
-        std::vector<float> data{0.1f, 0.2f, 1.f, 2.f, 0.6f, 100.f};
-        auto x = p.add_literal(s, data);
+        auto x = p.add_instruction(migraphx::op::abs{}, p.add_parameter("x", s));
         p.add_instruction(migraphx::op::log{}, x);
         return p;
     }
@@ -327,6 +325,19 @@ struct test_tanh : verify_program<test_tanh>
     }
 };
 
+struct test_trans_tanh : verify_program<test_trans_tanh>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto x  = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
+        auto tx = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, x);
+        auto tanhx = p.add_instruction(migraphx::op::tanh{}, tx);
+        p.add_instruction(migraphx::op::add{}, tanhx, tanhx);
+        return p;
+    }
+};
+
 struct test_asin : verify_program<test_asin>
 {
     migraphx::program create_program() const
@@ -674,6 +685,19 @@ struct test_abs : verify_program<test_abs>
     }
 };
 
+struct test_trans_abs : verify_program<test_trans_abs>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto x  = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
+        auto tx = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, x);
+        auto tanhx = p.add_instruction(migraphx::op::abs{}, tx);
+        p.add_instruction(migraphx::op::add{}, tanhx, tanhx);
+        return p;
+    }
+};
+
 struct test_leaky_relu : verify_program<test_leaky_relu>
 {
     migraphx::program create_program() const

test/print_graph_test.cpp

+#include <migraphx/program.hpp>
+#include <migraphx/ranges.hpp>
+#include <sstream>
+#include "test.hpp"
+#include <basic_ops.hpp>
+
+migraphx::program create_program()
+{
+    migraphx::program p;
+
+    auto x = p.add_parameter("x", {migraphx::shape::int64_type});
+    auto y = p.add_parameter("y", {migraphx::shape::int64_type});
+
+    auto sum = p.add_instruction(sum_op{}, x, y);
+    auto one = p.add_literal(1);
+    p.add_instruction(sum_op{}, sum, one);
+
+    return p;
+}
+
+TEST_CASE(basic_graph_test)
+{
+    migraphx::program p = create_program();
+    std::stringstream ss;
+    p.print_graph(ss);
+    std::string test = ss.str();
+    EXPECT(migraphx::contains(test, "digraph"));
+    EXPECT(migraphx::contains(test, "rankdir=LR"));
+    EXPECT(migraphx::contains(test, "\"@0\"[label=\"@literal\"]"));
+    EXPECT(migraphx::contains(test, "\"y\"[label=\"@param:y\"]"));
+    EXPECT(migraphx::contains(test, "\"x\"[label=\"@param:x\"]"));
+    EXPECT(migraphx::contains(test, "\"@3\"[label=\"sum\"]"));
+    EXPECT(migraphx::contains(test, "\"@4\"[label=\"sum\"]"));
+    EXPECT(migraphx::contains(test, "\"x\" -> \"@3\""));
+    EXPECT(migraphx::contains(test, "\"y\" -> \"@3\""));
+    EXPECT(migraphx::contains(test, "\"@3\" -> \"@4\""));
+    EXPECT(migraphx::contains(test, "\"@0\" -> \"@4\""));
+    EXPECT(migraphx::contains(test, "[label=\"int64_type, {1}, {0}\"]"));
+}
+
+int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/constant_propagate_test.cpp → test/propagate_constant_test.cpp

-#include <migraphx/constant_propagate.hpp>
+#include <migraphx/propagate_constant.hpp>
 #include <migraphx/dead_code_elimination.hpp>
 #include <migraphx/op/add.hpp>
+#include <migraphx/op/scalar.hpp>
+#include <migraphx/op/mul.hpp>
 #include <basic_ops.hpp>
 #include <test.hpp>
 
@@ -9,12 +11,12 @@ struct const_prop_target
     std::string name() const { return "const_prop"; }
     std::vector<migraphx::pass> get_passes(migraphx::context&) const
     {
-        return {migraphx::constant_propagate{}, migraphx::dead_code_elimination{}};
+        return {migraphx::propagate_constant{}, migraphx::dead_code_elimination{}};
     }
     migraphx::context get_context() const { return {}; }
 };
 
-TEST_CASE(const_add1)
+TEST_CASE(const_add)
 {
     migraphx::program p1;
     auto one = p1.add_literal(1);
@@ -29,7 +31,7 @@ TEST_CASE(const_add1)
     EXPECT(p1 == p2);
 }
 
-TEST_CASE(const_add2)
+TEST_CASE(const_add_parameter)
 {
     migraphx::program p1;
     auto one = p1.add_parameter("one", {migraphx::shape::int32_type, {1}});
@@ -44,7 +46,7 @@ TEST_CASE(const_add2)
     EXPECT(p1 != p2);
 }
 
-TEST_CASE(const_add3)
+TEST_CASE(const_multiadd)
 {
     migraphx::program p1;
     auto one  = p1.add_literal(1);
@@ -60,4 +62,54 @@ TEST_CASE(const_add3)
     EXPECT(p1 == p2);
 }
 
+TEST_CASE(const_add_mul)
+{
+    migraphx::program p1;
+    auto one  = p1.add_literal(1);
+    auto two  = p1.add_literal(2);
+    auto mul  = p1.add_instruction(migraphx::op::mul{}, two, two);
+    auto sum1 = p1.add_instruction(migraphx::op::add{}, one, mul);
+    auto sum2 = p1.add_instruction(migraphx::op::add{}, sum1, two);
+    p1.add_instruction(pass_op{}, sum2);
+    p1.compile(const_prop_target{});
+
+    migraphx::program p2;
+    auto total = p2.add_literal(7);
+    p2.add_instruction(pass_op{}, total);
+    EXPECT(p1 == p2);
+}
+
+TEST_CASE(const_add_scalar)
+{
+    migraphx::program p1;
+    auto one = p1.add_instruction(migraphx::op::scalar{{2, 2}}, p1.add_literal(1));
+    auto two = p1.add_instruction(migraphx::op::scalar{{2, 2}}, p1.add_literal(2));
+    auto sum = p1.add_instruction(migraphx::op::add{}, one, two);
+    p1.add_instruction(pass_op{}, sum);
+    p1.compile(const_prop_target{});
+
+    migraphx::program p2;
+    auto total =
+        p2.add_literal(migraphx::literal{{migraphx::shape::int32_type, {2, 2}}, {3, 3, 3, 3}});
+    p2.add_instruction(pass_op{}, total);
+    EXPECT(p1 == p2);
+}
+
+TEST_CASE(const_scalar)
+{
+    migraphx::program p1;
+    {
+        auto one = p1.add_instruction(migraphx::op::scalar{{2, 2}}, p1.add_literal(1));
+        p1.add_instruction(pass_op{}, one);
+    }
+    p1.compile(const_prop_target{});
+
+    migraphx::program p2;
+    {
+        auto one = p2.add_instruction(migraphx::op::scalar{{2, 2}}, p2.add_literal(1));
+        p2.add_instruction(pass_op{}, one);
+    }
+    EXPECT(p1 == p2);
+}
+
 int main(int argc, const char* argv[]) { test::run(argc, argv); }