Merge branch 'gpu_div' into test_bert

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/pow.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 y / x; });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

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/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>
@@ -102,6 +103,7 @@ struct miopen_apply
         add_generic_op<hip_atan>("atan");
         add_generic_op<hip_sqrt>("sqrt");
         add_generic_op<hip_mul>("mul");
+        add_generic_op<hip_div>("div");
         add_generic_op<hip_max>("max");
         add_generic_op<hip_min>("min");
         add_generic_op<hip_pow>("pow");

test/gpu/miopen.cpp

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