Merge branch 'develop' into pack_op

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

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_CONVOLUTION_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/convolution.hpp>
 #include <migraphx/gpu/miopen.hpp>
 
 namespace migraphx {

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_GATHER_HPP
 
 #include <migraphx/shape.hpp>
+#include <migraphx/op/gather.hpp>
 #include <migraphx/gpu/miopen.hpp>
 
 namespace migraphx {

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

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_GEMM_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/dot.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

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

 #ifndef MIGRAPHX_GUARD_MIGRAPHLIB_HIP_HPP
 #define MIGRAPHX_GUARD_MIGRAPHLIB_HIP_HPP
 
-#include <migraphx/operators.hpp>
 #include <migraphx/config.hpp>
+#include <migraphx/argument.hpp>
+#include <migraphx/check_shapes.hpp>
 #include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
+struct context;
+
 argument allocate_gpu(const shape& s, bool host = false);
 
 argument to_gpu(const argument& arg, bool host = false);

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

@@ -4,7 +4,7 @@
 #include <migraphx/gpu/lowering.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/logsoftmax.hpp>
 #include <migraphx/generate.hpp>
 #include <migraphx/shape_for_each.hpp>
 #include <migraphx/config.hpp>

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

@@ -2,7 +2,9 @@
 #define MIGRAPHX_GUARD_MIGRAPHLIB_MIOPEN_HPP
 
 #include <migraphx/manage_ptr.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/convolution.hpp>
+#include <migraphx/op/pooling.hpp>
+#include <migraphx/op/lrn.hpp>
 #include <miopen/miopen.h>
 #include <migraphx/config.hpp>
 

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

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_PAD_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/pad.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

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

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_POOLING_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/pooling.hpp>
 #include <migraphx/gpu/miopen.hpp>
 
 namespace migraphx {

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

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_SOFTMAX_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/softmax.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/logsoftmax.cpp

 #include <migraphx/gpu/logsoftmax.hpp>
 #include <migraphx/gpu/device/logsoftmax.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/logsoftmax.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/gpu/miopen.hpp>
 #include <utility>

src/targets/gpu/target.cpp

@@ -20,6 +20,7 @@
 #include <migraphx/eliminate_identity.hpp>
 #include <migraphx/gpu/concat_gpu_opt.hpp>
 #include <migraphx/gpu/schedule_model.hpp>
+#include <migraphx/eliminate_pad.hpp>
 #include <migraphx/schedule.hpp>
 
 namespace migraphx {
@@ -36,6 +37,8 @@ std::vector<pass> target::get_passes(migraphx::context& gctx) const
     {
         dead_code_elimination{},
         eliminate_identity{},
+        eliminate_pad{},
+        dead_code_elimination{},
         fwd_conv_batchnorm_rewrite{},
         dead_code_elimination{},
         rewrite_rnn{},

test/auto_contiguous_test.cpp

 #include <migraphx/auto_contiguous.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/transpose.hpp>
+#include <migraphx/op/broadcast.hpp>
 #include <migraphx/instruction.hpp>
 #include <basic_ops.hpp>
 #include <test.hpp>

test/common_subexpression_elimination_test.cpp

 #include <migraphx/common_subexpression_elimination.hpp>
 #include <migraphx/dead_code_elimination.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/add.hpp>
 #include <basic_ops.hpp>
 #include <test.hpp>
 

test/constant_propagate_test.cpp

 #include <migraphx/constant_propagate.hpp>
 #include <migraphx/dead_code_elimination.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/add.hpp>
 #include <basic_ops.hpp>
 #include <test.hpp>
 

test/cpu_rnn_ops_test.cpp

 #include <iostream>
 #include <vector>
 #include <migraphx/literal.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/rnn.hpp>
+#include <migraphx/op/gru.hpp>
+#include <migraphx/op/lstm.hpp>
+#include <migraphx/op/rnn_last_output.hpp>
+#include <migraphx/op/rnn_last_cell_output.hpp>
+#include <migraphx/op/abnormal_ops.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/cpu/target.hpp>
 #include <migraphx/verify.hpp>

test/dead_code_elimination_test.cpp

 #include <migraphx/dead_code_elimination.hpp>
 #include <basic_ops.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/abnormal_ops.hpp>
+#include <migraphx/op/add.hpp>
+#include <migraphx/op/identity.hpp>
 #include <test.hpp>
 
 struct dce_target
@@ -129,4 +131,55 @@ TEST_CASE(undefined_test)
     EXPECT(result != migraphx::literal{4});
 }
 
+TEST_CASE(duplicate_args1)
+{
+    migraphx::program p;
+
+    auto l0 = p.add_literal(0);
+    auto l3 = p.add_literal(3);
+    p.add_instruction(migraphx::op::add{}, l3, l3);
+    p.add_instruction(migraphx::op::identity{}, l0);
+    auto count = std::distance(p.begin(), p.end());
+    p.compile(dce_target{});
+    EXPECT(std::distance(p.begin(), p.end()) != count);
+    EXPECT(std::distance(p.begin(), p.end()) == 2);
+    auto result = p.eval({});
+    EXPECT(result == migraphx::literal{0});
+}
+
+TEST_CASE(duplicate_args2)
+{
+    migraphx::program p;
+
+    auto l0   = p.add_literal(0);
+    auto l3   = p.add_literal(3);
+    auto sum1 = p.add_instruction(migraphx::op::add{}, l0, l3);
+    p.add_instruction(migraphx::op::add{}, sum1, l3);
+    p.add_instruction(migraphx::op::identity{}, l0);
+    auto count = std::distance(p.begin(), p.end());
+    p.compile(dce_target{});
+    EXPECT(std::distance(p.begin(), p.end()) != count);
+    EXPECT(std::distance(p.begin(), p.end()) == 2);
+    auto result = p.eval({});
+    EXPECT(result == migraphx::literal{0});
+}
+
+TEST_CASE(duplicate_args3)
+{
+    migraphx::program p;
+
+    auto l0   = p.add_literal(0);
+    auto l3   = p.add_literal(3);
+    auto sum1 = p.add_instruction(migraphx::op::add{}, l0, l3);
+    auto sum2 = p.add_instruction(migraphx::op::add{}, l0, sum1);
+    p.add_instruction(migraphx::op::add{}, sum2, l3);
+    p.add_instruction(migraphx::op::identity{}, l0);
+    auto count = std::distance(p.begin(), p.end());
+    p.compile(dce_target{});
+    EXPECT(std::distance(p.begin(), p.end()) != count);
+    EXPECT(std::distance(p.begin(), p.end()) == 2);
+    auto result = p.eval({});
+    EXPECT(result == migraphx::literal{0});
+}
+
 int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/eliminate_allocation_test.cpp

 #include <migraphx/eliminate_allocation.hpp>
 #include <migraphx/dead_code_elimination.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/check_shapes.hpp>
+#include <migraphx/argument.hpp>
 #include <basic_ops.hpp>
 #include <test.hpp>
 

test/eliminate_concat_test.cpp

 #include <migraphx/eliminate_concat.hpp>
 #include <migraphx/dead_code_elimination.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/concat.hpp>
+#include <migraphx/op/load.hpp>
+#include <migraphx/op/identity.hpp>
 #include <basic_ops.hpp>
 #include <test.hpp>
 

test/eliminate_contiguous_test.cpp

 #include <migraphx/eliminate_contiguous.hpp>
 #include <migraphx/dead_code_elimination.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/transpose.hpp>
+#include <migraphx/op/contiguous.hpp>
 #include <basic_ops.hpp>
 #include <test.hpp>
 

test/eliminate_identity_test.cpp

@@ -2,7 +2,7 @@
 #include <migraphx/eliminate_identity.hpp>
 #include <migraphx/instruction.hpp>
 #include <basic_ops.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/identity.hpp>
 #include <test.hpp>
 
 struct eliminate_identity_target