[What] use ck::math::sqrt instead of __builtin_amdgcn_sqrtf

[Why] __builtin_amdgcn_sqrtf is only support float, double will cause casting

include/ck/tensor_operation/gpu/grid/gemm_layernorm/gridwise_welford_second_half_layernorm2d.hpp

@@ -350,7 +350,7 @@ struct GridwiseWelfordSecondHalfLayernorm2d
                                   e_thread_buf);
 
         static_for<0, MThreadSliceSize, 1>{}([&](auto m) {
-            auto divisor = 1 / __builtin_amdgcn_sqrtf(welford_var_thread_buf(m) + epsilon);
+            auto divisor = 1 / ck::math::sqrt(welford_var_thread_buf(m) + epsilon);
             static_for<0, NThreadSliceSize, 1>{}([&](auto n) {
                 constexpr auto m_n = thread_buffer_desc_m_n.CalculateOffset(make_tuple(m, n));
                 h_thread_buf(Number<m_n>{}) =

include/ck/tensor_operation/gpu/grid/gridwise_elementwise_layernorm_welford_variance.hpp

@@ -434,7 +434,7 @@ struct GridwiseElementwiseLayernormWelfordVariance_mk_to_mk
             });
 
             static_for<0, MThreadSliceSize, 1>{}([&](auto iM) {
-                auto divisor = 1 / __builtin_amdgcn_sqrtf(var_thread_buf(iM) + epsilon);
+                auto divisor = 1 / ck::math::sqrt(var_thread_buf(iM) + epsilon);
                 static_for<0, XThreadBufferNumber, 1>{}([&](auto iK0) {
                     static_for<0, XSrcVectorSize, 1>{}([&](auto iK1) {
                         constexpr auto offset_m_k =

include/ck/tensor_operation/gpu/grid/gridwise_normalization_welford_variance.hpp

@@ -319,7 +319,7 @@ struct GridwiseNormalizationWelfordVariance_mk_to_mk
             });
 
             static_for<0, MThreadSliceSize, 1>{}([&](auto iM) {
-                auto divisor = 1 / __builtin_amdgcn_sqrtf(var_thread_buf(iM) + epsilon);
+                auto divisor = 1 / ck::math::sqrt(var_thread_buf(iM) + epsilon);
                 static_for<0, XThreadBufferNumber, 1>{}([&](auto iK0) {
                     static_for<0, XSrcVectorSize, 1>{}([&](auto iK1) {
                         constexpr auto offset_m_k =