Merge pull request #168 from ROCmSoftwarePlatform/sub_gpu_implementation

Sub gpu implementation

src/targets/gpu/CMakeLists.txt

@@ -30,6 +30,7 @@ add_library(migraphx_device
     device/concat.cpp
     device/pad.cpp
     device/gather.cpp
+    device/sub.cpp
 )
 set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
 rocm_clang_tidy_check(migraphx_device)

src/targets/gpu/device/sub.cpp

+#include <migraphx/gpu/device/sub.hpp>
+#include <migraphx/gpu/device/nary.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void sub(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2)
+{
+    nary(stream, result, arg1, arg2)([](auto x, auto y) { return y - x; });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

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

+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_SUB_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_SUB_HPP
+
+#include <migraphx/argument.hpp>
+#include <migraphx/config.hpp>
+#include <hip/hip_runtime_api.h>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void sub(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPHX_GUARD_RTGLIB_SUB_HPP
+#define MIGRAPHX_GUARD_RTGLIB_SUB_HPP
+
+#include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/oper.hpp>
+#include <migraphx/manage_ptr.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/operators.hpp>
+#include <migraphx/generate.hpp>
+#include <migraphx/shape_for_each.hpp>
+#include <migraphx/gpu/miopen.hpp>
+#include <migraphx/gpu/hip.hpp>
+#include <migraphx/dfor.hpp>
+#include <migraphx/gpu/device/contiguous.hpp>
+#include <migraphx/gpu/device/sub.hpp>
+#include <migraphx/iterator_for.hpp>
+#include <migraphx/gpu/rocblas.hpp>
+#include <migraphx/gpu/context.hpp>
+#include <migraphx/config.hpp>
+#include <utility>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+struct hip_sub : binary_device<hip_sub, device::sub>
+{
+};
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/lowering.cpp

@@ -22,6 +22,7 @@
 #include <migraphx/gpu/elu.hpp>
 #include <migraphx/gpu/softmax.hpp>
 #include <migraphx/gpu/add.hpp>
+#include <migraphx/gpu/sub.hpp>
 #include <migraphx/gpu/exp.hpp>
 #include <migraphx/gpu/log.hpp>
 #include <migraphx/gpu/sin.hpp>
@@ -74,6 +75,7 @@ struct miopen_apply
         add_miopen_extend_op<miopen_elu, op::elu>("elu", make_elu);
 
         add_generic_op<hip_add>("add");
+        add_generic_op<hip_sub>("sub");
         add_generic_op<hip_exp>("exp");
         add_generic_op<hip_log>("log");
         add_generic_op<hip_sin>("sin");

test/gpu/miopen.cpp

@@ -483,6 +483,38 @@ struct test_triadd_broadcast
     }
 };
 
+struct test_sub
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {3}};
+        auto x    = p.add_parameter("x", s);
+        auto y    = p.add_parameter("y", s);
+        auto z    = p.add_parameter("z", s);
+        auto diff = p.add_instruction(migraphx::op::sub{}, x, y);
+        p.add_instruction(migraphx::op::sub{}, diff, z);
+        return p;
+    }
+};
+
+struct test_sub2
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        migraphx::shape b{migraphx::shape::float_type, {3}};
+        auto x    = p.add_parameter("x", s);
+        auto y    = p.add_parameter("y", s);
+        auto z    = p.add_parameter("z", b);
+        auto zb   = p.add_instruction(migraphx::op::broadcast{1, s}, z);
+        auto diff = p.add_instruction(migraphx::op::sub{}, x, y);
+        p.add_instruction(migraphx::op::sub{}, diff, zb);
+        return p;
+    }
+};
+
 struct test_softmax
 {
     migraphx::program create_program() const
@@ -1115,6 +1147,8 @@ int main()
     verify_program<test_add_broadcast4>();
     verify_program<test_add_broadcast5>();
     verify_program<test_triadd_broadcast>();
+    verify_program<test_sub>();
+    verify_program<test_sub2>();
     verify_program<test_softmax>();
     verify_program<test_softmax2>();
     verify_program<test_conv>();

test/onnx/implicit_sub_bcast_test.onnx

+subtraction2:q
+
+
0
+
1out"Subsubtraction2Z
+
0
+

+
+
+
+Z
+
1
+

+
+b
+out
+

+
+
+
+B
\ No newline at end of file

test/onnx/onnx_test.cpp

@@ -346,7 +346,7 @@ TEST_CASE(add_bcast_test)
     EXPECT(p == prog);
 }
 
-TEST_CASE(implicit_bcast_test)
+TEST_CASE(implicit_add_bcast_test)
 {
     migraphx::program p;
     auto l0 = p.add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 3, 4, 5}});
@@ -360,6 +360,33 @@ TEST_CASE(implicit_bcast_test)
     EXPECT(p == prog);
 }
 
+TEST_CASE(sub_bcast_test)
+{
+    migraphx::program p;
+    auto l0 = p.add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 3, 4, 5}});
+    auto l1 = p.add_parameter("1", migraphx::shape{migraphx::shape::float_type, {3, 4}});
+    auto l2 = p.add_instruction(migraphx::op::broadcast{1, l0->get_shape()}, l1);
+    p.add_instruction(migraphx::op::sub{}, l0, l2);
+
+    auto prog = migraphx::parse_onnx("sub_bcast_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+TEST_CASE(implicit_sub_bcast_test)
+{
+    migraphx::program p;
+    auto l0 = p.add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 3, 4, 5}});
+    auto l1 = p.add_parameter("1", migraphx::shape{migraphx::shape::float_type, {3, 4}});
+    auto l2 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 4, 5}}, l0);
+    auto l3 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 4, 5}}, l1);
+    p.add_instruction(migraphx::op::sub{}, l2, l3);
+
+    auto prog = migraphx::parse_onnx("implicit_sub_bcast_test.onnx");
+
+    EXPECT(p == prog);
+}
+
 TEST_CASE(unknown_test)
 {
     migraphx::program p;
@@ -550,6 +577,20 @@ TEST_CASE(add_scalar_test)
     EXPECT(p == prog);
 }
 
+TEST_CASE(sub_scalar_test)
+{
+    migraphx::program p;
+    auto l0 = p.add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 3, 4, 5}});
+    auto l1 =
+        p.add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {1}}, {1}});
+    auto m0 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 4, 5}}, l0);
+    auto m1 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 4, 5}}, l1);
+    p.add_instruction(migraphx::op::sub{}, m0, m1);
+    auto prog = migraphx::parse_onnx("sub_scalar_test.onnx");
+
+    EXPECT(p == prog);
+}
+
 TEST_CASE(group_conv_test)
 {
     migraphx::program p;

test/onnx/sub_bcast_test.onnx

+subtraction2:

+/
+
0
+
1out"Sub*
+axis

*
+	broadcast

subtraction2Z
+
0
+

+
+
+
+Z
+
1
+

+
+b
+out
+

+
+
+
+B
\ No newline at end of file