Merge operators implemenation to current branch.

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/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/lowering.cpp

@@ -17,7 +17,6 @@
 #include <migraphx/gpu/contiguous.hpp>
 #include <migraphx/gpu/relu.hpp>
 #include <migraphx/gpu/sigmoid.hpp>
-#include <migraphx/gpu/tanh.hpp>
 #include <migraphx/gpu/abs.hpp>
 #include <migraphx/gpu/leaky_relu.hpp>
 #include <migraphx/gpu/elu.hpp>
@@ -26,6 +25,12 @@
 #include <migraphx/gpu/sin.hpp>
 #include <migraphx/gpu/cos.hpp>
 #include <migraphx/gpu/tan.hpp>
+#include <migraphx/gpu/sinh.hpp>
+#include <migraphx/gpu/cosh.hpp>
+#include <migraphx/gpu/tanh.hpp>
+#include <migraphx/gpu/asin.hpp>
+#include <migraphx/gpu/acos.hpp>
+#include <migraphx/gpu/atan.hpp>
 #include <migraphx/gpu/mul.hpp>
 #include <migraphx/gpu/batchnorm.hpp>
 #include <migraphx/gpu/pooling.hpp>
@@ -57,7 +62,6 @@ struct miopen_apply
         apply_map["convolution"]          = &miopen_apply::apply_convolution;
         apply_map["relu"]                 = &miopen_apply::apply_relu;
         apply_map["sigmoid"]              = &miopen_apply::apply_sigmoid;
-        apply_map["tanh"]                 = &miopen_apply::apply_tanh;
         apply_map["abs"]                  = &miopen_apply::apply_abs;
         apply_map["leaky_relu"]           = &miopen_apply::apply_leaky_relu;
         apply_map["elu"]                  = &miopen_apply::apply_elu;
@@ -66,6 +70,12 @@ struct miopen_apply
         apply_map["sin"]                  = &miopen_apply::apply_generic_op<hip_sin>;
         apply_map["cos"]                  = &miopen_apply::apply_generic_op<hip_cos>;
         apply_map["tan"]                  = &miopen_apply::apply_generic_op<hip_tan>;
+        apply_map["sinh"]                 = &miopen_apply::apply_generic_op<hip_sinh>;
+        apply_map["cosh"]                 = &miopen_apply::apply_generic_op<hip_cosh>;
+        apply_map["tanh"]                 = &miopen_apply::apply_tanh;
+        apply_map["asin"]                 = &miopen_apply::apply_generic_op<hip_asin>;
+        apply_map["acos"]                 = &miopen_apply::apply_generic_op<hip_acos>;
+        apply_map["atan"]                 = &miopen_apply::apply_generic_op<hip_atan>;
         apply_map["mul"]                  = &miopen_apply::apply_generic_op<hip_mul>;
         apply_map["dot"]                  = &miopen_apply::apply_generic_op<miopen_gemm>;
         apply_map["contiguous"]           = &miopen_apply::apply_contiguous;
@@ -188,23 +198,6 @@ struct miopen_apply
         return prog->replace_instruction(ins, miopen_softmax{op}, ins->inputs().at(0), output);
     }
 
-    /*
-    instruction_ref apply_add(instruction_ref ins)
-    {
-        auto output = insert_allocation(ins, ins->get_shape());
-        return prog->replace_instruction(
-            ins, hip_add{}, ins->inputs().at(0), ins->inputs().at(1), output);
-    }
-    */
-
-    /*
-    instruction_ref apply_sin(instruction_ref ins)
-    {
-        auto output = insert_allocation(ins, ins->get_shape());
-        return prog->replace_instruction(ins, hip_sin{}, ins->inputs().at(0), output);
-    }
-    */
-
     template <class T>
     instruction_ref apply_generic_op(instruction_ref ins)
     {

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

test/gpu/miopen.cpp

@@ -239,6 +239,77 @@ struct test_tan
     }
 };
 
+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;
+        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
+        p.add_instruction(migraphx::op::tanh{}, x);
+        return p;
+    }
+};
+
+struct test_asin
+{
+    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::asin{}, x);
+        return p;
+    }
+};
+
+struct test_acos
+{
+    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::acos{}, x);
+        return p;
+    }
+};
+
+struct test_atan
+{
+    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::atan{}, x);
+        return p;
+    }
+};
+
 struct test_scale
 {
     migraphx::program create_program() const
@@ -491,17 +562,6 @@ struct test_sigmoid
     }
 };
 
