diff --git a/common/matrix/csr_kernels.hpp.inc b/common/matrix/csr_kernels.hpp.inc index 0ee4c34dad6..c800be4c826 100644 --- a/common/matrix/csr_kernels.hpp.inc +++ b/common/matrix/csr_kernels.hpp.inc @@ -901,6 +901,18 @@ __global__ __launch_bounds__(default_block_size) void check_unsorted( } +template +__global__ __launch_bounds__(default_block_size) void absolute_kernel( + size_type num, const ValueType *__restrict__ source, + remove_complex *__restrict__ result) +{ + const auto tidx = thread::get_thread_id_flat(); + if (tidx < num) { + result[tidx] = gko::abs(source[tidx]); + } +} + + } // namespace kernel diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index 6b9e8b4409e..b70dfed646a 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -615,6 +615,11 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_IS_SORTED_BY_COLUMN_INDEX); +template +GKO_DECLARE_CSR_ABSOLUTE(ValueType, IndexType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_ABSOLUTE); + } // namespace csr @@ -995,19 +1000,19 @@ GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE(GKO_DECLARE_AMGX_PGM_RENUMBER_KERNEL); template GKO_DECLARE_AMGX_PGM_EXTRACT_DIAG(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_EXTRACT_DIAG); template GKO_DECLARE_AMGX_PGM_FIND_STRONGEST_NEIGHBOR(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_FIND_STRONGEST_NEIGHBOR); template GKO_DECLARE_AMGX_PGM_ASSIGN_TO_EXIST_AGG(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_ASSIGN_TO_EXIST_AGG); template diff --git a/core/matrix/csr.cpp b/core/matrix/csr.cpp index 2d0012c7cb6..06ef2ea6e32 100644 --- a/core/matrix/csr.cpp +++ b/core/matrix/csr.cpp @@ -78,7 +78,7 @@ GKO_REGISTER_OPERATION(calculate_nonzeros_per_row, GKO_REGISTER_OPERATION(sort_by_column_index, csr::sort_by_column_index); GKO_REGISTER_OPERATION(is_sorted_by_column_index, csr::is_sorted_by_column_index); - +GKO_REGISTER_OPERATION(absolute, csr::absolute); } // namespace csr @@ -464,6 +464,22 @@ bool Csr::is_sorted_by_column_index() const } +template +std::unique_ptr Csr::absolute() const +{ + using abs_type = remove_complex; + using abs_csr = Csr; + auto exec = this->get_executor(); + auto abs_cpy = abs_csr::create(exec, this->get_size(), + this->get_num_stored_elements()); + abs_cpy->col_idxs_ = this->col_idxs_; + abs_cpy->row_ptrs_ = this->row_ptrs_; + exec->run(csr::make_absolute(this, lend(abs_cpy))); + this->convert_strategy_helper(lend(abs_cpy)); + return std::move(abs_cpy); +} + + #define GKO_DECLARE_CSR_MATRIX(ValueType, IndexType) \ class Csr GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_MATRIX); diff --git a/core/matrix/csr_kernels.hpp b/core/matrix/csr_kernels.hpp index f901c0f0952..82e04a9a255 100644 --- a/core/matrix/csr_kernels.hpp +++ b/core/matrix/csr_kernels.hpp @@ -174,6 +174,11 @@ namespace kernels { std::shared_ptr exec, \ const matrix::Csr *to_check, bool *is_sorted) +#define GKO_DECLARE_CSR_ABSOLUTE(ValueType, IndexType) \ + void absolute(std::shared_ptr exec, \ + const matrix::Csr *source, \ + matrix::Csr, IndexType> *result) + #define GKO_DECLARE_ALL_AS_TEMPLATES \ template \ GKO_DECLARE_CSR_SPMV_KERNEL(ValueType, IndexType); \ @@ -216,7 +221,9 @@ namespace kernels { template \ GKO_DECLARE_CSR_SORT_BY_COLUMN_INDEX(ValueType, IndexType); \ template \ - GKO_DECLARE_CSR_IS_SORTED_BY_COLUMN_INDEX(ValueType, IndexType) + GKO_DECLARE_CSR_IS_SORTED_BY_COLUMN_INDEX(ValueType, IndexType); \ + template \ + GKO_DECLARE_CSR_ABSOLUTE(ValueType, IndexType) namespace omp { diff --git a/core/multigrid/amgx_pgm.cpp b/core/multigrid/amgx_pgm.cpp index 4a2f9fdf56a..5a12d08fc80 100644 --- a/core/multigrid/amgx_pgm.cpp +++ b/core/multigrid/amgx_pgm.cpp @@ -41,6 +41,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "core/matrix/csr_builder.hpp" @@ -92,23 +93,35 @@ template void AmgxPgm::generate() { using matrix_type = matrix::Csr; + using rmc_value_type = remove_complex; + using weight_matrix_type = matrix::Csr; auto exec = this->get_executor(); const auto num = this->system_matrix_->get_size()[0]; - Array diag(this->get_executor(), num); + Array diag(this->get_executor(), num); Array strongest_neighbor(this->get_executor(), num); Array intermediate_agg(this->get_executor(), parameters_.deterministic * num); const auto amgxpgm_op = gko::as(this->system_matrix_.get()); // Initial agg = -1 exec->run(amgx_pgm::make_initial(agg_)); - // Extract the diagonal value of matrix - exec->run(amgx_pgm::make_extract_diag(amgxpgm_op, diag)); size_type num_unagg{0}; size_type num_unagg_prev{0}; + // TODO: if mtx is a hermitian matrix, weight_mtx = abs(mtx) + // compute weight_mtx = (abs(mtx) + abs(mtx'))/2; + auto abs_mtx = gko::as(amgxpgm_op->absolute()); + // abs_mtx is already real valuetype, so transpose is enough + auto weight_mtx = gko::as(abs_mtx->transpose()); + auto half_scaler = initialize>({0.5}, exec); + auto identity = matrix::Identity::create(exec, num); + // W = (abs_mtx + transpose(abs_mtx))/2 + abs_mtx->apply(lend(half_scaler), lend(identity), lend(half_scaler), + lend(weight_mtx)); + // Extract the diagonal value of matrix + exec->run(amgx_pgm::make_extract_diag(weight_mtx.get(), diag)); for (int i = 0; i < parameters_.max_iterations; i++) { // Find the strongest neighbor of each row - exec->run(amgx_pgm::make_find_strongest_neighbor(amgxpgm_op, diag, agg_, - strongest_neighbor)); + exec->run(amgx_pgm::make_find_strongest_neighbor( + weight_mtx.get(), diag, agg_, strongest_neighbor)); // Match edges exec->run(amgx_pgm::make_match_edge(strongest_neighbor, agg_)); // Get the num_unagg @@ -116,8 +129,7 @@ void AmgxPgm::generate() // no new match, all match, or the ratio of num_unagg/num is lower // than parameter.max_unassigned_percentage if (num_unagg == 0 || num_unagg == num_unagg_prev || - static_cast(num_unagg) / num < - parameters_.max_unassigned_percentage) { + num_unagg < parameters_.max_unassigned_percentage * num) { break; } num_unagg_prev = num_unagg; @@ -128,8 +140,8 @@ void AmgxPgm::generate() intermediate_agg = agg_; } while (num_unagg != 0) { - exec->run(amgx_pgm::make_assign_to_exist_agg(amgxpgm_op, diag, agg_, - intermediate_agg)); + exec->run(amgx_pgm::make_assign_to_exist_agg(weight_mtx.get(), diag, + agg_, intermediate_agg)); exec->run(amgx_pgm::make_count_unagg(agg_, &num_unagg)); } size_type num_agg; diff --git a/core/multigrid/amgx_pgm_kernels.hpp b/core/multigrid/amgx_pgm_kernels.hpp index 26031d570fd..b3ea2f3d4bc 100644 --- a/core/multigrid/amgx_pgm_kernels.hpp +++ b/core/multigrid/amgx_pgm_kernels.hpp @@ -85,16 +85,16 @@ namespace amgx_pgm { #define GKO_DECLARE_AMGX_PGM_FIND_STRONGEST_NEIGHBOR(ValueType, IndexType) \ void find_strongest_neighbor( \ std::shared_ptr exec, \ - const matrix::Csr *source, \ + const matrix::Csr *weight_mtx, \ const Array &diag, Array &agg, \ Array &strongest_neighbor) -#define GKO_DECLARE_AMGX_PGM_ASSIGN_TO_EXIST_AGG(ValueType, IndexType) \ - void assign_to_exist_agg(std::shared_ptr exec, \ - const matrix::Csr *source, \ - const Array &diag, \ - Array &agg, \ - Array &intermediate_agg) +#define GKO_DECLARE_AMGX_PGM_ASSIGN_TO_EXIST_AGG(ValueType, IndexType) \ + void assign_to_exist_agg( \ + std::shared_ptr exec, \ + const matrix::Csr *weight_mtx, \ + const Array &diag, Array &agg, \ + Array &intermediate_agg) #define GKO_DECLARE_AMGX_PGM_GENERATE(ValueType, IndexType) \ void amgx_pgm_generate(std::shared_ptr exec, \ diff --git a/cuda/matrix/csr_kernels.cu b/cuda/matrix/csr_kernels.cu index 3229aab74c8..44b62a23406 100644 --- a/cuda/matrix/csr_kernels.cu +++ b/cuda/matrix/csr_kernels.cu @@ -635,16 +635,17 @@ void spgeam(std::shared_ptr exec, auto total_nnz = a->get_num_stored_elements() + b->get_num_stored_elements(); auto nnz_per_row = total_nnz / a->get_size()[0]; - select_spgeam(spgeam_kernels(), - [&](int compiled_subwarp_size) { - return compiled_subwarp_size >= nnz_per_row || - compiled_subwarp_size == config::warp_size; - }, - syn::value_list(), syn::type_list<>(), exec, - alpha->get_const_values(), a->get_const_row_ptrs(), - a->get_const_col_idxs(), a->get_const_values(), - beta->get_const_values(), b->get_const_row_ptrs(), - b->get_const_col_idxs(), b->get_const_values(), c); + select_spgeam( + spgeam_kernels(), + [&](int compiled_subwarp_size) { + return compiled_subwarp_size >= nnz_per_row || + compiled_subwarp_size == config::warp_size; + }, + syn::value_list(), syn::type_list<>(), exec, + alpha->get_const_values(), a->get_const_row_ptrs(), + a->get_const_col_idxs(), a->get_const_values(), + beta->get_const_values(), b->get_const_row_ptrs(), + b->get_const_col_idxs(), b->get_const_values(), c); } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_SPGEAM_KERNEL); @@ -1101,6 +1102,23 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_IS_SORTED_BY_COLUMN_INDEX); +template +void absolute(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Csr, IndexType> *result) +{ + auto result_val = result->get_values(); + auto source_val = source->get_const_values(); + const auto num = source->get_num_stored_elements(); + const dim3 grid(ceildiv(num, default_block_size)); + + kernel::absolute_kernel<<>>(num, source_val, + result_val); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_ABSOLUTE); + + } // namespace csr } // namespace cuda } // namespace kernels diff --git a/cuda/multigrid/amgx_pgm_kernels.cu b/cuda/multigrid/amgx_pgm_kernels.cu index 310883ecae3..bec3b7aa840 100644 --- a/cuda/multigrid/amgx_pgm_kernels.cu +++ b/cuda/multigrid/amgx_pgm_kernels.cu @@ -120,34 +120,34 @@ void extract_diag(std::shared_ptr exec, GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_EXTRACT_DIAG); template -void find_strongest_neighbor(std::shared_ptr exec, - const matrix::Csr *source, - const Array &diag, - Array &agg, - Array &strongest_neighbor) +void find_strongest_neighbor( + std::shared_ptr exec, + const matrix::Csr *weight_mtx, + const Array &diag, Array &agg, + Array &strongest_neighbor) { GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_FIND_STRONGEST_NEIGHBOR); template void assign_to_exist_agg(std::shared_ptr exec, - const matrix::Csr *source, + const matrix::Csr *weight_mtx, const Array &diag, Array &agg, Array &intermediate_agg) { GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_ASSIGN_TO_EXIST_AGG); diff --git a/cuda/test/matrix/csr_kernels.cpp b/cuda/test/matrix/csr_kernels.cpp index 39608505cbd..4928e5b122d 100644 --- a/cuda/test/matrix/csr_kernels.cpp +++ b/cuda/test/matrix/csr_kernels.cpp @@ -60,6 +60,7 @@ namespace { class Csr : public ::testing::Test { protected: using Mtx = gko::matrix::Csr<>; + using AbsMtx = gko::matrix::Csr<>; using Vec = gko::matrix::Dense<>; using ComplexVec = gko::matrix::Dense>; using ComplexMtx = gko::matrix::Csr>; @@ -725,4 +726,16 @@ TEST_F(Csr, OneAutomaticalWorksWithDifferentMatrices) } +TEST_F(Csr, AbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); + + auto abs_mtx = mtx->absolute(); + auto dabs_mtx = dmtx->absolute(); + + GKO_ASSERT_MTX_NEAR(static_cast(abs_mtx.get()), + static_cast(dabs_mtx.get()), 1e-14); +} + + } // namespace diff --git a/hip/matrix/csr_kernels.hip.cpp b/hip/matrix/csr_kernels.hip.cpp index bf03ec4c7b7..c0eb5946558 100644 --- a/hip/matrix/csr_kernels.hip.cpp +++ b/hip/matrix/csr_kernels.hip.cpp @@ -651,15 +651,16 @@ void advanced_spgemm(std::shared_ptr exec, auto total_nnz = c_nnz + d->get_num_stored_elements(); auto nnz_per_row = total_nnz / m; - select_spgeam(spgeam_kernels(), - [&](int compiled_subwarp_size) { - return compiled_subwarp_size >= nnz_per_row || - compiled_subwarp_size == config::warp_size; - }, - syn::value_list(), syn::type_list<>(), exec, - alpha->get_const_values(), c_tmp_row_ptrs, c_tmp_col_idxs, - c_tmp_vals, beta->get_const_values(), d_row_ptrs, - d_col_idxs, d_vals, c); + select_spgeam( + spgeam_kernels(), + [&](int compiled_subwarp_size) { + return compiled_subwarp_size >= nnz_per_row || + compiled_subwarp_size == config::warp_size; + }, + syn::value_list(), syn::type_list<>(), exec, + alpha->get_const_values(), c_tmp_row_ptrs, c_tmp_col_idxs, + c_tmp_vals, beta->get_const_values(), d_row_ptrs, d_col_idxs, + d_vals, c); } else { GKO_NOT_IMPLEMENTED; } @@ -680,16 +681,17 @@ void spgeam(std::shared_ptr exec, auto total_nnz = a->get_num_stored_elements() + b->get_num_stored_elements(); auto nnz_per_row = total_nnz / a->get_size()[0]; - select_spgeam(spgeam_kernels(), - [&](int compiled_subwarp_size) { - return compiled_subwarp_size >= nnz_per_row || - compiled_subwarp_size == config::warp_size; - }, - syn::value_list(), syn::type_list<>(), exec, - alpha->get_const_values(), a->get_const_row_ptrs(), - a->get_const_col_idxs(), a->get_const_values(), - beta->get_const_values(), b->get_const_row_ptrs(), - b->get_const_col_idxs(), b->get_const_values(), c); + select_spgeam( + spgeam_kernels(), + [&](int compiled_subwarp_size) { + return compiled_subwarp_size >= nnz_per_row || + compiled_subwarp_size == config::warp_size; + }, + syn::value_list(), syn::type_list<>(), exec, + alpha->get_const_values(), a->get_const_row_ptrs(), + a->get_const_col_idxs(), a->get_const_values(), + beta->get_const_values(), b->get_const_row_ptrs(), + b->get_const_col_idxs(), b->get_const_values(), c); } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_SPGEAM_KERNEL); @@ -1166,6 +1168,24 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_IS_SORTED_BY_COLUMN_INDEX); +template +void absolute(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Csr, IndexType> *result) +{ + auto result_val = result->get_values(); + auto source_val = source->get_const_values(); + const auto num = source->get_num_stored_elements(); + const dim3 grid(ceildiv(num, default_block_size)); + + hipLaunchKernelGGL(HIP_KERNEL_NAME(kernel::absolute_kernel), dim3(grid), + dim3(default_block_size), 0, 0, num, source_val, + result_val); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_ABSOLUTE); + + } // namespace csr } // namespace hip } // namespace kernels diff --git a/hip/multigrid/amgx_pgm_kernels.hip.cpp b/hip/multigrid/amgx_pgm_kernels.hip.cpp index 0ff2784258b..24636180253 100644 --- a/hip/multigrid/amgx_pgm_kernels.hip.cpp +++ b/hip/multigrid/amgx_pgm_kernels.hip.cpp @@ -120,34 +120,34 @@ void extract_diag(std::shared_ptr exec, GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_EXTRACT_DIAG); template -void find_strongest_neighbor(std::shared_ptr exec, - const matrix::Csr *source, - const Array &diag, - Array &agg, - Array &strongest_neighbor) +void find_strongest_neighbor( + std::shared_ptr exec, + const matrix::Csr *weight_mtx, + const Array &diag, Array &agg, + Array &strongest_neighbor) { GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_FIND_STRONGEST_NEIGHBOR); template void assign_to_exist_agg(std::shared_ptr exec, - const matrix::Csr *source, + const matrix::Csr *weight_mtx, const Array &diag, Array &agg, Array &intermediate_agg) { GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_ASSIGN_TO_EXIST_AGG); diff --git a/hip/test/matrix/csr_kernels.hip.cpp b/hip/test/matrix/csr_kernels.hip.cpp index a1b2adfd794..64de57d329d 100644 --- a/hip/test/matrix/csr_kernels.hip.cpp +++ b/hip/test/matrix/csr_kernels.hip.cpp @@ -60,6 +60,7 @@ namespace { class Csr : public ::testing::Test { protected: using Mtx = gko::matrix::Csr<>; + using AbsMtx = gko::matrix::Csr<>; using Vec = gko::matrix::Dense<>; Csr() : mtx_size(532, 231), rand_engine(42) {} @@ -698,4 +699,16 @@ TEST_F(Csr, OneAutomaticalWorksWithDifferentMatrices) } +TEST_F(Csr, AbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); + + auto abs_mtx = mtx->absolute(); + auto dabs_mtx = dmtx->absolute(); + + GKO_ASSERT_MTX_NEAR(static_cast(abs_mtx.get()), + static_cast(dabs_mtx.get()), 1e-14); +} + + } // namespace diff --git a/include/ginkgo/core/base/math.hpp b/include/ginkgo/core/base/math.hpp index 1b7a2f9f18f..97a1eec80bb 100644 --- a/include/ginkgo/core/base/math.hpp +++ b/include/ginkgo/core/base/math.hpp @@ -132,6 +132,21 @@ struct remove_complex_impl> { using type = T; }; +/** + * Use the complex type if it is not complex. + */ +template +struct add_complex_impl { + using type = std::complex; +}; + +/** + * Use the same type if it is complex type. + */ +template +struct add_complex_impl> { + using type = std::complex; +}; template struct is_complex_impl : public std::integral_constant {}; @@ -175,6 +190,14 @@ template using remove_complex = typename detail::remove_complex_impl::type; +/** + * Obtains a complex type of non-complex type, and leaves the type + * unchanged if it is a complex type. + */ +template +using add_complex = typename detail::add_complex_impl::type; + + /** * Allows to check if T is a complex value during compile time by accessing the * `value` attribute of this struct. @@ -595,22 +618,6 @@ GKO_INLINE GKO_ATTRIBUTES constexpr T one(const T &) #undef GKO_BIND_ZERO_ONE -/** - * Returns the absolute value of the object. - * - * @tparam T the type of the object - * - * @param x the object - * - * @return x >= zero() ? x : -x; - */ -template -GKO_INLINE GKO_ATTRIBUTES constexpr T abs(const T &x) -{ - return x >= zero() ? x : -x; -} - - /** * Returns the larger of the arguments. * @@ -740,6 +747,33 @@ GKO_INLINE GKO_ATTRIBUTES constexpr auto squared_norm(const T &x) } +/** + * Returns the absolute value of the object. + * + * @tparam T the type of the object + * + * @param x the object + * + * @return x >= zero() ? x : -x; + */ +template +GKO_INLINE + GKO_ATTRIBUTES constexpr xstd::enable_if_t::value, T> + abs(const T &x) +{ + return x >= zero() ? x : -x; +} + + +template +GKO_INLINE GKO_ATTRIBUTES constexpr xstd::enable_if_t::value, + remove_complex> +abs(const T &x) +{ + return squared_norm(x); +} + + /** * Returns the position of the most significant bit of the number. * diff --git a/include/ginkgo/core/base/types.hpp b/include/ginkgo/core/base/types.hpp index 6c2ab2a50d2..85c883f3821 100644 --- a/include/ginkgo/core/base/types.hpp +++ b/include/ginkgo/core/base/types.hpp @@ -435,6 +435,22 @@ GKO_ATTRIBUTES constexpr bool operator!=(precision_reduction x, template _macro(int64) +/** + * Instantiates a template for each non-complex value and index type compiled by + * Ginkgo. + * + * @param _macro A macro which expands the template instantiation + * (not including the leading `template` specifier). + * Should take two arguments, which are replaced by the + * value and index types. + */ +#define GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE(_macro) \ + template _macro(float, int32); \ + template _macro(double, int32); \ + template _macro(float, int64); \ + template _macro(double, int64) + + /** * Instantiates a template for each value and index type compiled by Ginkgo. * @@ -443,14 +459,11 @@ GKO_ATTRIBUTES constexpr bool operator!=(precision_reduction x, * Should take two arguments, which are replaced by the * value and index types. */ -#define GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(_macro) \ - template _macro(float, int32); \ - template _macro(double, int32); \ - template _macro(std::complex, int32); \ - template _macro(std::complex, int32); \ - template _macro(float, int64); \ - template _macro(double, int64); \ - template _macro(std::complex, int64); \ +#define GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(_macro) \ + GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE(_macro); \ + template _macro(std::complex, int32); \ + template _macro(std::complex, int32); \ + template _macro(std::complex, int64); \ template _macro(std::complex, int64) diff --git a/include/ginkgo/core/matrix/csr.hpp b/include/ginkgo/core/matrix/csr.hpp index 227b6b06a28..62d88bbedf0 100644 --- a/include/ginkgo/core/matrix/csr.hpp +++ b/include/ginkgo/core/matrix/csr.hpp @@ -36,6 +36,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include +#include namespace gko { @@ -643,6 +644,7 @@ class Csr : public EnableLinOp>, detail::strategy_rebuild_helper(result); } } + friend class Csr, IndexType>; void convert_to( @@ -694,6 +696,15 @@ class Csr : public EnableLinOp>, std::unique_ptr inverse_column_permute( const Array *inverse_permutation_indices) const override; + friend class Csr, IndexType>; + + /** + * Return the matrix whose value are absolute of original matirx + * + * @return the absolute matrix + */ + std::unique_ptr absolute() const; + /** * Sorts all (value, col_idx) pairs in each row by column index */ diff --git a/omp/matrix/csr_kernels.cpp b/omp/matrix/csr_kernels.cpp index e4cb16d23a3..24fcb7c8421 100644 --- a/omp/matrix/csr_kernels.cpp +++ b/omp/matrix/csr_kernels.cpp @@ -348,14 +348,15 @@ void spgeam(std::shared_ptr exec, auto c_col_idxs = c_col_idxs_array.get_data(); auto c_vals = c_vals_array.get_data(); - abstract_spgeam(a, b, [&](IndexType row) { return c_row_ptrs[row]; }, - [&](IndexType, IndexType col, ValueType a_val, - ValueType b_val, IndexType &nz) { - c_vals[nz] = valpha * a_val + vbeta * b_val; - c_col_idxs[nz] = col; - ++nz; - }, - [](IndexType, IndexType) {}); + abstract_spgeam( + a, b, [&](IndexType row) { return c_row_ptrs[row]; }, + [&](IndexType, IndexType col, ValueType a_val, ValueType b_val, + IndexType &nz) { + c_vals[nz] = valpha * a_val + vbeta * b_val; + c_col_idxs[nz] = col; + ++nz; + }, + [](IndexType, IndexType) {}); } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_SPGEAM_KERNEL); @@ -804,6 +805,22 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_IS_SORTED_BY_COLUMN_INDEX); +template +void absolute(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Csr, IndexType> *result) +{ + auto result_val = result->get_values(); + auto source_val = source->get_const_values(); +#pragma omp parallel for + for (size_type i = 0; i < source->get_num_stored_elements(); i++) { + result_val[i] = abs(source_val[i]); + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_ABSOLUTE); + + } // namespace csr } // namespace omp } // namespace kernels diff --git a/omp/multigrid/amgx_pgm_kernels.cpp b/omp/multigrid/amgx_pgm_kernels.cpp index 982260d193b..502531354ad 100644 --- a/omp/multigrid/amgx_pgm_kernels.cpp +++ b/omp/multigrid/amgx_pgm_kernels.cpp @@ -117,34 +117,34 @@ void extract_diag(std::shared_ptr exec, GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_EXTRACT_DIAG); template -void find_strongest_neighbor(std::shared_ptr exec, - const matrix::Csr *source, - const Array &diag, - Array &agg, - Array &strongest_neighbor) +void find_strongest_neighbor( + std::shared_ptr exec, + const matrix::Csr *weight_mtx, + const Array &diag, Array &agg, + Array &strongest_neighbor) { GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_FIND_STRONGEST_NEIGHBOR); template void assign_to_exist_agg(std::shared_ptr exec, - const matrix::Csr *source, + const matrix::Csr *weight_mtx, const Array &diag, Array &agg, Array &intermediate_agg) { GKO_NOT_IMPLEMENTED; } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_ASSIGN_TO_EXIST_AGG); diff --git a/omp/test/matrix/csr_kernels.cpp b/omp/test/matrix/csr_kernels.cpp index bb607efd615..a00d1ddd144 100644 --- a/omp/test/matrix/csr_kernels.cpp +++ b/omp/test/matrix/csr_kernels.cpp @@ -61,6 +61,7 @@ class Csr : public ::testing::Test { protected: using Arr = gko::Array; using Mtx = gko::matrix::Csr<>; + using AbsMtx = gko::matrix::Csr<>; using Vec = gko::matrix::Dense<>; using ComplexVec = gko::matrix::Dense>; using ComplexMtx = gko::matrix::Csr>; @@ -534,4 +535,16 @@ TEST_F(Csr, SortUnsortedMatrixIsEquivalentToRef) } +TEST_F(Csr, AbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_mtx = mtx->absolute(); + auto dabs_mtx = dmtx->absolute(); + + GKO_ASSERT_MTX_NEAR(static_cast(abs_mtx.get()), + static_cast(dabs_mtx.get()), 1e-14); +} + + } // namespace diff --git a/reference/matrix/csr_kernels.cpp b/reference/matrix/csr_kernels.cpp index ce4e5471c66..a48c17c19b2 100644 --- a/reference/matrix/csr_kernels.cpp +++ b/reference/matrix/csr_kernels.cpp @@ -342,14 +342,15 @@ void spgeam(std::shared_ptr exec, auto c_col_idxs = c_col_idxs_array.get_data(); auto c_vals = c_vals_array.get_data(); - abstract_spgeam(a, b, [&](IndexType row) { return c_row_ptrs[row]; }, - [&](IndexType, IndexType col, ValueType a_val, - ValueType b_val, IndexType &nz) { - c_vals[nz] = valpha * a_val + vbeta * b_val; - c_col_idxs[nz] = col; - ++nz; - }, - [](IndexType, IndexType) {}); + abstract_spgeam( + a, b, [&](IndexType row) { return c_row_ptrs[row]; }, + [&](IndexType, IndexType col, ValueType a_val, ValueType b_val, + IndexType &nz) { + c_vals[nz] = valpha * a_val + vbeta * b_val; + c_col_idxs[nz] = col; + ++nz; + }, + [](IndexType, IndexType) {}); } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_SPGEAM_KERNEL); @@ -881,6 +882,20 @@ void is_sorted_by_column_index( GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_IS_SORTED_BY_COLUMN_INDEX); +template +void absolute(std::shared_ptr exec, + const matrix::Csr *source, + matrix::Csr, IndexType> *result) +{ + auto result_val = result->get_values(); + auto source_val = source->get_const_values(); + for (size_type i = 0; i < source->get_num_stored_elements(); i++) { + result_val[i] = abs(source_val[i]); + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_ABSOLUTE); + } // namespace csr } // namespace reference diff --git a/reference/multigrid/amgx_pgm_kernels.cpp b/reference/multigrid/amgx_pgm_kernels.cpp index 053cb46ec45..6cecc36a076 100644 --- a/reference/multigrid/amgx_pgm_kernels.cpp +++ b/reference/multigrid/amgx_pgm_kernels.cpp @@ -199,43 +199,33 @@ void extract_diag(std::shared_ptr exec, } } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_EXTRACT_DIAG); template -void find_strongest_neighbor(std::shared_ptr exec, - const matrix::Csr *source, - const Array &diag, - Array &agg, - Array &strongest_neighbor) +void find_strongest_neighbor( + std::shared_ptr exec, + const matrix::Csr *weight_mtx, + const Array &diag, Array &agg, + Array &strongest_neighbor) { - const auto row_ptrs = source->get_const_row_ptrs(); - const auto col_idxs = source->get_const_col_idxs(); - const auto vals = source->get_const_values(); + const auto row_ptrs = weight_mtx->get_const_row_ptrs(); + const auto col_idxs = weight_mtx->get_const_col_idxs(); + const auto vals = weight_mtx->get_const_values(); for (size_type row = 0; row < agg.get_num_elems(); row++) { - auto max_weight_unagg = zero>(); - auto max_weight_agg = zero>(); + auto max_weight_unagg = zero(); + auto max_weight_agg = zero(); IndexType strongest_unagg = -1; IndexType strongest_agg = -1; if (agg.get_const_data()[row] == -1) { - for (auto ind = row_ptrs[row]; ind < row_ptrs[row + 1]; ind++) { - auto weight = zero>(); - auto col = col_idxs[ind]; + for (auto idx = row_ptrs[row]; idx < row_ptrs[row + 1]; idx++) { + auto col = col_idxs[idx]; if (col == row) { continue; } - // Handle the unsymmetric values cases - for (auto trans_ind = row_ptrs[col]; - trans_ind < row_ptrs[col + 1]; trans_ind++) { - if (col_idxs[trans_ind] == row) { - weight = 0.5 * (abs(vals[ind]) + abs(vals[trans_ind])) / - max(abs(diag.get_const_data()[row]), - abs(diag.get_const_data()[col])); - break; - } - } - + auto weight = vals[idx] / max(abs(diag.get_const_data()[row]), + abs(diag.get_const_data()[col])); if (agg.get_const_data()[col] == -1 && (weight > max_weight_unagg || (weight == max_weight_unagg && col > strongest_unagg))) { @@ -264,21 +254,21 @@ void find_strongest_neighbor(std::shared_ptr exec, } } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_FIND_STRONGEST_NEIGHBOR); template void assign_to_exist_agg(std::shared_ptr exec, - const matrix::Csr *source, + const matrix::Csr *weight_mtx, const Array &diag, Array &agg, Array &intermediate_agg) { - const auto row_ptrs = source->get_const_row_ptrs(); - const auto col_idxs = source->get_const_col_idxs(); - const auto vals = source->get_const_values(); + const auto row_ptrs = weight_mtx->get_const_row_ptrs(); + const auto col_idxs = weight_mtx->get_const_col_idxs(); + const auto vals = weight_mtx->get_const_values(); const auto diag_vals = diag.get_const_data(); - auto max_weight_agg = zero>(); + auto max_weight_agg = zero(); const auto agg_const_val = agg.get_const_data(); auto agg_val = (intermediate_agg.get_num_elems() > 0) ? intermediate_agg.get_data() @@ -294,8 +284,8 @@ void assign_to_exist_agg(std::shared_ptr exec, if (col == row) { break; } - auto weight = - abs(vals[idx]) / max(abs(diag_vals[row]), abs(diag_vals[col])); + auto weight = vals[idx] / max(abs(diag.get_const_data()[row]), + abs(diag.get_const_data()[col])); if (agg_const_val[col] != -1 && (weight > max_weight_agg || (weight == max_weight_agg && col > strongest_agg))) { @@ -316,7 +306,7 @@ void assign_to_exist_agg(std::shared_ptr exec, } } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( +GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_AND_INDEX_TYPE( GKO_DECLARE_AMGX_PGM_ASSIGN_TO_EXIST_AGG); diff --git a/reference/test/matrix/csr_kernels.cpp b/reference/test/matrix/csr_kernels.cpp index 736f90349ad..1e545d398df 100644 --- a/reference/test/matrix/csr_kernels.cpp +++ b/reference/test/matrix/csr_kernels.cpp @@ -67,6 +67,8 @@ class Csr : public ::testing::Test { typename std::tuple_element<1, decltype(ValueIndexType())>::type; using Coo = gko::matrix::Coo; using Mtx = gko::matrix::Csr; + using AbsMtx = + gko::matrix::Csr, index_type>; using Sellp = gko::matrix::Sellp; using SparsityCsr = gko::matrix::SparsityCsr; using Ell = gko::matrix::Ell; @@ -1291,6 +1293,22 @@ TYPED_TEST(Csr, SortUnsortedMatrix) } +TYPED_TEST(Csr, Absolute) +{ + using Mtx = typename TestFixture::Mtx; + using AbsMtx = typename TestFixture::AbsMtx; + + auto mtx = gko::initialize( + {{1.0, 2.0, -2.0}, {3.0, -5.0, 0.0}, {0.0, 1.0, -1.5}}, this->exec); + auto abs_mtx = mtx->absolute(); + auto abs_mtx_csr = static_cast(abs_mtx.get()); + + GKO_ASSERT_MTX_NEAR(abs_mtx_csr, + l({{1.0, 2.0, 2.0}, {3.0, 5.0, 0.0}, {0.0, 1.0, 1.5}}), + 0.0); +} + + template class CsrComplex : public ::testing::Test { protected: @@ -1308,14 +1326,13 @@ TYPED_TEST(CsrComplex, MtxIsConjugateTransposable) { using Csr = typename TestFixture::Mtx; using T = typename TestFixture::value_type; - using value_type = typename TestFixture::value_type; auto exec = gko::ReferenceExecutor::create(); // clang-format off - auto mtx2 = gko::initialize( - {{T{1.0, 2.0}, T{3.0, 0.0}, T{2.0, 0.0}}, - {T{0.0, 0.0}, T{5.0, - 3.5}, T{0.0,0.0}}, - {T{0.0, 0.0}, T{0.0, 1.5}, T{2.0,0.0}}}, exec); + auto mtx2 = gko::initialize( + {{T{1.0, 2.0}, T{3.0, 0.0}, T{2.0, 0.0}}, + {T{0.0, 0.0}, T{5.0, - 3.5}, T{0.0,0.0}}, + {T{0.0, 0.0}, T{0.0, 1.5}, T{2.0,0.0}}}, exec); // clang-format on auto trans = mtx2->conj_transpose(); @@ -1330,4 +1347,28 @@ TYPED_TEST(CsrComplex, MtxIsConjugateTransposable) } +TYPED_TEST(CsrComplex, Absolute) +{ + using Mtx = typename TestFixture::Mtx; + using T = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using AbsMtx = gko::matrix::Csr, index_type>; + + auto exec = gko::ReferenceExecutor::create(); + // clang-format off + auto mtx = gko::initialize( + {{T{1.0, 0.0}, T{3.0, 4.0}, T{0.0, 2.0}}, + {T{-4.0, -3.0}, T{-1.0, 0}, T{0.0, 0.0}}, + {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); + // clang-format on + + auto abs_mtx = mtx->absolute(); + auto abs_mtx_csr = static_cast(abs_mtx.get()); + + GKO_ASSERT_MTX_NEAR(abs_mtx_csr, + l({{1.0, 5.0, 2.0}, {5.0, 1.0, 0.0}, {0.0, 1.5, 2.0}}), + 0.0); +} + + } // namespace diff --git a/reference/test/multigrid/amgx_pgm_kernels.cpp b/reference/test/multigrid/amgx_pgm_kernels.cpp index bbbdd5b1968..76911118c67 100644 --- a/reference/test/multigrid/amgx_pgm_kernels.cpp +++ b/reference/test/multigrid/amgx_pgm_kernels.cpp @@ -68,6 +68,8 @@ class AmgxPgm : public ::testing::Test { using Vec = gko::matrix::Dense; using RestrictProlong = gko::multigrid::AmgxPgm; using T = value_type; + using rmc_value_type = gko::remove_complex; + using WeightMtx = gko::matrix::Csr; AmgxPgm() : exec(gko::ReferenceExecutor::create()), amgxpgm_factory(RestrictProlong::build() @@ -92,21 +94,29 @@ class AmgxPgm : public ::testing::Test { exec)), mtx(Mtx::create(exec, gko::dim<2>(5, 5), 15, std::make_shared())), + weight(WeightMtx::create( + exec, gko::dim<2>(5, 5), 15, + std::make_shared())), coarse(Mtx::create(exec, gko::dim<2>(2, 2), 4, std::make_shared())), mtx_diag(exec, 5), agg(exec, 5) { - this->create_mtx(mtx.get(), &mtx_diag, &agg, coarse.get()); + this->create_mtx(mtx.get(), weight.get(), &mtx_diag, &agg, + coarse.get()); rstr_prlg = amgxpgm_factory->generate(mtx); } - void create_mtx(Mtx *fine, gko::Array *diag, + void create_mtx(Mtx *fine, WeightMtx *weight, + gko::Array *diag, gko::Array *agg, Mtx *coarse) { auto vals = fine->get_values(); auto cols = fine->get_col_idxs(); auto rows = fine->get_row_ptrs(); + auto w_vals = weight->get_values(); + auto w_cols = weight->get_col_idxs(); + auto w_rows = weight->get_row_ptrs(); auto diag_val = diag->get_data(); auto agg_val = agg->get_data(); auto c_vals = coarse->get_values(); @@ -158,6 +168,52 @@ class AmgxPgm : public ::testing::Test { cols[13] = 2; cols[14] = 4; + /* weight matrix is stored: + * 5 3 3 0 0 + * 3 5 0 2.5 1.5 + * 3 0 5 0 1.5 + * 0 2.5 0 5 0 + * 0 1.5 1.5 0 5 + */ + w_vals[0] = 5; + w_vals[1] = 3; + w_vals[2] = 3; + w_vals[3] = 3; + w_vals[4] = 5; + w_vals[5] = 2.5; + w_vals[6] = 1.5; + w_vals[7] = 3; + w_vals[8] = 5; + w_vals[9] = 1.5; + w_vals[10] = 2.5; + w_vals[11] = 5; + w_vals[12] = 1.5; + w_vals[13] = 1.5; + w_vals[14] = 5; + + w_rows[0] = 0; + w_rows[1] = 3; + w_rows[2] = 7; + w_rows[3] = 10; + w_rows[4] = 12; + w_rows[5] = 15; + + w_cols[0] = 0; + w_cols[1] = 1; + w_cols[2] = 2; + w_cols[3] = 0; + w_cols[4] = 1; + w_cols[5] = 3; + w_cols[6] = 4; + w_cols[7] = 0; + w_cols[8] = 2; + w_cols[9] = 4; + w_cols[10] = 1; + w_cols[11] = 3; + w_cols[12] = 1; + w_cols[13] = 2; + w_cols[14] = 4; + diag_val[0] = 5; diag_val[1] = 5; diag_val[2] = 5; @@ -214,7 +270,8 @@ class AmgxPgm : public ::testing::Test { std::shared_ptr exec; std::shared_ptr mtx; std::shared_ptr coarse; - gko::Array mtx_diag; + std::shared_ptr weight; + gko::Array mtx_diag; gko::Array agg; std::shared_ptr coarse_b; std::shared_ptr fine_b; @@ -448,11 +505,11 @@ TYPED_TEST(AmgxPgm, CoarseFineProlongApplyadd) TYPED_TEST(AmgxPgm, ExtractDiag) { - using value_type = typename TestFixture::value_type; - gko::Array diag(this->exec, 5); + using rmc_value_type = typename TestFixture::rmc_value_type; + gko::Array diag(this->exec, 5); - gko::kernels::reference::amgx_pgm::extract_diag(this->exec, this->mtx.get(), - diag); + gko::kernels::reference::amgx_pgm::extract_diag(this->exec, + this->weight.get(), diag); GKO_ASSERT_ARRAY_EQ(diag, this->mtx_diag); } @@ -470,7 +527,8 @@ TYPED_TEST(AmgxPgm, FindStrongestNeighbor) } gko::kernels::reference::amgx_pgm::find_strongest_neighbor( - this->exec, this->mtx.get(), this->mtx_diag, agg, strongest_neighbor); + this->exec, this->weight.get(), this->mtx_diag, agg, + strongest_neighbor); ASSERT_EQ(snb_vals[0], 2); ASSERT_EQ(snb_vals[1], 0); @@ -494,7 +552,7 @@ TYPED_TEST(AmgxPgm, AssignToExistAgg) agg_vals[4] = -1; gko::kernels::reference::amgx_pgm::assign_to_exist_agg( - this->exec, this->mtx.get(), this->mtx_diag, agg, intermediate_agg); + this->exec, this->weight.get(), this->mtx_diag, agg, intermediate_agg); ASSERT_EQ(agg_vals[0], 0); ASSERT_EQ(agg_vals[1], 1);