Merge branch 'develop' into lwpck-471

include/ck/tensor_operation/gpu/element/binary_element_wise_operation.hpp

@@ -150,6 +150,13 @@ struct Bilinear
     template <typename Y, typename X0, typename X1>
     __host__ __device__ constexpr void operator()(Y&, const X0&, const X1&) const;
 
+    template <>
+    __host__ __device__ constexpr void
+    operator()<double, double, double>(double& y, const double& x0, const double& x1) const
+    {
+        y = alpha_ * x0 + beta_ * x1;
+    };
+
     template <>
     __host__ __device__ constexpr void
     operator()<float, float, float>(float& y, const float& x0, const float& x1) const

include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp

@@ -95,6 +95,12 @@ struct Scale
         y = scale_ * x;
     };
 
+    template <>
+    __host__ __device__ void operator()<double, double>(double& y, const double& x) const
+    {
+        y = scale_ * x;
+    };
+
     float scale_;
 };
 

include/ck/utility/math_v2.hpp

@@ -83,6 +83,11 @@ static inline __host__ bool isnan(int4_t x)
 };
 #endif
 
+static inline __host__ half_t sqrt(half_t x)
+{
+    return static_cast<half_t>(std::sqrt(static_cast<float>(x)));
+};
+
 static inline __host__ float sqrt(float x) { return std::sqrt(x); };
 
 static inline __host__ double sqrt(double x) { return std::sqrt(x); };
@@ -158,6 +163,11 @@ static inline __device__ bool isnan(half_t x)
     return (xx & 0x7FFF) > 0x7C00;
 };
 
+static inline __device__ half_t sqrt(half_t x)
+{
+    return static_cast<half_t>(__builtin_amdgcn_sqrtf(static_cast<float>(x)));
+};
+
 static inline __device__ float sqrt(float x) { return __builtin_amdgcn_sqrtf(x); };
 
 static inline __device__ double sqrt(double x) { return __builtin_amdgcn_sqrt(x); };

script/cmake-ck-dev.sh

@@ -10,8 +10,8 @@ cmake
 -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc                                                         \
 -D CMAKE_CXX_FLAGS="-O3 -ftemplate-backtrace-limit=0 -gline-tables-only -save-temps=$PWD"         \
 -D CMAKE_BUILD_TYPE=Release                                                                       \
--D BUILD_DEV=ON                                                                                   \
--D GPU_TARGETS="gfx908;gfx90a"                                                                    \
+-D BUILD_DEV=OFF                                                                                   \
+-D GPU_TARGETS="gfx90a"                                                                    \
 -D CMAKE_VERBOSE_MAKEFILE:BOOL=ON                                                                 \
 -D USE_BITINT_EXTENSION_INT4=OFF                                                                  \
 ${MY_PROJECT_SOURCE}