Merge branch 'adjust_gpu_allocation' into ins_fp32_fp16

src/include/migraphx/op/unary.hpp

@@ -21,7 +21,7 @@ struct unary
     shape compute_shape(std::vector<shape> inputs) const
     {
         check_shapes{inputs}.has(1);
-        return {inputs.at(0).type(), inputs.at(0).lens()};
+        return inputs.at(0);
     }
 };
 

src/targets/cpu/lowering.cpp

@@ -593,26 +593,13 @@ struct cpu_unary
 {
     Op op;
     std::string name() const { return op.name(); }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        return {inputs.front().type(), inputs.front().lens()};
-    }
-
+    shape compute_shape(const std::vector<shape>& inputs) const { return inputs.front(); }
     argument compute(context&, const shape& output_shape, std::vector<argument> args) const
     {
         argument result{output_shape};
         result.visit([&](auto output) {
             args[0].visit([&](auto input) {
-                if(input.get_shape().packed())
-                {
-                    std::transform(input.begin(), input.end(), output.begin(), op.fcn());
-                }
-                else
-                {
-                    shape_for_each(output.get_shape(), [&](const auto& idx) {
-                        output(idx.begin(), idx.end()) = op.fcn()(input(idx.begin(), idx.end()));
-                    });
-                }
+                std::transform(input.begin(), input.end(), output.begin(), op.fcn());
             });
         });
 
@@ -798,11 +785,7 @@ struct cpu_binary
 {
     Op op;
     std::string name() const { return op.name(); }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        // operator will generate standard output shape
-        return {inputs.front().type(), inputs.front().lens()};
-    }
+    shape compute_shape(const std::vector<shape>& inputs) const { return inputs.front(); }
     argument compute(context&, const shape& output_shape, std::vector<argument> args) const
     {
         argument result{output_shape};

src/targets/gpu/abs.cpp

@@ -8,7 +8,7 @@ namespace gpu {
 shape miopen_abs::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_abs::compute(context& ctx,

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).type(), inputs.at(0).lens()};
+        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).type(), inputs.at(0).lens()};
+        return inputs.at(2);
     }
 
     argument compute(context& ctx, const shape&, const std::vector<argument>& args) const

src/targets/gpu/lrn.cpp

@@ -7,7 +7,7 @@ namespace gpu {
 
 shape miopen_lrn::compute_shape(const std::vector<shape>& inputs) const
 {
-    check_shapes{inputs, *this}.has(2).standard();
+    check_shapes{inputs, *this}.has(2).not_broadcasted();
     return inputs.at(1);
 }
 

src/targets/gpu/relu.cpp

@@ -7,7 +7,7 @@ namespace gpu {
 
 shape miopen_relu::compute_shape(const std::vector<shape>& inputs) const
 {
-    check_shapes{inputs, *this}.has(2).not_broadcasted().not_transposed();
+    check_shapes{inputs, *this}.has(2).not_broadcasted();
     return inputs.at(1);
 }
 

src/targets/gpu/tanh.cpp

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

@@ -327,6 +327,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 +687,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