Combine operator implementation to simplify code.

src/targets/gpu/CMakeLists.txt

@@ -48,18 +48,7 @@ add_library(migraphx_gpu
     relu.cpp
     leaky_relu.cpp
     add.cpp
-    sin.cpp
-    max.cpp
-    min.cpp
-    cos.cpp
-    tan.cpp
-    sinh.cpp
-    cosh.cpp
     tanh.cpp
-    asin.cpp
-    acos.cpp
-    atan.cpp
-    mul.cpp
     batchnorm.cpp
     write_literals.cpp
     rocblas.cpp

src/targets/gpu/acos.cpp

-#include <migraphx/gpu/acos.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-shape hip_acos::compute_shape(const std::vector<shape>& inputs) const
-{
-    check_shapes{inputs, *this}.has(2);
-    return inputs.at(0);
-}
-
-argument hip_acos::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::acos(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

src/targets/gpu/add.cpp

@@ -9,19 +9,6 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-shape hip_add::compute_shape(const std::vector<shape>& inputs) const
-{
-    // check_shapes{inputs, *this}.has(3).standard();
-    check_shapes{inputs, *this}.has(3);
-    return inputs.at(0);
-}
-
-argument hip_add::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::add(ctx.get_stream().get(), args[2], args[0], args[1]);
-    return args[2];
-}
-
 shape miopen_add::compute_shape(const std::vector<shape>& inputs) const
 {
     check_shapes{inputs, *this}.has(3).not_broadcasted();

src/targets/gpu/asin.cpp

-#include <migraphx/gpu/asin.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-shape hip_asin::compute_shape(const std::vector<shape>& inputs) const
-{
-    check_shapes{inputs, *this}.has(2);
-    return inputs.at(0);
-}
-
-argument hip_asin::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::asin(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

src/targets/gpu/atan.cpp

-#include <migraphx/gpu/atan.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-shape hip_atan::compute_shape(const std::vector<shape>& inputs) const
-{
-    check_shapes{inputs, *this}.has(2);
-    return inputs.at(0);
-}
-
-argument hip_atan::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::atan(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

src/targets/gpu/cos.cpp

-#include <migraphx/gpu/cos.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-shape hip_cos::compute_shape(const std::vector<shape>& inputs) const
-{
-    check_shapes{inputs, *this}.has(2);
-    return inputs.at(0);
-}
-
-argument hip_cos::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::cos(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

src/targets/gpu/cosh.cpp

-#include <migraphx/gpu/cosh.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-shape hip_cosh::compute_shape(const std::vector<shape>& inputs) const
-{
-    check_shapes{inputs, *this}.has(2);
-    return inputs.at(0);
-}
-
-argument hip_cosh::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::cosh(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_ACOS_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_acos
+struct hip_acos : unary_device<hip_acos, device::acos>
 {
-    std::string name() const { return "gpu::acos"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_ADD_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_add
+struct hip_add : binary_device<hip_add, device::add>
 {
-    std::string name() const { return "gpu::add"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 struct miopen_add

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_ASIN_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_asin
+struct hip_asin : unary_device<hip_asin, device::asin>
 {
-    std::string name() const { return "gpu::asin"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_ATAN_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_atan
+struct hip_atan : unary_device<hip_atan, device::atan>
 {
-    std::string name() const { return "gpu::atan"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_COS_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -11,7 +12,6 @@
 #include <migraphx/gpu/hip.hpp>
 #include <migraphx/dfor.hpp>
 #include <migraphx/gpu/device/contiguous.hpp>
-#include <migraphx/gpu/device/sin.hpp>
 #include <migraphx/gpu/device/cos.hpp>
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/gpu/rocblas.hpp>
@@ -23,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_cos
+struct hip_cos : unary_device<hip_cos, device::cos>
 {
-    std::string name() const { return "gpu::cos"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_COSH_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_cosh
+struct hip_cosh : unary_device<hip_cosh, device::cosh>
 {
-    std::string name() const { return "gpu::cosh"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

-#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_ADD_HPP
-#define MIGRAPHX_GUARD_RTGLIB_DEVICE_ADD_HPP
+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_SIN_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_SIN_HPP
 
 #include <migraphx/argument.hpp>
 #include <migraphx/config.hpp>

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_MAX_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_max
+struct hip_max : binary_device<hip_max, device::max>
 {
-    std::string name() const { return "gpu::max"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_MIN_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_min
+struct hip_min : binary_device<hip_min, device::min>
 {
-    std::string name() const { return "gpu::min"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_MUL_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_mul
+struct hip_mul : binary_device<hip_mul, device::mul>
 {
-    std::string name() const { return "gpu::mul"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_SIN_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_sin
+struct hip_sin : unary_device<hip_sin, device::sin>
 {
-    std::string name() const { return "gpu::sin"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_SINH_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_sinh
+struct hip_sinh : unary_device<hip_sinh, device::sinh>
 {
-    std::string name() const { return "gpu::sinh"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_TAN_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -11,8 +12,6 @@
 #include <migraphx/gpu/hip.hpp>
 #include <migraphx/dfor.hpp>
 #include <migraphx/gpu/device/contiguous.hpp>
-#include <migraphx/gpu/device/sin.hpp>
-#include <migraphx/gpu/device/cos.hpp>
 #include <migraphx/gpu/device/tan.hpp>
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/gpu/rocblas.hpp>
@@ -24,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_tan
+struct hip_tan : unary_device<hip_tan, device::tan>
 {
-    std::string name() const { return "gpu::tan"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu