Merge pull request #310 from ROCmSoftwarePlatform/gpu_div

add gpu implementation of the div operator

src/targets/gpu/CMakeLists.txt

@@ -37,6 +37,7 @@ add_library(migraphx_device
     device/pad.cpp
     device/gather.cpp
     device/sub.cpp
+    device/div.cpp
     device/clip.cpp
     device/reduce_sum.cpp
     device/reduce_mean.cpp

src/targets/gpu/device/div.cpp

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

src/targets/gpu/device/sub.cpp

@@ -8,7 +8,7 @@ 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; });
+    nary(stream, result, arg1, arg2)([](auto x, auto y) { return x - y; });
 }
 
 } // namespace device

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

+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_DIV_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_DIV_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 div(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/div.hpp

+#ifndef MIGRAPHX_GUARD_RTGLIB_DIV_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DIV_HPP
+
+#include <migraphx/gpu/oper.hpp>
+#include <migraphx/gpu/device/div.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+struct hip_div : binary_device<hip_div, device::div>
+{
+};
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

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

@@ -88,7 +88,7 @@ struct binary_device : oper<Derived>
 
     argument compute(context& ctx, const shape&, const std::vector<argument>& args) const
     {
-        F(ctx.get_stream().get(), args[2], args[1], args[0]);
+        F(ctx.get_stream().get(), args[2], args[0], args[1]);
         return args[2];
     }
 

src/targets/gpu/lowering.cpp

@@ -26,6 +26,7 @@
 #include <migraphx/gpu/logsoftmax.hpp>
 #include <migraphx/gpu/add.hpp>
 #include <migraphx/gpu/sub.hpp>
+#include <migraphx/gpu/div.hpp>
 #include <migraphx/gpu/exp.hpp>
 #include <migraphx/gpu/erf.hpp>
 #include <migraphx/gpu/log.hpp>
@@ -99,6 +100,7 @@ struct miopen_apply
         add_generic_op<hip_acos>("acos");
         add_generic_op<hip_atan>("atan");
         add_generic_op<hip_mul>("mul");
+        add_generic_op<hip_div>("div");
         add_generic_op<hip_max>("max");
         add_generic_op<hip_min>("min");
 

test/gpu/miopen.cpp

@@ -581,6 +581,38 @@ struct test_sub2 : verify_program<test_sub2>
     }
 };
 
+struct test_div : verify_program<test_div>
+{
+    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::div{}, x, y);
+        p.add_instruction(migraphx::op::div{}, diff, z);
+        return p;
+    }
+};
+
+struct test_div2 : verify_program<test_div2>
+{
+    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.lens()}, z);
+        auto diff = p.add_instruction(migraphx::op::div{}, x, y);
+        p.add_instruction(migraphx::op::div{}, diff, zb);
+        return p;
+    }
+};
+
 struct test_softmax1 : verify_program<test_softmax1>
 {
     migraphx::program create_program() const