manual merge

src/include/migraphx/operators.hpp

@@ -641,6 +641,11 @@ struct exp : unary
     std::string name() const { return "exp"; }
 };
 
+struct log : unary
+{
+    std::string name() const { return "log"; }
+};
+
 struct sin : unary
 {
     std::string name() const { return "sin"; }

src/onnx/onnx.cpp

@@ -38,6 +38,8 @@ struct onnx_parser
         add_generic_op("Relu", op::relu{});
         add_generic_op("Sigmoid", op::sigmoid{});
         add_generic_op("Abs", op::abs{});
+        add_generic_op("Exp", op::exp{});
+        add_generic_op("Log", op::log{});
         // disable dropout for inference
         add_generic_op("Dropout", op::identity{});
         add_generic_op("Identity", op::identity{});

src/targets/cpu/lowering.cpp

@@ -360,6 +360,15 @@ struct exp_op
     }
 };
 
+struct log_op
+{
+    std::string name() const { return "cpu::log"; }
+    auto fcn() const
+    {
+        return [](auto x) { return std::log(x); };
+    }
+};
+
 struct sin_op
 {
     std::string name() const { return "cpu::sin"; }
@@ -662,6 +671,7 @@ struct cpu_apply
         apply_map["tanh"]       = simple_op<cpu_unary<tanh_op>>();
         apply_map["sigmoid"]    = simple_op<cpu_unary<sigmoid_op>>();
         apply_map["exp"]        = simple_op<cpu_unary<exp_op>>();
+        apply_map["log"]        = simple_op<cpu_unary<log_op>>();
         apply_map["neg"]        = simple_op<cpu_unary<neg_op>>();
         apply_map["sin"]        = simple_op<cpu_unary<sin_op>>();
         apply_map["cos"]        = simple_op<cpu_unary<cos_op>>();

src/targets/gpu/CMakeLists.txt

@@ -14,6 +14,8 @@ add_library(migraphx_device
     device/add.cpp
     device/max.cpp
     device/min.cpp
+    device/exp.cpp
+    device/log.cpp
     device/sin.cpp
     device/cos.cpp
     device/tan.cpp
@@ -47,19 +49,7 @@ add_library(migraphx_gpu
     concat.cpp
     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/add.cpp

-#include <migraphx/gpu/add.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_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();
-    return inputs.at(0);
-}
-
-argument miopen_add::compute(context& ctx,
-                             const shape& output_shape,
-                             const std::vector<argument>& args) const
-{
-    float alpha = 1, beta = 0;
-    auto a_desc = make_tensor(args[0].get_shape());
-    auto b_desc = make_tensor(args[1].get_shape());
-    auto c_desc = make_tensor(output_shape);
-    miopenOpTensor(ctx.get_stream().get_miopen(),
-                   miopenTensorOpAdd,
-                   &alpha,
-                   a_desc.get(),
-                   args[0].implicit(),
-                   &alpha,
-                   b_desc.get(),
-                   args[1].implicit(),
-                   &beta,
-                   c_desc.get(),
-                   args[2].implicit());
-    return args[2];
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

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/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/device/exp.cpp

+#include <migraphx/gpu/device/exp.hpp>
+#include <migraphx/gpu/device/nary.hpp>
+#include <migraphx/gpu/device/types.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void exp(hipStream_t stream, const argument& result, const argument& arg)
+{
+    nary(stream, result, arg)([](auto x) { return ::exp(to_hip_type(x)); });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/targets/gpu/device/log.cpp

+#include <migraphx/gpu/device/log.hpp>
+#include <migraphx/gpu/device/nary.hpp>
+#include <migraphx/gpu/device/types.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void log(hipStream_t stream, const argument& result, const argument& arg)
+{
+    nary(stream, result, arg)([](auto x) { return ::log(to_hip_type(x)); });
+}
+
+} // namespace device
+} // 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/oper.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/oper.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,21 +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
-{
-    std::string name() const { return "gpu::add"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument
-    compute(context& ctx, const shape& output_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/asin.hpp

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_ASIN_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/oper.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/oper.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/oper.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/oper.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/exp.hpp

+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_EXP_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_EXP_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 exp(hipStream_t stream, const argument& result, const argument& arg);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_LOG_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_LOG_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 log(hipStream_t stream, const argument& result, const argument& arg);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

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/acos.cpp → src/targets/gpu/include/migraphx/gpu/exp.hpp

-#include <migraphx/gpu/acos.hpp>
-#include <migraphx/operators.hpp>
+#ifndef MIGRAPHX_GUARD_RTGLIB_EXP_HPP
+#define MIGRAPHX_GUARD_RTGLIB_EXP_HPP
+
+#include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/oper.hpp>
 #include <migraphx/manage_ptr.hpp>
-#include <migraphx/config.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/exp.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 {
 
-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
+struct hip_exp : unary_device<hip_exp, device::exp>
 {
-    device::acos(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
+};
 
 } // namespace gpu
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx
+
+#endif