-struct test_tanh
-{
-    migraphx::program create_program() const
-    {
-        migraphx::program p;
-        auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
-        p.add_instruction(migraphx::op::tanh{}, x);
-        return p;
-    }
-};
-
 struct test_abs
 {
     migraphx::program create_program() const
@@ -925,6 +985,14 @@ int main()
     verify_program<test_add_half>();
     verify_program<test_mul>();
     verify_program<test_sin>();
+    verify_program<test_cos>();
+    verify_program<test_tan>();
+    verify_program<test_sinh>();
+    verify_program<test_cosh>();
+    verify_program<test_tanh>();
+    verify_program<test_asin>();
+    verify_program<test_acos>();
+    verify_program<test_atan>();
     verify_program<test_scale>();
     verify_program<test_triadd>();
     verify_program<test_triadd2>();
@@ -943,7 +1011,6 @@ int main()
     verify_program<test_add_relu>();
     verify_program<test_leaky_relu>();
     verify_program<test_sigmoid>();
-    verify_program<test_tanh>();
     verify_program<test_elu>();
     verify_program<test_conv_pooling>();
     verify_program<test_global_avg_pooling>();

test/onnx/acos_test.onnx

+acos-example:;
+
+
x
y"Acos	test_acosZ
+
x
+
+

+
+b
+
y
+
+

+
+B	
\ No newline at end of file

test/onnx/asin_test.onnx

+asin-example:;
+
+
x
y"Asin	test_asinZ
+
x
+
+

+
+b
+
y
+
+

+
+B
\ No newline at end of file

test/onnx/atan_test.onnx

+atan-example:;
+
+
x
y"Atan	test_atanZ
+
x
+
+

+
+b
+
y
+
+

+
+B	
\ No newline at end of file

test/onnx/cosh_test.onnx

+cosh-example:;
+
+
x
y"Cosh	test_coshZ
+
x
+
+

+
b
+
y
+
+

+
B	
\ No newline at end of file

test/onnx/onnx_cosh.onnx

+cosh-example:;
+
+
x
y"Cosh	test_coshZ
+
x
+
+

+
b
+
y
+
+

+
B	
\ No newline at end of file

test/onnx/onnx_sinh.onnx

+sinh-example:;
+
+
x
y"Sinh	test_sinhZ
+
x
+
+

+
+b
+
y
+
+

+
+B	
\ No newline at end of file

test/onnx/onnx_test.cpp

@@ -176,10 +176,10 @@ void sin_test()
     p.add_instruction(migraphx::op::sin{}, input);
 
     auto prog = migraphx::parse_onnx("sin_test.onnx");
-
     EXPECT(p == prog);
 }
 
+
 void cos_test()
 {
     migraphx::program p;
@@ -187,10 +187,10 @@ void cos_test()
     p.add_instruction(migraphx::op::cos{}, input);
 
     auto prog = migraphx::parse_onnx("cos_test.onnx");
-
     EXPECT(p == prog);
 }
 
+
 void tan_test()
 {
     migraphx::program p;
@@ -198,6 +198,61 @@ void tan_test()
     p.add_instruction(migraphx::op::tan{}, input);
 
     auto prog = migraphx::parse_onnx("tan_test.onnx");
+    EXPECT(p == prog);
+}
+
+
+void sinh_test()
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::sinh{}, input);
+
+    auto prog = migraphx::parse_onnx("sinh_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+void cosh_test()
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1}});
+    p.add_instruction(migraphx::op::cosh{}, input);
+
+    auto prog = migraphx::parse_onnx("cosh_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+void asin_test()
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::asin{}, input);
+
+    auto prog = migraphx::parse_onnx("asin_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+void acos_test()
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::acos{}, input);
+
+    auto prog = migraphx::parse_onnx("acos_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+void atan_test()
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::atan{}, input);
+
+    auto prog = migraphx::parse_onnx("atan_test.onnx");
 
     EXPECT(p == prog);
 }
@@ -217,4 +272,9 @@ int main()
     sin_test();
     cos_test();
     tan_test();
+    sinh_test();
+    cosh_test();
+    asin_test();
+    acos_test();
+    atan_test();
 }

test/onnx/sinh_test.onnx

+sinh-example:;
+
+
x
y"Sinh	test_sinhZ
+
x
+
+

+
+b
+
y
+
+

+
+B	
\ No newline at end of file