merge sqrt operator to test_bert branch

src/include/migraphx/generate.hpp

@@ -81,7 +81,16 @@ template <class T>
 std::vector<T> generate_tensor_data(const migraphx::shape& s, unsigned long seed = 0)
 {
     std::vector<T> result(s.elements());
+    shape::type_t type = s.type();
+    if (type == shape::int64_type or type == shape::int32_type)
+    {
+        std::generate(result.begin(), result.end(), []{ return 1; });
+    }
+    else
+    {
         std::generate(result.begin(), result.end(), xorshf96_generator<T>{seed});    
+    }
+    
     // std::generate(result.begin(), result.end(), [&]{ return seed % 7; });
     // std::generate(result.begin(), result.end(), []{ return 1; });
     return result;

src/include/migraphx/op/sqrt.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_SQRT_HPP
+#define MIGRAPHX_GUARD_OPERATORS_SQRT_HPP
+
+#include <migraphx/op/unary.hpp>
+#include <migraphx/config.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct sqrt : unary<sqrt>
+{
+    auto apply() const
+    {
+        return [](auto x) { return std::sqrt(x); };
+    }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/operators.hpp

@@ -59,6 +59,7 @@
 #include <migraphx/op/sin.hpp>
 #include <migraphx/op/slice.hpp>
 #include <migraphx/op/softmax.hpp>
+#include <migraphx/op/sqrt.hpp>
 #include <migraphx/op/squeeze.hpp>
 #include <migraphx/op/sub.hpp>
 #include <migraphx/op/tanh.hpp>

src/onnx/onnx.cpp

@@ -54,6 +54,7 @@ struct onnx_parser
         add_generic_op("Asin", op::asin{});
         add_generic_op("Acos", op::acos{});
         add_generic_op("Atan", op::atan{});
+        add_generic_op("Sqrt", op::sqrt{});
 
         add_binary_op("Add", op::add{});
         add_binary_op("Div", op::div{});

src/targets/gpu/CMakeLists.txt

@@ -41,6 +41,7 @@ add_library(migraphx_device
     device/reduce_sum.cpp
     device/pow.cpp
     device/reduce_mean.cpp
+    device/sqrt.cpp
 )
 set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
 rocm_clang_tidy_check(migraphx_device)

src/targets/gpu/device/sqrt.cpp

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

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

+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_SQRT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_SQRT_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 sqrt(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/sqrt.hpp

+#ifndef MIGRAPHX_GUARD_RTGLIB_SQRT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_SQRT_HPP
+
+#include <migraphx/gpu/oper.hpp>
+#include <migraphx/gpu/device/sqrt.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+struct hip_sqrt : unary_device<hip_sqrt, device::sqrt>
+{
+};
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/lowering.cpp

@@ -53,6 +53,7 @@
 #include <migraphx/gpu/reduce_sum.hpp>
 #include <migraphx/gpu/pow.hpp>
 #include <migraphx/gpu/reduce_mean.hpp>
+#include <migraphx/gpu/sqrt.hpp>
 #include <utility>
 #include <functional>
 #include <algorithm>
@@ -99,6 +100,7 @@ struct miopen_apply
         add_generic_op<hip_asin>("asin");
         add_generic_op<hip_acos>("acos");
         add_generic_op<hip_atan>("atan");
+        add_generic_op<hip_sqrt>("sqrt");
         add_generic_op<hip_mul>("mul");
         add_generic_op<hip_max>("max");
         add_generic_op<hip_min>("min");