diff --git a/core/test/utils/assertions.hpp b/core/test/utils/assertions.hpp index 3a275b8ee53..174d4536657 100644 --- a/core/test/utils/assertions.hpp +++ b/core/test/utils/assertions.hpp @@ -259,7 +259,7 @@ void save_matrices_to_disk(Ostream& os, const MatrixData1& first, template 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; diff --git a/core/test/utils/matrix_generator_test.cpp b/core/test/utils/matrix_generator_test.cpp index af268af1471..7c798d2e835 100644 --- a/core/test/utils/matrix_generator_test.cpp +++ b/core/test/utils/matrix_generator_test.cpp @@ -278,8 +278,8 @@ TYPED_TEST(MatrixGenerator, CanGenerateTridiagInverseMatrix) auto lower = gko::test::detail::get_rand_value(dist, engine); auto upper = gko::test::detail::get_rand_value(dist, engine); // make diagonally dominant - auto diag = std::abs(gko::test::detail::get_rand_value(dist, engine)) + - std::abs(lower) + std::abs(upper); + auto diag = gko::abs(gko::test::detail::get_rand_value(dist, engine)) + + gko::abs(lower) + gko::abs(upper); gko::size_type size = 50; if (std::is_same>::value) { // half precision can only handle small matrix diff --git a/dpcpp/factorization/par_ilut_filter_kernels.hpp.inc b/dpcpp/factorization/par_ilut_filter_kernels.hpp.inc index d2345848d1f..6081bc0f417 100644 --- a/dpcpp/factorization/par_ilut_filter_kernels.hpp.inc +++ b/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); } diff --git a/dpcpp/factorization/par_ilut_select_kernels.hpp.inc b/dpcpp/factorization/par_ilut_select_kernels.hpp.inc index 1ebfe6ed320..430bf650e07 100644 --- a/dpcpp/factorization/par_ilut_select_kernels.hpp.inc +++ b/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(lidx * stride)]; - samples[i] = std::abs(val); + samples[i] = gko::abs(val); } bitonic_sort(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(&*counter, IndexType{found}); if (found) { - output[ofs] = std::abs(input[i]); + output[ofs] = gko::abs(input[i]); } } } diff --git a/dpcpp/solver/cb_gmres_kernels.dp.cpp b/dpcpp/solver/cb_gmres_kernels.dp.cpp index 8747dcb60a7..10e5996ef4f 100644 --- a/dpcpp/solver/cb_gmres_kernels.dp.cpp +++ b/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)); } } } diff --git a/dpcpp/solver/common_gmres_kernels.dp.inc b/dpcpp/solver/common_gmres_kernels.dp.inc index 0b5de8188f2..f8a54fe5116 100644 --- a/dpcpp/solver/common_gmres_kernels.dp.inc +++ b/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(); register_sin = one(); } 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; } diff --git a/dpcpp/solver/idr_kernels.dp.cpp b/dpcpp/solver/idr_kernels.dp.cpp index a5531f2dc40..e956eeb428f 100644 --- a/dpcpp/solver/idr_kernels.dp.cpp +++ b/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; diff --git a/dpcpp/stop/residual_norm_kernels.dp.cpp b/dpcpp/stop/residual_norm_kernels.dp.cpp index ddb617a1a84..e67b599dbb5 100644 --- a/dpcpp/stop/residual_norm_kernels.dp.cpp +++ b/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; diff --git a/include/ginkgo/core/base/math.hpp b/include/ginkgo/core/base/math.hpp index 7d45c58caba..67483a5ea2a 100644 --- a/include/ginkgo/core/base/math.hpp +++ b/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 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 a) { // Using float abs not sqrt on norm to avoid overflow return gko::half(abs(std::complex(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 sqrt(std::complex a) +GKO_ATTRIBUTES GKO_INLINE std::complex sqrt( + std::complex a) { return std::complex(sqrt(std::complex( static_cast(a.real()), static_cast(a.imag())))); } -} // namespace std +} // namespace gko namespace gko { diff --git a/reference/test/solver/gcr_kernels.cpp b/reference/test/solver/gcr_kernels.cpp index d3f29fffee9..6077af49af6 100644 --- a/reference/test/solver/gcr_kernels.cpp +++ b/reference/test/solver/gcr_kernels.cpp @@ -458,7 +458,7 @@ TYPED_TEST(Gcr, SolveWithImplicitResNormCritIsDisabled) template gko::remove_complex infNorm(gko::matrix::Dense* mat, size_t col = 0) { - using std::abs; + using gko::abs; using no_cpx_t = gko::remove_complex; no_cpx_t norm = 0.0; for (size_t i = 0; i < mat->get_size()[0]; ++i) { diff --git a/test/base/device_matrix_data_kernels.cpp b/test/base/device_matrix_data_kernels.cpp index 9f2abd2730e..5f1633aa268 100644 --- a/test/base/device_matrix_data_kernels.cpp +++ b/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( - 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. diff --git a/test/solver/cb_gmres_kernels.cpp b/test/solver/cb_gmres_kernels.cpp index 022899d21e6..98eb295091b 100644 --- a/test/solver/cb_gmres_kernels.cpp +++ b/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::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]; diff --git a/test/test_install/test_install.cpp b/test/test_install/test_install.cpp index 2f4cdeda6e4..3d990d3e5eb 100644 --- a/test/test_install/test_install.cpp +++ b/test/test_install/test_install.cpp @@ -25,7 +25,7 @@ void assert_similar_matrices(gko::ptr_param> 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); } } }