continue testing

src/onnx/onnx.cpp

@@ -123,45 +123,6 @@ struct onnx_parser
                 }
                 return prog.add_instruction(x, args);
             }
-<<<<<<< HEAD
-=======
-            else if(args[0]->get_shape() != args[1]->get_shape())
-            {
-                // Example:
-                // s0 = (3,2,4,5) and s1 = (2,1,1)
-                //
-                // In this case we need to broadcast (:,1,1) portion of
-                // s1 plus broadcast the 1st dimension of s1
-                // giving output_lens = (3,2,4,5)
-                //
-                // Another example:
-                // s0 = (3,2,1,5) and s1 = (2,7,5)
-                // In this case we need to broadcast the (:,:,1:,:) axis
-                // of s0 plus the 1st dimension of s1 giving
-                // output_lens = (3,2,7,5)
-                //
-                // Get lengths for both arguments
-                const std::vector<std::size_t>* s0 = &args[0]->get_shape().lens();
-                const std::vector<std::size_t>* s1 = &args[1]->get_shape().lens();
-
-                // Make sure s0 is the smaller size
-                if(s0->size() > s1->size())
-                    std::swap(s0, s1);
-
-                // Copy the larger vector to output_lens
-                std::vector<std::size_t> output_lens = *s1;
-                auto offset                          = s1->size() - s0->size();
-                std::transform(s0->begin(),
-                               s0->end(),
-                               s1->begin() + offset,
-                               output_lens.begin() + offset,
-                               [](auto a, auto b) { return std::max(a, b); });
-
-                auto l0 = prog.add_instruction(op::multibroadcast{output_lens}, args[0]);
-                auto l1 = prog.add_instruction(op::multibroadcast{output_lens}, args[1]);
-                return prog.add_instruction(x, l0, l1);
-            }
->>>>>>> 84e7335eb6088f9918dcf86f9fc1b58ef27c3360
             else
             {
                 return add_broadcastable_binary_op(args[0], args[1], x);

src/targets/gpu/CMakeLists.txt

@@ -12,13 +12,13 @@ endif()
 
 add_library(migraphx_device
     device/add.cpp
+    device/max.cpp
+    device/min.cpp
     device/sin.cpp
     device/add_relu.cpp
     device/contiguous.cpp
     device/mul.cpp
     device/concat.cpp
-    device/max.cpp
-    device/min.cpp
 )
 set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
 rocm_clang_tidy_check(migraphx_device)
@@ -42,6 +42,8 @@ add_library(migraphx_gpu
     leaky_relu.cpp
     add.cpp
     sin.cpp
+    max.cpp
+    min.cpp
     mul.cpp
     batchnorm.cpp
     write_literals.cpp
@@ -50,8 +52,6 @@ add_library(migraphx_gpu
     tanh.cpp
     abs.cpp
     elu.cpp
-    max.cpp
-    min.cpp
 )
 set_target_properties(migraphx_gpu PROPERTIES EXPORT_NAME gpu)
 rocm_clang_tidy_check(migraphx_gpu)

src/targets/gpu/include/migraphx/gpu/device/max.hpp

@@ -13,12 +13,6 @@ namespace device {
 
 void max(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2);
 
-void max(hipStream_t stream,
-         const argument& result,
-         const argument& arg1,
-         const argument& arg2,
-         const argument& arg3);
-
 } // namespace device
 } // namespace gpu
 } // namespace MIGRAPH_INLINE_NS

src/targets/gpu/include/migraphx/gpu/device/min.hpp

@@ -13,12 +13,6 @@ namespace device {
 
 void min(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2);
 
-void min(hipStream_t stream,
-         const argument& result,
-         const argument& arg1,
-         const argument& arg2,
-         const argument& arg3);
-
 } // namespace device
 } // namespace gpu
 } // namespace MIGRAPH_INLINE_NS

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

@@ -7,7 +7,6 @@
 #include <migraphx/operators.hpp>
 #include <migraphx/generate.hpp>
 #include <migraphx/shape_for_each.hpp>
-#include <migraphx/config.hpp>
 #include <migraphx/gpu/miopen.hpp>
 #include <migraphx/gpu/hip.hpp>
 #include <migraphx/dfor.hpp>
@@ -16,6 +15,7 @@
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/gpu/rocblas.hpp>
 #include <migraphx/gpu/context.hpp>
+#include <migraphx/config.hpp>
 #include <utility>
 
 namespace migraphx {

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

@@ -7,7 +7,6 @@
 #include <migraphx/operators.hpp>
 #include <migraphx/generate.hpp>
 #include <migraphx/shape_for_each.hpp>
-#include <migraphx/config.hpp>
 #include <migraphx/gpu/miopen.hpp>
 #include <migraphx/gpu/hip.hpp>
 #include <migraphx/dfor.hpp>
@@ -16,6 +15,7 @@
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/gpu/rocblas.hpp>
 #include <migraphx/gpu/context.hpp>
+#include <migraphx/config.hpp>
 #include <utility>
 
 namespace migraphx {

src/targets/gpu/max.cpp

 #include <migraphx/gpu/max.hpp>
 #include <migraphx/operators.hpp>
 #include <migraphx/manage_ptr.hpp>
+#include <migraphx/config.hpp>
 #include <migraphx/gpu/miopen.hpp>
 #include <utility>
 

src/targets/gpu/min.cpp

 #include <migraphx/gpu/min.hpp>
 #include <migraphx/operators.hpp>
 #include <migraphx/manage_ptr.hpp>
+#include <migraphx/config.hpp>
 #include <migraphx/gpu/miopen.hpp>
 #include <utility>
 

test/onnx/max_test.onnx

+max-example:e
+
+
0
+
1
+
2
3"Maxtest-dropoutZ
+
0
+
+

+Z
+
1
+
+

+Z
+
2
+
+

+b
+
2
+
+

+B
\ No newline at end of file

test/onnx/min_test.onnx

+min-example:e
+
+
0
+
1
+
2
3"Mintest-dropoutZ
+
0
+
+

+Z
+
1
+
+

+Z
+
2
+
+

+b
+
2
+
+

+B
\ No newline at end of file

test/onnx/onnx_test.cpp

@@ -169,6 +169,22 @@ void dropout_test()
     EXPECT(p == prog);
 }
 
+void sum_test()
+{
+    migraphx::program p;
+    auto input0 = p.add_parameter("0", migraphx::shape{migraphx::shape::float_type, {3}});
+    auto input1 = p.add_parameter("1", migraphx::shape{migraphx::shape::float_type, {3}});
+    auto input2 = p.add_parameter("2", migraphx::shape{migraphx::shape::float_type, {3}});
+}
+    auto l0 = p.add_instruction(migraphx::op::add{}, input0, input1);
+    p.add_instruction(migraphx::op::add{}, l0, input2);
+    
+
+    auto prog = migraph::parse_onnx("sum_test.onnx");
+    
+    EXPECT(p == prog);
+}
+
 int main()
 {
     pytorch_conv_bias_test();
@@ -181,4 +197,5 @@ int main()
     globalmaxpool_test();
     transpose_test();
     dropout_test();
+    sum_test();
 }

test/onnx/sum_test.onnx

+sum-example:e
+
+
0
+
1
+
2
3"Sumtest-dropoutZ
+
0
+
+

+Z
+
1
+
+

+Z
+
2
+
+

+b
+
2
+
+

+B
\ No newline at end of file