add operators, sinh cosh and tanh.

src/include/migraphx/operators.hpp

@@ -656,6 +656,16 @@ struct atan : unary
     std::string name() const { return "atan"; }
 };
 
+struct sinh : unary
+{
+    std::string name() const { return "sinh"; }
+};
+
+struct cosh : unary
+{
+    std::string name() const { return "cosh"; }
+};
+
 struct tanh : unary
 {
     std::string name() const { return "tanh"; }

src/targets/cpu/lowering.cpp

@@ -406,6 +406,24 @@ struct atan_op
     }
 };
 
+struct sinh_op
+{
+    std::string name() const { return "cpu::sinh"; }
+    auto fcn() const
+    {
+        return [](auto x) { return std::sinh(x); };
+    }
+};
+
+struct cosh_op
+{
+    std::string name() const { return "cpu::cosh"; }
+    auto fcn() const
+    {
+        return [](auto x) { return std::cosh(x); };
+    }
+};
+
 struct tanh_op
 {
     std::string name() const { return "cpu::tanh"; }
@@ -600,6 +618,8 @@ struct cpu_apply
         apply_map["concat"]     = extend_op<cpu_concat, op::concat>();
         apply_map["leaky_relu"] = extend_op<cpu_unary<leaky_relu_op>, op::leaky_relu>();
         apply_map["identity"]   = simple_op<cpu_unary<identity_op>>();
+        apply_map["sinh"]       = simple_op<cpu_unary<sinh_op>>();
+        apply_map["cosh"]       = simple_op<cpu_unary<cosh_op>>();
         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>>();

src/targets/gpu/CMakeLists.txt

@@ -13,6 +13,9 @@ endif()
 add_library(migraphx_device
     device/add.cpp
     device/sin.cpp
+    device/sinh.cpp
+    device/cosh.cpp
+    device/tanh.cpp
     device/add_relu.cpp
     device/contiguous.cpp
     device/mul.cpp
@@ -40,6 +43,9 @@ add_library(migraphx_gpu
     leaky_relu.cpp
     add.cpp
     sin.cpp
+    sinh.cpp
+    cosh.cpp
+    tanh.cpp
     mul.cpp
     batchnorm.cpp
     write_literals.cpp

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 MIGRAPH_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 MIGRAPH_INLINE_NS
+} // namespace migraphx

src/targets/gpu/device/cosh.cpp

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

src/targets/gpu/device/sinh.cpp

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

src/targets/gpu/device/tanh.cpp

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

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

+#ifndef MIGRAPH_GUARD_RTGLIB_COSH_HPP
+#define MIGRAPH_GUARD_RTGLIB_COSH_HPP
+
+#include <migraphx/gpu/lowering.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/cosh.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 MIGRAPH_INLINE_NS {
+namespace gpu {
+
+struct hip_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
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_COSH_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_COSH_HPP
+
+#include <migraphx/argument.hpp>
+#include <migraphx/config.hpp>
+#include <hip/hip_runtime_api.h>
+
+namespace migraphx {
+inline namespace MIGRAPH_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void cosh(hipStream_t stream, const argument& result, const argument& arg);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_SINH_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_SINH_HPP
+
+#include <migraphx/argument.hpp>
+#include <migraphx/config.hpp>
+#include <hip/hip_runtime_api.h>
+
+namespace migraphx {
+inline namespace MIGRAPH_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void sinh(hipStream_t stream, const argument& result, const argument& arg);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_TANH_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_TANH_HPP
+
+#include <migraphx/argument.hpp>
+#include <migraphx/config.hpp>
+#include <hip/hip_runtime_api.h>
+
+namespace migraphx {
+inline namespace MIGRAPH_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void tanh(hipStream_t stream, const argument& result, const argument& arg);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPH_GUARD_RTGLIB_SINH_HPP
+#define MIGRAPH_GUARD_RTGLIB_SINH_HPP
+
+#include <migraphx/gpu/lowering.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/sinh.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 MIGRAPH_INLINE_NS {
+namespace gpu {
+
+struct hip_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
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPH_GUARD_RTGLIB_TANH_HPP
+#define MIGRAPH_GUARD_RTGLIB_TANH_HPP
+
+#include <migraphx/gpu/lowering.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/tanh.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 MIGRAPH_INLINE_NS {
+namespace gpu {
+
+struct hip_tanh
+{
+    std::string name() const { return "gpu::tanh"; }
+    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
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/lowering.cpp

@@ -20,6 +20,9 @@
 #include <migraphx/gpu/softmax.hpp>
 #include <migraphx/gpu/add.hpp>
 #include <migraphx/gpu/sin.hpp>
+#include <migraphx/gpu/sinh.hpp>
+#include <migraphx/gpu/cosh.hpp>
+#include <migraphx/gpu/tanh.hpp>
 #include <migraphx/gpu/mul.hpp>
 #include <migraphx/gpu/batchnorm.hpp>
 #include <migraphx/gpu/pooling.hpp>
@@ -72,6 +75,18 @@ struct miopen_apply
             {
                 check_shape(s, apply_sin(it));
             }
+            else if(it->name() == "sinh")
+            {
+                check_shape(s, apply_sinh(it));
+            }
+            else if(it->name() == "cosh")
+            {
+                check_shape(s, apply_cosh(it));
+            }
+            else if(it->name() == "tanh")
+            {
+                check_shape(s, apply_tanh(it));
+            }
             else if(it->name() == "mul")
             {
                 check_shape(s, apply_mul(it));
@@ -176,6 +191,24 @@ struct miopen_apply
         return prog->replace_instruction(ins, hip_sin{}, ins->inputs().at(0), output);
     }
 
+    instruction_ref apply_sinh(instruction_ref ins)
+    {
+        auto output = insert_allocation(ins, ins->get_shape());
+        return prog->replace_instruction(ins, hip_sinh{}, ins->inputs().at(0), output);
+    }
+
+    instruction_ref apply_cosh(instruction_ref ins)
+    {
+        auto output = insert_allocation(ins, ins->get_shape());
+        return prog->replace_instruction(ins, hip_cosh{}, ins->inputs().at(0), output);
+    }
+
+    instruction_ref apply_tanh(instruction_ref ins)
+    {
+        auto output = insert_allocation(ins, ins->get_shape());
+        return prog->replace_instruction(ins, hip_tanh{}, ins->inputs().at(0), output);
+    }
+
     instruction_ref apply_mul(instruction_ref ins)
     {
         auto output = insert_allocation(ins, ins->get_shape());

src/targets/gpu/sinh.cpp

+#include <migraphx/gpu/sinh.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 MIGRAPH_INLINE_NS {
+namespace gpu {
+
+shape hip_sinh::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(2);
+    return inputs.at(0);
+}
+
+argument hip_sinh::compute(context& ctx, const shape&, const std::vector<argument>& args) const
+{
+    device::sinh(ctx.get_stream().get(), args[1], args[0]);
+    return args[1];
+}
+
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx

src/targets/gpu/tanh.cpp

+#include <migraphx/gpu/tanh.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 MIGRAPH_INLINE_NS {
+namespace gpu {
+
+shape hip_tanh::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(2);
+    return inputs.at(0);
+}
+
+argument hip_tanh::compute(context& ctx, const shape&, const std::vector<argument>& args) const
+{
+    device::tanh(ctx.get_stream().get(), args[1], args[0]);
+    return args[1];
+}
+
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx

test/gpu/miopen.cpp

@@ -215,6 +215,43 @@ struct test_sin
     }
 };
 
+struct test_sinh
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::double_type, {16}};
+        auto x = p.add_parameter("x", s);
+        p.add_instruction(migraphx::op::sinh{}, x);
+        return p;
+    }
+};
+
+struct test_cosh
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::double_type, {16}};
+        auto x = p.add_parameter("x", s);
+        p.add_instruction(migraphx::op::cosh{}, x);
+        return p;
+    }
+};
+
+struct test_tanh
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::double_type, {16}};
+        auto x = p.add_parameter("x", s);
+        p.add_instruction(migraphx::op::tanh{}, x);
+        return p;
+    }
+};
+
+
 struct test_scale
 {
     migraphx::program create_program() const
@@ -856,6 +893,9 @@ int main()
     verify_program<test_add_half>();
     verify_program<test_mul>();
     verify_program<test_sin>();
+    verify_program<test_sinh>();
+    verify_program<test_cosh>();
+    verify_program<test_tanh>();
     verify_program<test_scale>();
     verify_program<test_triadd>();
     verify_program<test_triadd2>();