Merge cos and tan operators to current branch.

src/onnx/onnx.cpp

@@ -57,6 +57,9 @@ struct onnx_parser
         // disable dropout for inference
         add_generic_op("Dropout", op::identity{});
         add_generic_op("Identity", op::identity{});
+        add_generic_op("Sin", op::sin{});
+        add_generic_op("Cos", op::cos{});
+        add_generic_op("Tan", op::tan{});
 
         add_broadcastable_binary_op("Add", op::add{});
         add_broadcastable_binary_op("Div", op::div{});

src/targets/gpu/CMakeLists.txt

@@ -13,6 +13,8 @@ endif()
 add_library(migraphx_device
     device/add.cpp
     device/sin.cpp
+    device/cos.cpp
+    device/tan.cpp
     device/add_relu.cpp
     device/contiguous.cpp
     device/mul.cpp
@@ -40,6 +42,8 @@ add_library(migraphx_gpu
     leaky_relu.cpp
     add.cpp
     sin.cpp
+    cos.cpp
+    tan.cpp
     mul.cpp
     batchnorm.cpp
     write_literals.cpp

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

src/targets/gpu/device/cos.cpp

+#include <migraphx/gpu/device/cos.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 cos(hipStream_t stream, const argument& result, const argument& arg)
+{
+    nary(stream, result, arg)([](auto x) { return ::cos(to_hip_type(x)); });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx

src/targets/gpu/device/tan.cpp

+#include <migraphx/gpu/device/tan.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 tan(hipStream_t stream, const argument& result, const argument& arg)
+{
+    nary(stream, result, arg)([](auto x) { return ::tan(to_hip_type(x)); });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx

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

+#ifndef MIGRAPH_GUARD_RTGLIB_COS_HPP
+#define MIGRAPH_GUARD_RTGLIB_COS_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/sin.hpp>
+#include <migraphx/gpu/device/cos.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_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
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+
+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_COS_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_COS_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 cos(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/sin.hpp

 
-#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_ADD_HPP
-#define MIGRAPH_GUARD_RTGLIB_DEVICE_ADD_HPP
+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_SIN_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_SIN_HPP
 
 #include <migraphx/argument.hpp>
 #include <migraphx/config.hpp>

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

+
+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_TAN_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_TAN_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 tan(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/tan.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_TAN_HPP
+#define MIGRAPH_GUARD_RTGLIB_TAN_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/sin.hpp>
+#include <migraphx/gpu/device/cos.hpp>
+#include <migraphx/gpu/device/tan.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_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
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/lowering.cpp

@@ -24,6 +24,8 @@
 #include <migraphx/gpu/softmax.hpp>
 #include <migraphx/gpu/add.hpp>
 #include <migraphx/gpu/sin.hpp>
+#include <migraphx/gpu/cos.hpp>
+#include <migraphx/gpu/tan.hpp>
 #include <migraphx/gpu/mul.hpp>
 #include <migraphx/gpu/batchnorm.hpp>
 #include <migraphx/gpu/pooling.hpp>
@@ -62,6 +64,8 @@ struct miopen_apply
         apply_map["pooling"]              = &miopen_apply::apply_pooling;
         apply_map["add"]                  = &miopen_apply::apply_generic_op<hip_add>;
         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["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;
@@ -211,25 +215,6 @@ struct miopen_apply
         return prog->replace_instruction(ins, T{}, refs);
     }
 
-    /*
-    instruction_ref apply_mul(instruction_ref ins)
-    {
-        auto output = insert_allocation(ins, ins->get_shape());
-        return prog->replace_instruction(
-            ins, hip_mul{}, ins->inputs().at(0), ins->inputs().at(1), output);
-    }
-    */
-
-    /*
-    instruction_ref apply_dot(instruction_ref ins)
-    {
-        auto&& op   = any_cast<op::dot>(ins->get_operator());
-        auto output = insert_allocation(ins, ins->get_shape());
-        return prog->replace_instruction(
-            ins, miopen_gemm{op}, ins->inputs().at(0), ins->inputs().at(1), output);
-    }
-    */
-
     instruction_ref apply_contiguous(instruction_ref ins)
     {
         auto&& op   = any_cast<op::contiguous>(ins->get_operator());

src/targets/gpu/tan.cpp

+#include <migraphx/gpu/tan.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_tan::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(2);
+    return inputs.at(0);
+}
+
+argument hip_tan::compute(context& ctx, const shape&, const std::vector<argument>& args) const
+{
+    device::tan(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,30 @@ struct test_sin
     }
 };
 
+struct test_cos
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::double_type, {8}};
+        auto x = p.add_parameter("x", s);
+        p.add_instruction(migraphx::op::cos{}, x);
+        return p;
+    }
+};
+
+struct test_tan
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        migraphx::shape s{migraphx::shape::float_type, {16}};
+        auto x = p.add_parameter("x", s);
+        p.add_instruction(migraphx::op::tan{}, x);
+        return p;
+    }
+};
+
 struct test_scale
 {
     migraphx::program create_program() const

test/onnx/cos_test.onnx

+cos-example:9
+
+
x
y"Costest_cosZ
+
x
+
+

+
+b
+
y
+
+

+
+B
\ No newline at end of file

test/onnx/onnx_test.cpp

@@ -169,6 +169,39 @@ void dropout_test()
     EXPECT(p == prog);
 }
 
+void sin_test()
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::sin{}, input);
+
+    auto prog = migraphx::parse_onnx("sin_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+void cos_test()
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::cos{}, input);
+
+    auto prog = migraphx::parse_onnx("cos_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+void tan_test()
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::tan{}, input);
+
+    auto prog = migraphx::parse_onnx("tan_test.onnx");
+
+    EXPECT(p == prog);
+}
+
 int main()
 {
     pytorch_conv_bias_test();
@@ -181,4 +214,7 @@ int main()
     globalmaxpool_test();
     transpose_test();
     dropout_test();
+    sin_test();
+    cos_test();
+    tan_test();
 }

test/onnx/sin_test.onnx

+sin-example:9
+
+
x
y"Sintest_sinZ
+
x
+
+

+
+b
+
y
+
+

+
+B
\ No newline at end of file

test/onnx/tan_test.onnx

+tan-example:9
+
+
x
y"Tantest_tanZ
+
x
+
+

+
+b
+
y
+
+

+
+B
\ No newline at end of file