priliminary implementation of the sin operator on GPU.

.gitignore

 *.swp
+.metadata/
+.vscode/
+build/
+deps/
+

src/targets/gpu/CMakeLists.txt

@@ -12,6 +12,7 @@ endif()
 
 add_library(migraphx_device
     device/add.cpp
+    device/sin.cpp
     device/add_relu.cpp
     device/contiguous.cpp
     device/mul.cpp
@@ -38,6 +39,7 @@ add_library(migraphx_gpu
     relu.cpp
     leaky_relu.cpp
     add.cpp
+    sin.cpp
     mul.cpp
     batchnorm.cpp
     write_literals.cpp

src/targets/gpu/device/sin.cpp

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

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

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

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

+#ifndef MIGRAPH_GUARD_RTGLIB_SIN_HPP
+#define MIGRAPH_GUARD_RTGLIB_SIN_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/sin.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <migraph/config.hpp>
+#include <utility>
+
+namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
+namespace gpu {
+
+struct hip_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
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraph
+
+#endif

src/targets/gpu/sin.cpp

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