put function in gko not std

Co-authored-by: Thomas Grützmacher <[email protected]>

core/test/utils/assertions.hpp

@@ -259,7 +259,7 @@ void save_matrices_to_disk(Ostream& os, const MatrixData1& first,
 template <typename MatrixData1, typename MatrixData2>
 double get_relative_error(const MatrixData1& first, const MatrixData2& second)
 {
-    using std::abs;
+    using gko::abs;
     using vt = typename detail::biggest_valuetype<
         typename MatrixData1::value_type,
         typename MatrixData2::value_type>::type;

core/test/utils/matrix_generator_test.cpp

@@ -278,8 +278,8 @@ TYPED_TEST(MatrixGenerator, CanGenerateTridiagInverseMatrix)
     auto lower = gko::test::detail::get_rand_value<T>(dist, engine);
     auto upper = gko::test::detail::get_rand_value<T>(dist, engine);
     // make diagonally dominant
-    auto diag = std::abs(gko::test::detail::get_rand_value<T>(dist, engine)) +
-                std::abs(lower) + std::abs(upper);
+    auto diag = gko::abs(gko::test::detail::get_rand_value<T>(dist, engine)) +
+                gko::abs(lower) + gko::abs(upper);
     gko::size_type size = 50;
     if (std::is_same<gko::half, gko::remove_complex<T>>::value) {
         // half precision can only handle small matrix

dpcpp/factorization/par_ilut_filter_kernels.hpp.inc

@@ -102,7 +102,7 @@ void threshold_filter_nnz(const IndexType* __restrict__ row_ptrs,
         row_ptrs, num_rows,
         [&](IndexType idx, IndexType row_begin, IndexType row_end) {
             auto diag_idx = lower ? row_end - 1 : row_begin;
-            return std::abs(vals[idx]) >= threshold || idx == diag_idx;
+            return gko::abs(vals[idx]) >= threshold || idx == diag_idx;
         },
         nnz, item_ct1);
 }
@@ -140,7 +140,7 @@ void threshold_filter(const IndexType* __restrict__ old_row_ptrs,
         old_row_ptrs, old_col_idxs, old_vals, num_rows,
         [&](IndexType idx, IndexType row_begin, IndexType row_end) {
             auto diag_idx = lower ? row_end - 1 : row_begin;
-            return std::abs(old_vals[idx]) >= threshold || idx == diag_idx;
+            return gko::abs(old_vals[idx]) >= threshold || idx == diag_idx;
         },
         new_row_ptrs, new_row_idxs, new_col_idxs, new_vals, item_ct1);
 }

dpcpp/factorization/par_ilut_select_kernels.hpp.inc

@@ -38,7 +38,7 @@ void build_searchtree(const ValueType* __restrict__ input, IndexType size,
     for (int i = 0; i < sampleselect_oversampling; ++i) {
         auto lidx = idx * sampleselect_oversampling + i;
         auto val = input[static_cast<IndexType>(lidx * stride)];
-        samples[i] = std::abs(val);
+        samples[i] = gko::abs(val);
     }
 
     bitonic_sort<sample_size, sampleselect_oversampling>(samples, sh_samples,
@@ -113,7 +113,7 @@ void count_buckets(const ValueType* __restrict__ input, IndexType size,
     auto end = min(block_end, size);
     for (IndexType i = begin; i < end; i += default_block_size) {
         // traverse the search tree with the input element
-        auto el = std::abs(input[i]);
+        auto el = gko::abs(input[i]);
         IndexType tree_idx{};
 #pragma unroll
         for (int level = 0; level < sampleselect_searchtree_height; ++level) {
@@ -297,7 +297,7 @@ void filter_bucket(const ValueType* __restrict__ input, IndexType size,
         auto found = bucket == oracles[i];
         auto ofs = atomic_add<atomic::local_space>(&*counter, IndexType{found});
         if (found) {
-            output[ofs] = std::abs(input[i]);
+            output[ofs] = gko::abs(input[i]);
         }
     }
 }

dpcpp/solver/cb_gmres_kernels.dp.cpp

@@ -285,9 +285,9 @@ void multinorminf_without_stop_kernel(
              i += default_dot_dim) {
             const auto next_krylov_idx = i * stride_next_krylov + col_idx;
             local_max =
-                (local_max >= std::abs(next_krylov_basis[next_krylov_idx]))
+                (local_max >= gko::abs(next_krylov_basis[next_krylov_idx]))
                     ? local_max
-                    : std::abs(next_krylov_basis[next_krylov_idx]);
+                    : gko::abs(next_krylov_basis[next_krylov_idx]);
         }
     }
     reduction_helper[tidx * (default_dot_dim + 1) + tidy] = local_max;
@@ -373,7 +373,7 @@ void multinorm2_inf_kernel(
             local_res += squared_norm(num);
             if (compute_inf) {
                 local_max =
-                    ((local_max >= std::abs(num)) ? local_max : std::abs(num));
+                    ((local_max >= gko::abs(num)) ? local_max : gko::abs(num));
             }
         }
     }

dpcpp/solver/common_gmres_kernels.dp.inc

@@ -72,12 +72,12 @@ void calculate_sin_and_cos_kernel(size_type col_idx, size_type num_cols,
         register_cos = zero<ValueType>();
         register_sin = one<ValueType>();
     } else {
-        const auto scale = std::abs(this_hess) + std::abs(next_hess);
+        const auto scale = gko::abs(this_hess) + gko::abs(next_hess);
         const auto hypotenuse =
             scale *
-            std::sqrt(
-                std::abs(this_hess / scale) * std::abs(this_hess / scale) +
-                std::abs(next_hess / scale) * std::abs(next_hess / scale));
+            gko::sqrt(
+                gko::abs(this_hess / scale) * gko::abs(this_hess / scale) +
+                gko::abs(next_hess / scale) * gko::abs(next_hess / scale));
         register_cos = conj(this_hess) / hypotenuse;
         register_sin = conj(next_hess) / hypotenuse;
     }
@@ -102,7 +102,7 @@ void calculate_residual_norm_kernel(size_type col_idx, size_type num_cols,
     const auto next_rnc = -conj(register_sin) * this_rnc;
     residual_norm_collection[iter * stride_residual_norm_collection + col_idx] =
         register_cos * this_rnc;
-    residual_norm[col_idx] = std::abs(next_rnc);
+    residual_norm[col_idx] = gko::abs(next_rnc);
     residual_norm_collection[(iter + 1) * stride_residual_norm_collection +
                              col_idx] = next_rnc;
 }

dpcpp/solver/idr_kernels.dp.cpp

@@ -542,8 +542,8 @@ void compute_omega_kernel(
     if (!stop_status[global_id].has_stopped()) {
         auto thr = omega[global_id];
         omega[global_id] /= tht[global_id];
-        auto absrho = std::abs(
-            thr / (std::sqrt(real(tht[global_id])) * residual_norm[global_id]));
+        auto absrho = gko::abs(
+            thr / (gko::sqrt(real(tht[global_id])) * residual_norm[global_id]));
 
         if (absrho < kappa) {
             omega[global_id] *= kappa / absrho;

dpcpp/stop/residual_norm_kernels.dp.cpp

@@ -108,7 +108,7 @@ void implicit_residual_norm(
         cgh.parallel_for(
             sycl::range<1>{tau->get_size()[1]}, [=](sycl::id<1> idx_id) {
                 const auto tidx = idx_id[0];
-                if (std::sqrt(std::abs(tau_val[tidx])) <=
+                if (gko::sqrt(gko::abs(tau_val[tidx])) <=
                     rel_residual_goal * orig_tau_val[tidx]) {
                     stop_status_val[tidx].converge(stoppingId, setFinalized);
                     device_storage_val[1] = true;

include/ginkgo/core/base/math.hpp

@@ -30,28 +30,40 @@ class complex;
 
 
 }
-namespace std {
 
 
-inline gko::half abs(gko::half a) { return gko::half((a > 0) ? a : -a); }
+// when using gko, abs will be ambiguous. delete that, get_relative_error can
+// not find proper half
+namespace gko {
+using std::abs;
+using std::sqrt;
 
-inline gko::half abs(std::complex<gko::half> a)
+GKO_ATTRIBUTES GKO_INLINE gko::half abs(gko::half a)
+{
+    return gko::half((a > 0) ? a : -a);
+}
+
+GKO_ATTRIBUTES GKO_INLINE gko::half abs(std::complex<gko::half> a)
 {
     // Using float abs not sqrt on norm to avoid overflow
     return gko::half(abs(std::complex<float>(a)));
 }
 
 
-inline gko::half sqrt(gko::half a) { return gko::half(sqrt(float(a))); }
+GKO_ATTRIBUTES GKO_INLINE gko::half sqrt(gko::half a)
+{
+    return gko::half(std::sqrt(float(a)));
+}
 
-inline std::complex<gko::half> sqrt(std::complex<gko::half> a)
+GKO_ATTRIBUTES GKO_INLINE std::complex<gko::half> sqrt(
+    std::complex<gko::half> a)
 {
     return std::complex<gko::half>(sqrt(std::complex<float>(
         static_cast<float>(a.real()), static_cast<float>(a.imag()))));
 }
 
 
-}  // namespace std
+}  // namespace gko
 
 
 namespace gko {

reference/test/solver/gcr_kernels.cpp

@@ -458,7 +458,7 @@ TYPED_TEST(Gcr, SolveWithImplicitResNormCritIsDisabled)
 template <typename T>
 gko::remove_complex<T> infNorm(gko::matrix::Dense<T>* mat, size_t col = 0)
 {
-    using std::abs;
+    using gko::abs;
     using no_cpx_t = gko::remove_complex<T>;
     no_cpx_t norm = 0.0;
     for (size_t i = 0; i < mat->get_size()[0]; ++i) {

test/base/device_matrix_data_kernels.cpp

@@ -316,7 +316,7 @@ TYPED_TEST(DeviceMatrixData, SumsDuplicates)
     arrays.values.set_executor(this->exec->get_master());
     for (int i = 0; i < arrays.values.get_size(); i++) {
         max_error = std::max<double>(
-            max_error, std::abs(arrays.values.get_const_data()[i] -
+            max_error, gko::abs(arrays.values.get_const_data()[i] -
                                 ref_arrays.values.get_const_data()[i]));
     }
     // when Hip with GNU < 7, it will give a little difference.

test/solver/cb_gmres_kernels.cpp

@@ -146,7 +146,7 @@ class CbGmres : public CommonTestFixture {
         auto& krylov_bases = range_helper.get_bases();
         d_to_host = d_range_helper.get_bases();
         const auto tolerance = r<storage_type>::value;
-        using std::abs;
+        using gko::abs;
         for (gko::size_type i = 0; i < krylov_bases.get_size(); ++i) {
             const auto ref_value = krylov_bases.get_const_data()[i];
             const auto dev_value = d_to_host.get_const_data()[i];

test/test_install/test_install.cpp

@@ -25,7 +25,7 @@ void assert_similar_matrices(gko::ptr_param<const gko::matrix::Dense<>> m1,
     assert(m1->get_size()[1] == m2->get_size()[1]);
     for (gko::size_type i = 0; i < m1->get_size()[0]; ++i) {
         for (gko::size_type j = 0; j < m2->get_size()[1]; ++j) {
-            assert(std::abs(m1->at(i, j) - m2->at(i, j)) < prec);
+            assert(gko::abs(m1->at(i, j) - m2->at(i, j)) < prec);
         }
     }
 }