diff --git a/common/cuda_hip/stop/residual_norm_kernels.cpp b/common/cuda_hip/stop/residual_norm_kernels.cpp index 9d6db5211e8..584783dea78 100644 --- a/common/cuda_hip/stop/residual_norm_kernels.cpp +++ b/common/cuda_hip/stop/residual_norm_kernels.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -41,12 +41,12 @@ __global__ __launch_bounds__(default_block_size) void residual_norm_kernel( if (tidx < num_cols) { if (tau[tidx] <= rel_residual_goal * orig_tau[tidx]) { stop_status[tidx].converge(stoppingId, setFinalized); - device_storage[1] = true; + device_storage[0] = true; } // because only false is written to all_converged, write conflicts // should not cause any problem else if (!stop_status[tidx].has_stopped()) { - device_storage[0] = false; + device_storage[1] = false; } } } @@ -55,8 +55,10 @@ __global__ __launch_bounds__(default_block_size) void residual_norm_kernel( __global__ __launch_bounds__(1) void init_kernel( bool* __restrict__ device_storage) { - device_storage[0] = true; - device_storage[1] = false; + // one_changed + device_storage[0] = false; + // all_converged + device_storage[1] = true; } @@ -67,7 +69,7 @@ void residual_norm(std::shared_ptr exec, ValueType rel_residual_goal, uint8 stoppingId, bool setFinalized, array* stop_status, array* device_storage, bool* all_converged, - bool* one_changed) + bool* indicators) { static_assert(is_complex_s::value == false, "ValueType must not be complex in this function!"); @@ -86,9 +88,10 @@ void residual_norm(std::shared_ptr exec, as_device_type(device_storage->get_data())); } - /* Represents all_converged, one_changed */ - *all_converged = get_element(*device_storage, 0); - *one_changed = get_element(*device_storage, 1); + /* Represents all_converged(1), one_changed(0) */ + exec->get_master()->copy_from(exec, 2, device_storage->get_const_data(), + indicators); + *all_converged = indicators[1]; } GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_TYPE( @@ -122,25 +125,17 @@ __launch_bounds__(default_block_size) void implicit_residual_norm_kernel( if (tidx < num_cols) { if (sqrt(abs(tau[tidx])) <= rel_residual_goal * orig_tau[tidx]) { stop_status[tidx].converge(stoppingId, setFinalized); - device_storage[1] = true; + device_storage[0] = true; } // because only false is written to all_converged, write conflicts // should not cause any problem else if (!stop_status[tidx].has_stopped()) { - device_storage[0] = false; + device_storage[1] = false; } } } -__global__ __launch_bounds__(1) void init_kernel( - bool* __restrict__ device_storage) -{ - device_storage[0] = true; - device_storage[1] = false; -} - - template void implicit_residual_norm( std::shared_ptr exec, @@ -148,9 +143,9 @@ void implicit_residual_norm( const matrix::Dense>* orig_tau, remove_complex rel_residual_goal, uint8 stoppingId, bool setFinalized, array* stop_status, - array* device_storage, bool* all_converged, bool* one_changed) + array* device_storage, bool* all_converged, bool* indicators) { - init_kernel<<<1, 1, 0, exec->get_stream()>>>( + residual_norm::init_kernel<<<1, 1, 0, exec->get_stream()>>>( as_device_type(device_storage->get_data())); const auto block_size = default_block_size; @@ -166,9 +161,10 @@ void implicit_residual_norm( as_device_type(device_storage->get_data())); } - /* Represents all_converged, one_changed */ - *all_converged = get_element(*device_storage, 0); - *one_changed = get_element(*device_storage, 1); + /* Represents all_converged(1), one_changed(0) */ + exec->get_master()->copy_from(exec, 2, device_storage->get_const_data(), + indicators); + *all_converged = indicators[1]; } GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_IMPLICIT_RESIDUAL_NORM_KERNEL); diff --git a/core/solver/bicg.cpp b/core/solver/bicg.cpp index 8719d556eaf..65e7a144eb8 100644 --- a/core/solver/bicg.cpp +++ b/core/solver/bicg.cpp @@ -134,7 +134,6 @@ void Bicg::apply_dense_impl(const matrix::Dense* dense_b, GKO_SOLVER_ONE_MINUS_ONE(); - bool one_changed{}; GKO_SOLVER_STOP_REDUCTION_ARRAYS(); // rho = 0.0 @@ -195,13 +194,13 @@ void Bicg::apply_dense_impl(const matrix::Dense* dense_b, z->compute_conj_dot(r2, rho, reduction_tmp); ++iter; - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(r) - .implicit_sq_residual_norm(rho) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(r) + .implicit_sq_residual_norm(rho) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, iter, r, nullptr, rho, &stop_status, all_stopped); @@ -251,7 +250,7 @@ void Bicg::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 2; + return 3; } diff --git a/core/solver/bicgstab.cpp b/core/solver/bicgstab.cpp index 2b42982af8a..87e3dd872ac 100644 --- a/core/solver/bicgstab.cpp +++ b/core/solver/bicgstab.cpp @@ -120,7 +120,6 @@ void Bicgstab::apply_dense_impl(const VectorType* dense_b, GKO_SOLVER_ONE_MINUS_ONE(); - bool one_changed{}; GKO_SOLVER_STOP_REDUCTION_ARRAYS(); // r = dense_b @@ -160,13 +159,13 @@ void Bicgstab::apply_dense_impl(const VectorType* dense_b, ++iter; rr->compute_conj_dot(r, rho, reduction_tmp); - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(r) - .implicit_sq_residual_norm(rho) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(r) + .implicit_sq_residual_norm(rho) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, iter, r, nullptr, rho, &stop_status, all_stopped); @@ -193,14 +192,14 @@ void Bicgstab::apply_dense_impl(const VectorType* dense_b, gko::detail::get_local(r), gko::detail::get_local(s), gko::detail::get_local(v), rho, alpha, beta, &stop_status)); - all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(s) - .implicit_sq_residual_norm(rho) - // .solution(dense_x) // outdated at this point - .check(RelativeStoppingId, false, &stop_status, &one_changed); - if (one_changed) { + all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(s) + .implicit_sq_residual_norm(rho) + // .solution(dense_x) // outdated at this point + .check(RelativeStoppingId, false, &stop_status, + stop_indicators.get_data()); + if (stop_indicators.get_const_data()[0]) { exec->run(bicgstab::make_finalize(gko::detail::get_local(dense_x), gko::detail::get_local(y), alpha, &stop_status)); @@ -254,7 +253,7 @@ void Bicgstab::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 2; + return 3; } diff --git a/core/solver/cb_gmres.cpp b/core/solver/cb_gmres.cpp index 1496e37f77d..e253fb34d37 100644 --- a/core/solver/cb_gmres.cpp +++ b/core/solver/cb_gmres.cpp @@ -268,7 +268,8 @@ void CbGmres::apply_dense_impl( array final_iter_nums(this->get_executor(), num_rhs); auto y = Vector::create(exec, dim<2>{krylov_dim, num_rhs}); - bool one_changed{}; + array stop_indicators(this->get_executor()->get_master(), 2); + stop_indicators.get_data()[0] = false; array reduction_tmp{this->get_executor()}; array stop_status(this->get_executor(), num_rhs); // reorth_status and num_reorth are both helper variables for GPU @@ -341,17 +342,18 @@ void CbGmres::apply_dense_impl( residual_norm.get(), nullptr, &stop_status, false); ++forced_iterations; } else { - bool all_changed = stop_criterion->update() - .num_iterations(total_iter) - .residual(residual) - .residual_norm(residual_norm) - .solution(dense_x) - .check(RelativeStoppingId, true, - &stop_status, &one_changed); + bool all_changed = + stop_criterion->update() + .num_iterations(total_iter) + .residual(residual) + .residual_norm(residual_norm) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, total_iter, residual.get(), residual_norm.get(), nullptr, &stop_status, all_changed); - if (one_changed || all_changed) { + if (stop_indicators.get_const_data()[0] || all_changed) { host_stop_status = stop_status; bool host_array_changed{false}; for (size_type i = 0; i < host_stop_status.get_size(); diff --git a/core/solver/cg.cpp b/core/solver/cg.cpp index 04fd3c4de22..dc19617065b 100644 --- a/core/solver/cg.cpp +++ b/core/solver/cg.cpp @@ -112,7 +112,6 @@ void Cg::apply_dense_impl(const VectorType* dense_b, GKO_SOLVER_ONE_MINUS_ONE(); - bool one_changed{}; GKO_SOLVER_STOP_REDUCTION_ARRAYS(); // r = dense_b @@ -146,13 +145,13 @@ void Cg::apply_dense_impl(const VectorType* dense_b, r->compute_conj_dot(z, rho, reduction_tmp); ++iter; - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(r) - .implicit_sq_residual_norm(rho) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(r) + .implicit_sq_residual_norm(rho) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, iter, r, nullptr, rho, &stop_status, all_stopped); @@ -202,7 +201,7 @@ void Cg::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 2; + return 3; } diff --git a/core/solver/cgs.cpp b/core/solver/cgs.cpp index 983140c18e3..32e004ebdae 100644 --- a/core/solver/cgs.cpp +++ b/core/solver/cgs.cpp @@ -118,7 +118,6 @@ void Cgs::apply_dense_impl(const VectorType* dense_b, GKO_SOLVER_ONE_MINUS_ONE(); - bool one_changed{}; GKO_SOLVER_STOP_REDUCTION_ARRAYS(); // r = dense_b @@ -155,13 +154,13 @@ void Cgs::apply_dense_impl(const VectorType* dense_b, r->compute_conj_dot(r_tld, rho, reduction_tmp); ++iter; - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(r) - .implicit_sq_residual_norm(rho) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(r) + .implicit_sq_residual_norm(rho) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, iter, r, nullptr, rho, &stop_status, all_stopped); @@ -222,7 +221,7 @@ void Cgs::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 2; + return 3; } diff --git a/core/solver/chebyshev.cpp b/core/solver/chebyshev.cpp index 7f409484e6f..c5d9d042342 100644 --- a/core/solver/chebyshev.cpp +++ b/core/solver/chebyshev.cpp @@ -225,6 +225,11 @@ void Chebyshev::apply_dense_impl(const VectorType* dense_b, auto& stop_status = this->template create_workspace_array( ws::stop, dense_b->get_size()[1]); exec->run(ir::make_initialize(&stop_status)); + exec->run(ir::make_initialize(&stop_status)); + auto& stop_indicators = + this->template create_workspace_array(ws::indicators, 2); + stop_indicators.set_executor(this->get_executor()->get_master()); + stop_indicators.get_data()[0] = false; if (guess != initial_guess_mode::zero) { residual->copy_from(dense_b); this->get_system_matrix()->apply(neg_one_op, dense_x, one_op, residual); @@ -251,7 +256,7 @@ void Chebyshev::apply_dense_impl(const VectorType* dense_b, }; bool all_stopped = update_residual( this, iter, dense_b, dense_x, residual, residual_ptr, - stop_criterion, stop_status, log_func); + stop_criterion, stop_status, &stop_indicators, log_func); if (all_stopped) { break; } @@ -321,7 +326,7 @@ void Chebyshev::apply_with_initial_guess_impl( template int workspace_traits>::num_arrays(const Solver&) { - return 1; + return 2; } diff --git a/core/solver/fcg.cpp b/core/solver/fcg.cpp index 7f04b8ba2fd..d29a52941a8 100644 --- a/core/solver/fcg.cpp +++ b/core/solver/fcg.cpp @@ -112,7 +112,6 @@ void Fcg::apply_dense_impl(const VectorType* dense_b, GKO_SOLVER_ONE_MINUS_ONE(); - bool one_changed{}; GKO_SOLVER_STOP_REDUCTION_ARRAYS(); // r = dense_b @@ -148,13 +147,13 @@ void Fcg::apply_dense_impl(const VectorType* dense_b, t->compute_conj_dot(z, rho_t, reduction_tmp); ++iter; - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(r) - .implicit_sq_residual_norm(rho) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(r) + .implicit_sq_residual_norm(rho) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, iter, r, nullptr, rho, &stop_status, all_stopped); @@ -204,7 +203,7 @@ void Fcg::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 2; + return 3; } diff --git a/core/solver/gcr.cpp b/core/solver/gcr.cpp index 629cbe7ddb2..40e215e8c94 100644 --- a/core/solver/gcr.cpp +++ b/core/solver/gcr.cpp @@ -136,8 +136,6 @@ void Gcr::apply_dense_impl(const VectorType* dense_b, auto& final_iter_nums = this->template create_workspace_array( ws::final_iter_nums, num_rhs); - // indicates if the status of a vector has changed - bool one_changed{}; GKO_SOLVER_ONE_MINUS_ONE(); GKO_SOLVER_STOP_REDUCTION_ARRAYS(); @@ -196,13 +194,13 @@ void Gcr::apply_dense_impl(const VectorType* dense_b, residual->compute_norm2(residual_norm, reduction_tmp); // Should the iteration stop? - auto all_stopped = - stop_criterion->update() - .num_iterations(total_iter) - .residual(residual) - .residual_norm(residual_norm) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + auto all_stopped = stop_criterion->update() + .num_iterations(total_iter) + .residual(residual) + .residual_norm(residual_norm) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); // Log current iteration this->template log( @@ -320,7 +318,7 @@ void Gcr::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 3; + return 4; } diff --git a/core/solver/gmres.cpp b/core/solver/gmres.cpp index b5bceb462de..d246c163b8f 100644 --- a/core/solver/gmres.cpp +++ b/core/solver/gmres.cpp @@ -395,7 +395,6 @@ void Gmres::apply_dense_impl(const VectorType* dense_b, GKO_SOLVER_ONE_MINUS_ONE(); - bool one_changed{}; GKO_SOLVER_STOP_REDUCTION_ARRAYS(); auto& final_iter_nums = this->template create_workspace_array( ws::final_iter_nums, num_rhs); @@ -448,13 +447,13 @@ void Gmres::apply_dense_impl(const VectorType* dense_b, */ while (true) { ++total_iter; - bool all_stopped = - stop_criterion->update() - .num_iterations(total_iter) - .residual(residual) - .residual_norm(residual_norm) - .solution(dense_x) - .check(RelativeStoppingId, false, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(total_iter) + .residual(residual) + .residual_norm(residual_norm) + .solution(dense_x) + .check(RelativeStoppingId, false, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, total_iter, residual, residual_norm, nullptr, &stop_status, all_stopped); @@ -650,7 +649,7 @@ void Gmres::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 3; + return 4; } diff --git a/core/solver/idr.cpp b/core/solver/idr.cpp index e1dfa4ef349..f9b443468fc 100644 --- a/core/solver/idr.cpp +++ b/core/solver/idr.cpp @@ -157,7 +157,6 @@ void Idr::iterate(const VectorType* dense_b, ws::subspace_minus_one, 1); subspace_neg_one_op->fill(-one()); - bool one_changed{}; GKO_SOLVER_STOP_REDUCTION_ARRAYS(); // Initialization @@ -212,13 +211,13 @@ void Idr::iterate(const VectorType* dense_b, while (true) { ++total_iter; - bool all_stopped = - stop_criterion->update() - .num_iterations(total_iter) - .residual(residual) - .residual_norm(residual_norm) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(total_iter) + .residual(residual) + .residual_norm(residual_norm) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, total_iter, residual, nullptr, nullptr, &stop_status, all_stopped); @@ -348,7 +347,7 @@ void Idr::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 2; + return 3; } diff --git a/core/solver/ir.cpp b/core/solver/ir.cpp index 62f618f1d7a..f6bf4038d24 100644 --- a/core/solver/ir.cpp +++ b/core/solver/ir.cpp @@ -208,6 +208,10 @@ void Ir::apply_dense_impl(const VectorType* dense_b, auto& stop_status = this->template create_workspace_array( ws::stop, dense_b->get_size()[1]); exec->run(ir::make_initialize(&stop_status)); + auto& stop_indicators = + this->template create_workspace_array(ws::indicators, 2); + stop_indicators.set_executor(this->get_executor()->get_master()); + stop_indicators.get_data()[0] = false; if (guess != initial_guess_mode::zero) { residual->copy_from(dense_b); this->get_system_matrix()->apply(neg_one_op, dense_x, one_op, residual); @@ -235,7 +239,7 @@ void Ir::apply_dense_impl(const VectorType* dense_b, }; bool all_stopped = update_residual( this, iter, dense_b, dense_x, residual, residual_ptr, - stop_criterion, stop_status, log_func); + stop_criterion, stop_status, &stop_indicators, log_func); if (all_stopped) { break; } @@ -289,7 +293,7 @@ void Ir::apply_with_initial_guess_impl( template int workspace_traits>::num_arrays(const Solver&) { - return 1; + return 2; } diff --git a/core/solver/minres.cpp b/core/solver/minres.cpp index 80ed81e1ee4..baf1361f78a 100644 --- a/core/solver/minres.cpp +++ b/core/solver/minres.cpp @@ -153,7 +153,6 @@ void Minres::apply_dense_impl(const VectorType* dense_b, GKO_SOLVER_ONE_MINUS_ONE(); - bool one_changed{}; GKO_SOLVER_STOP_REDUCTION_ARRAYS(); // r = dense_b @@ -198,13 +197,13 @@ void Minres::apply_dense_impl(const VectorType* dense_b, */ while (true) { ++iter; - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(nullptr) - .implicit_sq_residual_norm(tau) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(nullptr) + .implicit_sq_residual_norm(tau) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, iter, r, nullptr, tau, &stop_status, all_stopped); @@ -331,7 +330,7 @@ int workspace_traits>::num_vectors(const Solver&) template int workspace_traits>::num_arrays(const Solver&) { - return 2; + return 3; } diff --git a/core/solver/multigrid.cpp b/core/solver/multigrid.cpp index 867c6e94c4a..2dc7edc000b 100644 --- a/core/solver/multigrid.cpp +++ b/core/solver/multigrid.cpp @@ -966,7 +966,9 @@ void Multigrid::apply_dense_impl(const VectorType* b, VectorType* x, auto& stop_status = this->template create_workspace_array( ws::stop, b->get_size()[1]); - bool one_changed{}; + auto& stop_indicators = + this->template create_workspace_array(ws::indicators, 2); + stop_indicators.set_executor(this->get_executor()->get_master()); exec->run(multigrid::make_initialize(&stop_status)); auto stop_criterion = this->get_stop_criterion_factory()->generate( this->get_system_matrix(), @@ -985,7 +987,7 @@ void Multigrid::apply_dense_impl(const VectorType* b, VectorType* x, // residual check. .solution(x) .check(RelativeStoppingId, true, &stop_status, - &one_changed); + stop_indicators.get_data()); this->template log( this, b, x, iter, nullptr, nullptr, nullptr, &stop_status, all_stopped); @@ -1109,7 +1111,7 @@ Multigrid::Multigrid(std::shared_ptr exec) {} -int workspace_traits::num_arrays(const Solver&) { return 1; } +int workspace_traits::num_arrays(const Solver&) { return 2; } int workspace_traits::num_vectors(const Solver&) { return 0; } diff --git a/core/solver/pipe_cg.cpp b/core/solver/pipe_cg.cpp index 11d74ed9db6..eeb9249fe30 100644 --- a/core/solver/pipe_cg.cpp +++ b/core/solver/pipe_cg.cpp @@ -119,8 +119,6 @@ void PipeCg::apply_dense_impl(const VectorType* dense_b, GKO_SOLVER_ONE_MINUS_ONE(); - bool one_changed{}; - GKO_SOLVER_STOP_REDUCTION_ARRAYS(); // r = b @@ -149,13 +147,13 @@ void PipeCg::apply_dense_impl(const VectorType* dense_b, this->get_system_matrix(), std::shared_ptr(dense_b, [](const LinOp*) {}), dense_x, r); int iter = 0; - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(r) - .implicit_sq_residual_norm(rho) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(r) + .implicit_sq_residual_norm(rho) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, iter, r, nullptr, rho, &stop_status, all_stopped); @@ -202,13 +200,13 @@ void PipeCg::apply_dense_impl(const VectorType* dense_b, w->compute_conj_dot(z, delta, reduction_tmp); // check ++iter; - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(r) - .implicit_sq_residual_norm(rho) - .solution(dense_x) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(r) + .implicit_sq_residual_norm(rho) + .solution(dense_x) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); this->template log( this, dense_b, dense_x, iter, r, nullptr, rho, &stop_status, all_stopped); @@ -252,7 +250,7 @@ void PipeCg::apply_impl(const LinOp* alpha, const LinOp* b, template int workspace_traits>::num_arrays(const Solver&) { - return 2; + return 3; } diff --git a/core/solver/solver_boilerplate.hpp b/core/solver/solver_boilerplate.hpp index db380833ca5..712196cf059 100644 --- a/core/solver/solver_boilerplate.hpp +++ b/core/solver/solver_boilerplate.hpp @@ -25,9 +25,13 @@ auto neg_one_op = this->template create_workspace_fixed_scalar( \ GKO_SOLVER_TRAITS::minus_one, 1, -one()) -#define GKO_SOLVER_STOP_REDUCTION_ARRAYS() \ - auto& stop_status = \ - this->template create_workspace_array( \ - GKO_SOLVER_TRAITS::stop, dense_b->get_size()[1]); \ - auto& reduction_tmp = \ - this->template create_workspace_array(GKO_SOLVER_TRAITS::tmp) +#define GKO_SOLVER_STOP_REDUCTION_ARRAYS() \ + auto& stop_status = \ + this->template create_workspace_array( \ + GKO_SOLVER_TRAITS::stop, dense_b->get_size()[1]); \ + auto& reduction_tmp = \ + this->template create_workspace_array(GKO_SOLVER_TRAITS::tmp); \ + auto& stop_indicators = this->template create_workspace_array( \ + GKO_SOLVER_TRAITS::indicators, 2); \ + stop_indicators.set_executor(this->get_executor()->get_master()); \ + stop_indicators.get_data()[0] = false diff --git a/core/solver/update_residual.hpp b/core/solver/update_residual.hpp index 688dcdb4b18..40d1a55f865 100644 --- a/core/solver/update_residual.hpp +++ b/core/solver/update_residual.hpp @@ -21,7 +21,8 @@ bool update_residual(SolverType* solver, int iter, const VectorType* dense_b, VectorType* dense_x, VectorType* residual, const VectorType*& residual_ptr, std::unique_ptr& stop_criterion, - array& stop_status, LogFunc log) + array& stop_status, + array* stop_indicators, LogFunc log) { using ws = workspace_traits>; constexpr uint8 relative_stopping_id{1}; @@ -30,27 +31,26 @@ bool update_residual(SolverType* solver, int iter, const VectorType* dense_b, auto one_op = solver->get_workspace_op(ws::one); auto neg_one_op = solver->get_workspace_op(ws::minus_one); - bool one_changed{}; if (iter == 0) { // In iter 0, the iteration and residual are updated. - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(residual_ptr) - .solution(dense_x) - .check(relative_stopping_id, true, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(residual_ptr) + .solution(dense_x) + .check(relative_stopping_id, true, &stop_status, + stop_indicators->get_data()); log(solver, dense_b, dense_x, iter, residual_ptr, stop_status, all_stopped); return all_stopped; } else { // In the other iterations, the residual can be updated separately. - bool all_stopped = - stop_criterion->update() - .num_iterations(iter) - .solution(dense_x) - // we have the residual check later - .ignore_residual_check(true) - .check(relative_stopping_id, false, &stop_status, &one_changed); + bool all_stopped = stop_criterion->update() + .num_iterations(iter) + .solution(dense_x) + // we have the residual check later + .ignore_residual_check(true) + .check(relative_stopping_id, false, &stop_status, + stop_indicators->get_data()); if (all_stopped) { log(solver, dense_b, dense_x, iter, nullptr, stop_status, all_stopped); @@ -61,12 +61,12 @@ bool update_residual(SolverType* solver, int iter, const VectorType* dense_b, residual->copy_from(dense_b); solver->get_system_matrix()->apply(neg_one_op, dense_x, one_op, residual); - all_stopped = - stop_criterion->update() - .num_iterations(iter) - .residual(residual_ptr) - .solution(dense_x) - .check(relative_stopping_id, true, &stop_status, &one_changed); + all_stopped = stop_criterion->update() + .num_iterations(iter) + .residual(residual_ptr) + .solution(dense_x) + .check(relative_stopping_id, true, &stop_status, + stop_indicators->get_data()); log(solver, dense_b, dense_x, iter, residual_ptr, stop_status, all_stopped); return all_stopped; diff --git a/core/stop/combined.cpp b/core/stop/combined.cpp index d29d65f73bc..66db87b50f0 100644 --- a/core/stop/combined.cpp +++ b/core/stop/combined.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -32,22 +32,26 @@ Combined::Combined(const Combined::Factory* factory, const CriterionArgs& args) bool Combined::check_impl(uint8 stoppingId, bool setFinalized, - array* stop_status, - bool* one_changed, const Updater& updater) + array* stop_status, bool* indicators, + const Updater& updater) { bool one_converged = false; gko::uint8 ids{1}; - *one_changed = false; + auto one_changed = false; for (auto& c : criteria_) { - bool local_one_changed = false; - one_converged |= c->check(ids, setFinalized, stop_status, - &local_one_changed, updater); - *one_changed |= local_one_changed; + // use indicators in each critirion such that we can reduce the + // #copy_call + indicators[0] = false; + one_converged |= + c->check(ids, setFinalized, stop_status, indicators, updater); + one_changed |= indicators[0]; if (one_converged) { break; } ids++; } + // set the result back to return + indicators[0] = one_changed; return one_converged; } diff --git a/core/stop/iteration.cpp b/core/stop/iteration.cpp index 2f712865eda..73365300d12 100644 --- a/core/stop/iteration.cpp +++ b/core/stop/iteration.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -11,12 +11,12 @@ namespace stop { bool Iteration::check_impl(uint8 stoppingId, bool setFinalized, array* stop_status, - bool* one_changed, const Updater& updater) + bool* indicators, const Updater& updater) { bool result = updater.num_iterations_ >= parameters_.max_iters; if (result) { this->set_all_statuses(stoppingId, setFinalized, stop_status); - *one_changed = true; + *indicators = true; } return result; } diff --git a/core/stop/residual_norm.cpp b/core/stop/residual_norm.cpp index c962784033a..844d00aaeb0 100644 --- a/core/stop/residual_norm.cpp +++ b/core/stop/residual_norm.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -162,7 +162,7 @@ ResidualNormBase::ResidualNormBase( template bool ResidualNormBase::check_impl( uint8 stopping_id, bool set_finalized, array* stop_status, - bool* one_changed, const Criterion::Updater& updater) + bool* indicators, const Criterion::Updater& updater) { const NormVector* dense_tau; if (updater.residual_norm_ != nullptr) { @@ -197,7 +197,7 @@ bool ResidualNormBase::check_impl( this->get_executor()->run(residual_norm::make_residual_norm( dense_tau, starting_tau_.get(), reduction_factor_, stopping_id, set_finalized, stop_status, &device_storage_, &all_converged, - one_changed)); + indicators)); return all_converged; } @@ -206,7 +206,7 @@ bool ResidualNormBase::check_impl( template bool ImplicitResidualNorm::check_impl( uint8 stopping_id, bool set_finalized, array* stop_status, - bool* one_changed, const Criterion::Updater& updater) + bool* indicators, const Criterion::Updater& updater) { const Vector* dense_tau; if (updater.implicit_sq_residual_norm_ != nullptr) { @@ -220,7 +220,7 @@ bool ImplicitResidualNorm::check_impl( implicit_residual_norm::make_implicit_residual_norm( dense_tau, this->starting_tau_.get(), this->reduction_factor_, stopping_id, set_finalized, stop_status, &this->device_storage_, - &all_converged, one_changed)); + &all_converged, indicators)); return all_converged; } diff --git a/core/stop/residual_norm_kernels.hpp b/core/stop/residual_norm_kernels.hpp index 7625dadefeb..cc9d5209dac 100644 --- a/core/stop/residual_norm_kernels.hpp +++ b/core/stop/residual_norm_kernels.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -24,7 +24,7 @@ namespace residual_norm { const matrix::Dense<_type>* tau, const matrix::Dense<_type>* orig_tau, \ _type rel_residual_goal, uint8 stoppingId, bool setFinalized, \ array* stop_status, array* device_storage, \ - bool* all_converged, bool* one_changed) + bool* all_converged, bool* indicators) #define GKO_DECLARE_ALL_AS_TEMPLATES \ @@ -45,7 +45,7 @@ namespace implicit_residual_norm { const matrix::Dense>* orig_tau, \ remove_complex<_type> rel_residual_goal, uint8 stoppingId, \ bool setFinalized, array* stop_status, \ - array* device_storage, bool* all_converged, bool* one_changed) + array* device_storage, bool* all_converged, bool* indicators) #define GKO_DECLARE_ALL_AS_TEMPLATES2 \ diff --git a/core/stop/time.cpp b/core/stop/time.cpp index 0481b9c91d3..70ad0ca88ff 100644 --- a/core/stop/time.cpp +++ b/core/stop/time.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -10,13 +10,13 @@ namespace stop { bool Time::check_impl(uint8 stoppingId, bool setFinalized, - array* stop_status, bool* one_changed, + array* stop_status, bool* indicators, const Updater& updater) { bool result = clock::now() - start_ >= time_limit_; if (result) { this->set_all_statuses(stoppingId, setFinalized, stop_status); - *one_changed = true; + *indicators = true; } return result; } diff --git a/core/test/stop/criterion.cpp b/core/test/stop/criterion.cpp index ce555d01969..8014b6daf77 100644 --- a/core/test/stop/criterion.cpp +++ b/core/test/stop/criterion.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -56,7 +56,7 @@ class DummyCriterion protected: bool check_impl(gko::uint8 stopping_id, bool set_finalized, gko::array* stop_status, - bool* one_changed, const Updater& updater) override + bool* indicators, const Updater& updater) override { return true; } @@ -95,10 +95,10 @@ TEST_F(Criterion, DefaultUpdateStatus) TEST_F(Criterion, CanLogCheck) { auto before_logger = *logger; - bool one_changed = false; + gko::array stop_indicators(exec->get_master(), 2); - criterion->check(gko::uint8(0), false, &stopping_status, &one_changed, - criterion->update()); + criterion->check(gko::uint8(0), false, &stopping_status, + stop_indicators.get_data(), criterion->update()); ASSERT_EQ(logger->criterion_check_started, before_logger.criterion_check_started + 1); diff --git a/dpcpp/stop/residual_norm_kernels.dp.cpp b/dpcpp/stop/residual_norm_kernels.dp.cpp index 129c26e08f4..cace56a6c74 100644 --- a/dpcpp/stop/residual_norm_kernels.dp.cpp +++ b/dpcpp/stop/residual_norm_kernels.dp.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -35,15 +35,17 @@ void residual_norm(std::shared_ptr exec, ValueType rel_residual_goal, uint8 stoppingId, bool setFinalized, array* stop_status, array* device_storage, bool* all_converged, - bool* one_changed) + bool* indicators) { static_assert(is_complex_s::value == false, "ValueType must not be complex in this function!"); auto device_storage_val = device_storage->get_data(); exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for(sycl::range<1>{1}, [=](sycl::id<1>) { - device_storage_val[0] = true; - device_storage_val[1] = false; + // one_changed + device_storage_val[0] = false; + // all_converged + device_storage_val[1] = true; }); }); @@ -56,19 +58,20 @@ void residual_norm(std::shared_ptr exec, const auto tidx = idx_id[0]; if (tau_val[tidx] <= rel_residual_goal * orig_tau_val[tidx]) { stop_status_val[tidx].converge(stoppingId, setFinalized); - device_storage_val[1] = true; + device_storage_val[0] = true; } // because only false is written to all_converged, write // conflicts should not cause any problem else if (!stop_status_val[tidx].has_stopped()) { - device_storage_val[0] = false; + device_storage_val[1] = false; } }); }); - /* Represents all_converged, one_changed */ - *all_converged = get_element(*device_storage, 0); - *one_changed = get_element(*device_storage, 1); + /* Represents all_converged(1), one_changed(0) */ + exec->get_master()->copy_from(exec, 2, device_storage->get_const_data(), + indicators); + *all_converged = indicators[1]; } GKO_INSTANTIATE_FOR_EACH_NON_COMPLEX_VALUE_TYPE( @@ -93,13 +96,15 @@ void implicit_residual_norm( const matrix::Dense>* orig_tau, remove_complex rel_residual_goal, uint8 stoppingId, bool setFinalized, array* stop_status, - array* device_storage, bool* all_converged, bool* one_changed) + array* device_storage, bool* all_converged, bool* indicators) { auto device_storage_val = device_storage->get_data(); exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for(sycl::range<1>{1}, [=](sycl::id<1>) { - device_storage_val[0] = true; - device_storage_val[1] = false; + // one_changed + device_storage_val[0] = false; + // all_converged + device_storage_val[1] = true; }); }); @@ -113,19 +118,20 @@ void implicit_residual_norm( 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; + device_storage_val[0] = true; } // because only false is written to all_converged, write // conflicts should not cause any problem else if (!stop_status_val[tidx].has_stopped()) { - device_storage_val[0] = false; + device_storage_val[1] = false; } }); }); - /* Represents all_converged, one_changed */ - *all_converged = get_element(*device_storage, 0); - *one_changed = get_element(*device_storage, 1); + /* Represents all_converged(1), one_changed(0) */ + exec->get_master()->copy_from(exec, 2, device_storage->get_const_data(), + indicators); + *all_converged = indicators[1]; } GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_IMPLICIT_RESIDUAL_NORM_KERNEL); diff --git a/examples/custom-stopping-criterion/custom-stopping-criterion.cpp b/examples/custom-stopping-criterion/custom-stopping-criterion.cpp index 030e11323af..4155b62ac5b 100644 --- a/examples/custom-stopping-criterion/custom-stopping-criterion.cpp +++ b/examples/custom-stopping-criterion/custom-stopping-criterion.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -37,12 +37,12 @@ class ByInteraction protected: bool check_impl(gko::uint8 stoppingId, bool setFinalized, gko::array* stop_status, - bool* one_changed, const Criterion::Updater&) override + bool* indicators, const Criterion::Updater&) override { bool result = *(parameters_.stop_iteration_process); if (result) { this->set_all_statuses(stoppingId, setFinalized, stop_status); - *one_changed = true; + *indicators = true; } return result; } diff --git a/include/ginkgo/core/solver/bicg.hpp b/include/ginkgo/core/solver/bicg.hpp index 2a43c1ca3f8..03d1f36159a 100644 --- a/include/ginkgo/core/solver/bicg.hpp +++ b/include/ginkgo/core/solver/bicg.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -170,6 +170,8 @@ struct workspace_traits> { constexpr static int stop = 0; // reduction tmp array constexpr static int tmp = 1; + // stopping indicator array + constexpr static int indicators = 2; }; diff --git a/include/ginkgo/core/solver/bicgstab.hpp b/include/ginkgo/core/solver/bicgstab.hpp index a57a6c27aa4..292c062b536 100644 --- a/include/ginkgo/core/solver/bicgstab.hpp +++ b/include/ginkgo/core/solver/bicgstab.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -174,6 +174,8 @@ struct workspace_traits> { constexpr static int stop = 0; // reduction tmp array constexpr static int tmp = 1; + // stopping indicator array + constexpr static int indicators = 2; }; diff --git a/include/ginkgo/core/solver/cg.hpp b/include/ginkgo/core/solver/cg.hpp index 984d5d1f104..32976c37488 100644 --- a/include/ginkgo/core/solver/cg.hpp +++ b/include/ginkgo/core/solver/cg.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -156,6 +156,8 @@ struct workspace_traits> { constexpr static int stop = 0; // reduction tmp array constexpr static int tmp = 1; + // stopping indicator array + constexpr static int indicators = 2; }; diff --git a/include/ginkgo/core/solver/cgs.hpp b/include/ginkgo/core/solver/cgs.hpp index bde23d76910..9fc4adc6286 100644 --- a/include/ginkgo/core/solver/cgs.hpp +++ b/include/ginkgo/core/solver/cgs.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -165,6 +165,8 @@ struct workspace_traits> { constexpr static int stop = 0; // reduction tmp array constexpr static int tmp = 1; + // stopping indicator array + constexpr static int indicators = 2; }; diff --git a/include/ginkgo/core/solver/chebyshev.hpp b/include/ginkgo/core/solver/chebyshev.hpp index cc04282e624..4ef4b54aade 100644 --- a/include/ginkgo/core/solver/chebyshev.hpp +++ b/include/ginkgo/core/solver/chebyshev.hpp @@ -224,6 +224,8 @@ struct workspace_traits> { // stopping status array constexpr static int stop = 0; + // stopping indicator array + constexpr static int indicators = 1; }; diff --git a/include/ginkgo/core/solver/fcg.hpp b/include/ginkgo/core/solver/fcg.hpp index dfaf252b557..dc3b5dfeccc 100644 --- a/include/ginkgo/core/solver/fcg.hpp +++ b/include/ginkgo/core/solver/fcg.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -165,6 +165,8 @@ struct workspace_traits> { constexpr static int stop = 0; // reduction tmp array constexpr static int tmp = 1; + // stopping indicator array + constexpr static int indicators = 2; }; diff --git a/include/ginkgo/core/solver/gcr.hpp b/include/ginkgo/core/solver/gcr.hpp index 62ce9c9c93c..cecd5bcad26 100644 --- a/include/ginkgo/core/solver/gcr.hpp +++ b/include/ginkgo/core/solver/gcr.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -183,6 +183,8 @@ struct workspace_traits> { constexpr static int tmp = 1; // final iteration number array constexpr static int final_iter_nums = 2; + // stopping indicator array + constexpr static int indicators = 3; }; diff --git a/include/ginkgo/core/solver/gmres.hpp b/include/ginkgo/core/solver/gmres.hpp index 3ba3acf94bb..1f05b7aa12c 100644 --- a/include/ginkgo/core/solver/gmres.hpp +++ b/include/ginkgo/core/solver/gmres.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -225,6 +225,8 @@ struct workspace_traits> { constexpr static int tmp = 1; // final iteration number array constexpr static int final_iter_nums = 2; + // stopping indicator array + constexpr static int indicators = 3; }; diff --git a/include/ginkgo/core/solver/idr.hpp b/include/ginkgo/core/solver/idr.hpp index 9f167d9b2eb..c25d0dbbce4 100644 --- a/include/ginkgo/core/solver/idr.hpp +++ b/include/ginkgo/core/solver/idr.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -292,6 +292,8 @@ struct workspace_traits> { constexpr static int stop = 0; // reduction tmp array constexpr static int tmp = 1; + // stopping indicator array + constexpr static int indicators = 2; }; diff --git a/include/ginkgo/core/solver/ir.hpp b/include/ginkgo/core/solver/ir.hpp index 91949261a79..4c896a5fc5e 100644 --- a/include/ginkgo/core/solver/ir.hpp +++ b/include/ginkgo/core/solver/ir.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -285,6 +285,8 @@ struct workspace_traits> { // stopping status array constexpr static int stop = 0; + // stopping indicator array + constexpr static int indicators = 1; }; diff --git a/include/ginkgo/core/solver/minres.hpp b/include/ginkgo/core/solver/minres.hpp index 808787e9583..a4d439bc14c 100644 --- a/include/ginkgo/core/solver/minres.hpp +++ b/include/ginkgo/core/solver/minres.hpp @@ -180,6 +180,8 @@ struct workspace_traits> { constexpr static int stop = 0; // reduction tmp array constexpr static int tmp = 1; + // stopping indicator array + constexpr static int indicators = 2; }; diff --git a/include/ginkgo/core/solver/multigrid.hpp b/include/ginkgo/core/solver/multigrid.hpp index 2d0278b538e..31187df02ba 100644 --- a/include/ginkgo/core/solver/multigrid.hpp +++ b/include/ginkgo/core/solver/multigrid.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -503,6 +503,8 @@ struct workspace_traits { // stopping status array constexpr static int stop = 0; + // stopping indicator array + constexpr static int indicators = 1; }; diff --git a/include/ginkgo/core/solver/pipe_cg.hpp b/include/ginkgo/core/solver/pipe_cg.hpp index 53597487e86..29f6f3d84a8 100644 --- a/include/ginkgo/core/solver/pipe_cg.hpp +++ b/include/ginkgo/core/solver/pipe_cg.hpp @@ -182,6 +182,8 @@ struct workspace_traits> { constexpr static int stop = 0; // reduction tmp array constexpr static int tmp = 1; + // stopping indicator array + constexpr static int indicators = 2; }; diff --git a/include/ginkgo/core/stop/combined.hpp b/include/ginkgo/core/stop/combined.hpp index 62451538431..4e371f32c32 100644 --- a/include/ginkgo/core/stop/combined.hpp +++ b/include/ginkgo/core/stop/combined.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -72,7 +72,7 @@ class Combined : public EnablePolymorphicObject { protected: bool check_impl(uint8 stoppingId, bool setFinalized, - array* stop_status, bool* one_changed, + array* stop_status, bool* indicators, const Updater&) override; explicit Combined(std::shared_ptr exec); diff --git a/include/ginkgo/core/stop/criterion.hpp b/include/ginkgo/core/stop/criterion.hpp index edf589d4ca5..b462521fd2d 100644 --- a/include/ginkgo/core/stop/criterion.hpp +++ b/include/ginkgo/core/stop/criterion.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -71,10 +71,10 @@ class Criterion : public EnableAbstractPolymorphicObject { * @copydoc Criterion::check(uint8, bool, array, bool) */ bool check(uint8 stopping_id, bool set_finalized, - array* stop_status, bool* one_changed) const + array* stop_status, bool* indicators) const { auto converged = parent_->check(stopping_id, set_finalized, - stop_status, one_changed, *this); + stop_status, indicators, *this); return converged; } @@ -129,26 +129,28 @@ class Criterion : public EnableAbstractPolymorphicObject { * @param set_finalized Controls if the current version should count as * finalized or not * @param stop_status status of the stopping criterion - * @param one_changed indicates if the status of a vector has changed + * @param indicators it should be a bool pointer to the host array + * containing at least two elements. The first entry will + * indicate if the status of a vector has changed. * @param updater the Updater object containing all the information * * @returns whether convergence was completely reached */ bool check(uint8 stopping_id, bool set_finalized, - array* stop_status, bool* one_changed, + array* stop_status, bool* indecators, const Updater& updater) { this->template log( this, updater.num_iterations_, updater.residual_, updater.residual_norm_, updater.solution_, stopping_id, set_finalized); - auto all_converged = this->check_impl( - stopping_id, set_finalized, stop_status, one_changed, updater); + auto all_converged = this->check_impl(stopping_id, set_finalized, + stop_status, indecators, updater); this->template log( this, updater.num_iterations_, updater.residual_, updater.residual_norm_, updater.implicit_sq_residual_norm_, updater.solution_, stopping_id, set_finalized, stop_status, - *one_changed, all_converged); + *indecators, all_converged); return all_converged; } @@ -164,14 +166,16 @@ class Criterion : public EnableAbstractPolymorphicObject { * @param set_finalized Controls if the current version should count as * finalized or not * @param stop_status status of the stopping criterion - * @param one_changed indicates if the status of a vector has changed + * @param indicators it should be a bool pointer to the host array + * containing at least two elements. The first entry will + * indicate if the status of a vector has changed. * @param updater the Updater object containing all the information * * @returns whether convergence was completely reached */ virtual bool check_impl(uint8 stopping_id, bool set_finalized, array* stop_status, - bool* one_changed, const Updater& updater) = 0; + bool* indicators, const Updater& updater) = 0; /** * This is a helper function which properly sets all elements of the diff --git a/include/ginkgo/core/stop/iteration.hpp b/include/ginkgo/core/stop/iteration.hpp index 49304b0f5df..0c33f4dfaf7 100644 --- a/include/ginkgo/core/stop/iteration.hpp +++ b/include/ginkgo/core/stop/iteration.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -37,7 +37,7 @@ class Iteration : public EnablePolymorphicObject { protected: bool check_impl(uint8 stoppingId, bool setFinalized, - array* stop_status, bool* one_changed, + array* stop_status, bool* indicators, const Updater& updater) override; explicit Iteration(std::shared_ptr exec) diff --git a/include/ginkgo/core/stop/residual_norm.hpp b/include/ginkgo/core/stop/residual_norm.hpp index c7f240950fa..15341c345a5 100644 --- a/include/ginkgo/core/stop/residual_norm.hpp +++ b/include/ginkgo/core/stop/residual_norm.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -59,7 +59,7 @@ class ResidualNormBase using NormVector = matrix::Dense; using Vector = matrix::Dense; bool check_impl(uint8 stoppingId, bool setFinalized, - array* stop_status, bool* one_changed, + array* stop_status, bool* indicators, const Criterion::Updater& updater) override; explicit ResidualNormBase(std::shared_ptr exec) @@ -196,7 +196,7 @@ class ImplicitResidualNorm : public ResidualNormBase { // check_impl needs to be overwritten again since we focus on the implicit // residual here bool check_impl(uint8 stoppingId, bool setFinalized, - array* stop_status, bool* one_changed, + array* stop_status, bool* indicators, const Criterion::Updater& updater) override; explicit ImplicitResidualNorm(std::shared_ptr exec) diff --git a/include/ginkgo/core/stop/time.hpp b/include/ginkgo/core/stop/time.hpp index a41b9c49499..55fb8dd0c72 100644 --- a/include/ginkgo/core/stop/time.hpp +++ b/include/ginkgo/core/stop/time.hpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -39,7 +39,7 @@ class Time : public EnablePolymorphicObject { protected: bool check_impl(uint8 stoppingId, bool setFinalized, - array* stop_status, bool* one_changed, + array* stop_status, bool* indicators, const Updater&) override; explicit Time(std::shared_ptr exec) diff --git a/omp/stop/residual_norm_kernels.cpp b/omp/stop/residual_norm_kernels.cpp index 0ec4395a16b..1018552b1ec 100644 --- a/omp/stop/residual_norm_kernels.cpp +++ b/omp/stop/residual_norm_kernels.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -28,7 +28,7 @@ void residual_norm(std::shared_ptr exec, ValueType rel_residual_goal, uint8 stoppingId, bool setFinalized, array* stop_status, array* device_storage, bool* all_converged, - bool* one_changed) + bool* indicators) { static_assert(is_complex_s::value == false, "ValueType must not be complex in this function!"); @@ -40,7 +40,7 @@ void residual_norm(std::shared_ptr exec, local_one_changed = true; } } - *one_changed = local_one_changed; + *indicators = local_one_changed; // No early stopping here because one cannot use break with parallel for // But it's parallel so does it matter? bool local_all_converged = true; @@ -75,7 +75,7 @@ void implicit_residual_norm( const matrix::Dense>* orig_tau, remove_complex rel_residual_goal, uint8 stoppingId, bool setFinalized, array* stop_status, - array* device_storage, bool* all_converged, bool* one_changed) + array* device_storage, bool* all_converged, bool* indicators) { bool local_one_changed = false; #pragma omp parallel for reduction(|| : local_one_changed) @@ -85,7 +85,7 @@ void implicit_residual_norm( local_one_changed = true; } } - *one_changed = local_one_changed; + *indicators = local_one_changed; // No early stopping here because one cannot use break with parallel for // But it's parallel so does it matter? bool local_all_converged = true; diff --git a/reference/stop/residual_norm_kernels.cpp b/reference/stop/residual_norm_kernels.cpp index ba2672edc28..541a5b292b8 100644 --- a/reference/stop/residual_norm_kernels.cpp +++ b/reference/stop/residual_norm_kernels.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -30,16 +30,16 @@ void residual_norm(std::shared_ptr exec, ValueType rel_residual_goal, uint8 stoppingId, bool setFinalized, array* stop_status, array* device_storage, bool* all_converged, - bool* one_changed) + bool* indicators) { static_assert(is_complex_s::value == false, "ValueType must not be complex in this function!"); *all_converged = true; - *one_changed = false; + *indicators = false; for (size_type i = 0; i < tau->get_size()[1]; ++i) { if (tau->at(i) <= rel_residual_goal * orig_tau->at(i)) { stop_status->get_data()[i].converge(stoppingId, setFinalized); - *one_changed = true; + *indicators = true; } } for (size_type i = 0; i < stop_status->get_size(); ++i) { @@ -72,14 +72,14 @@ void implicit_residual_norm( const matrix::Dense>* orig_tau, remove_complex rel_residual_goal, uint8 stoppingId, bool setFinalized, array* stop_status, - array* device_storage, bool* all_converged, bool* one_changed) + array* device_storage, bool* all_converged, bool* indicators) { *all_converged = true; - *one_changed = false; + *indicators = false; for (size_type i = 0; i < tau->get_size()[1]; ++i) { if (sqrt(abs(tau->at(i))) <= rel_residual_goal * orig_tau->at(i)) { stop_status->get_data()[i].converge(stoppingId, setFinalized); - *one_changed = true; + *indicators = true; } } for (size_type i = 0; i < stop_status->get_size(); ++i) { diff --git a/reference/test/stop/combined.cpp b/reference/test/stop/combined.cpp index 900e8131aba..917004bf3c9 100644 --- a/reference/test/stop/combined.cpp +++ b/reference/test/stop/combined.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -63,25 +63,26 @@ class Combined : public ::testing::Test { * the huge time picked. */ TEST_F(Combined, WaitsTillIteration) { - bool one_changed{}; + gko::array stop_indicators(exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(exec_, 1); stop_status.get_data()[0].reset(); constexpr gko::uint8 RelativeStoppingId{1}; auto criterion = factory_->generate(nullptr, nullptr, nullptr); gko::array converged(exec_, 1); - ASSERT_FALSE( - criterion->update() - .num_iterations(test_iterations - 1) - .check(RelativeStoppingId, true, &stop_status, &one_changed)); - ASSERT_TRUE( - criterion->update() - .num_iterations(test_iterations) - .check(RelativeStoppingId, true, &stop_status, &one_changed)); - ASSERT_TRUE( - criterion->update() - .num_iterations(test_iterations + 1) - .check(RelativeStoppingId, true, &stop_status, &one_changed)); + ASSERT_FALSE(criterion->update() + .num_iterations(test_iterations - 1) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); + ASSERT_TRUE(criterion->update() + .num_iterations(test_iterations) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); + ASSERT_TRUE(criterion->update() + .num_iterations(test_iterations + 1) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(static_cast(stop_status.get_data()[0].get_id()), 1); } @@ -102,7 +103,8 @@ TEST_F(Combined, WaitsTillTime) .on(exec_)) .on(exec_); unsigned int iters = 0; - bool one_changed{}; + gko::array stop_indicators(exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(exec_, 1); stop_status.get_data()[0].reset(); constexpr gko::uint8 RelativeStoppingId{1}; @@ -112,7 +114,8 @@ TEST_F(Combined, WaitsTillTime) for (int i = 0; i < testiters; i++) { sleep_millisecond(timelimit_ms / testiters); if (criterion->update().num_iterations(i).check( - RelativeStoppingId, true, &stop_status, &one_changed)) + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())) break; } auto time = std::chrono::steady_clock::now() - start; diff --git a/reference/test/stop/criterion_kernels.cpp b/reference/test/stop/criterion_kernels.cpp index 39ea9c72098..40ae9381771 100644 --- a/reference/test/stop/criterion_kernels.cpp +++ b/reference/test/stop/criterion_kernels.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -33,7 +33,8 @@ class Criterion : public ::testing::Test { TEST_F(Criterion, SetsOneStopStatus) { - bool one_changed{}; + gko::array stop_indicators(exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; auto criterion = factory_->generate(nullptr, nullptr, nullptr); gko::array stop_status(exec_, 1); @@ -41,7 +42,8 @@ TEST_F(Criterion, SetsOneStopStatus) criterion->update() .num_iterations(test_iterations) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); ASSERT_EQ(stop_status.get_data()[0].has_stopped(), true); } @@ -49,7 +51,8 @@ TEST_F(Criterion, SetsOneStopStatus) TEST_F(Criterion, SetsMultipleStopStatuses) { - bool one_changed{}; + gko::array stop_indicators(exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; auto criterion = factory_->generate(nullptr, nullptr, nullptr); gko::array stop_status(exec_, 3); @@ -59,7 +62,8 @@ TEST_F(Criterion, SetsMultipleStopStatuses) criterion->update() .num_iterations(test_iterations) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); ASSERT_EQ(stop_status.get_data()[0].has_stopped(), true); ASSERT_EQ(stop_status.get_data()[1].has_stopped(), true); diff --git a/reference/test/stop/iteration.cpp b/reference/test/stop/iteration.cpp index fbe53888c61..e58e0b6a369 100644 --- a/reference/test/stop/iteration.cpp +++ b/reference/test/stop/iteration.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -29,24 +29,25 @@ class Iteration : public ::testing::Test { TEST_F(Iteration, WaitsTillIteration) { - bool one_changed{}; + gko::array stop_indicators(exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(exec_, 1); stop_status.get_data()[0].reset(); constexpr gko::uint8 RelativeStoppingId{1}; auto criterion = factory_->generate(nullptr, nullptr, nullptr); - ASSERT_FALSE( - criterion->update() - .num_iterations(test_iterations - 1) - .check(RelativeStoppingId, true, &stop_status, &one_changed)); - ASSERT_TRUE( - criterion->update() - .num_iterations(test_iterations) - .check(RelativeStoppingId, true, &stop_status, &one_changed)); - ASSERT_TRUE( - criterion->update() - .num_iterations(test_iterations + 1) - .check(RelativeStoppingId, true, &stop_status, &one_changed)); + ASSERT_FALSE(criterion->update() + .num_iterations(test_iterations - 1) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); + ASSERT_TRUE(criterion->update() + .num_iterations(test_iterations) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); + ASSERT_TRUE(criterion->update() + .num_iterations(test_iterations + 1) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); } diff --git a/reference/test/stop/residual_norm_kernels.cpp b/reference/test/stop/residual_norm_kernels.cpp index e7eef0565d2..9fb2d8002d8 100644 --- a/reference/test/stop/residual_norm_kernels.cpp +++ b/reference/test/stop/residual_norm_kernels.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -93,14 +93,15 @@ TYPED_TEST(ResidualNorm, CheckIfResZeroConverges) .on(this->exec_) ->generate(mtx, rhs, x.get(), nullptr); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); EXPECT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); EXPECT_TRUE(stop_status.get_data()[0].has_converged()); - EXPECT_TRUE(one_changed); + EXPECT_TRUE(stop_indicators.get_const_data()[0]); } } @@ -145,15 +146,16 @@ TYPED_TEST(ResidualNorm, CanIgorneResidualNorm) auto criterion = this->rhs_factory_->generate(nullptr, scalar, nullptr, nullptr); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(criterion->update().ignore_residual_check(true).check( - RelativeStoppingId, true, &stop_status, &one_changed)); - ASSERT_THROW(criterion->update().check(RelativeStoppingId, true, - &stop_status, &one_changed), - gko::NotSupported); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); + ASSERT_THROW( + criterion->update().check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()), + gko::NotSupported); } @@ -175,69 +177,81 @@ TYPED_TEST(ResidualNorm, WaitsTillResidualGoal) auto rhs_norm = gko::initialize({100.0}, this->exec_); gko::as(rhs)->compute_norm2(rhs_norm); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(rhs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); res_norm->at(0) = r::value * 1.1 * rhs_norm->at(0); ASSERT_FALSE(rhs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); res_norm->at(0) = r::value * 0.9 * rhs_norm->at(0); ASSERT_TRUE(rhs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } { auto res_norm = gko::initialize({100.0}, this->exec_); auto init_res_val = res_norm->at(0, 0); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); res_norm->at(0) = r::value * 1.1 * init_res_val; ASSERT_FALSE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); res_norm->at(0) = r::value * 0.9 * init_res_val; ASSERT_TRUE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } { auto res_norm = gko::initialize({100.0}, this->exec_); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); res_norm->at(0) = r::value * 1.1; ASSERT_FALSE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); res_norm->at(0) = r::value * 0.9; ASSERT_TRUE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } } @@ -264,12 +278,14 @@ TYPED_TEST(ResidualNorm, SelfCalculatesThrowWithoutMatrix) auto rhs_norm = gko::initialize({100.0}, this->exec_); gko::as(rhs)->compute_norm2(rhs_norm); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_THROW(rhs_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed), + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()), gko::NotSupported); } { @@ -277,23 +293,27 @@ TYPED_TEST(ResidualNorm, SelfCalculatesThrowWithoutMatrix) auto solution = gko::initialize({rhs_val - initial_norm}, this->exec_); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_THROW(rel_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed), + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()), gko::NotSupported); } { auto solution = gko::initialize({rhs_val - T{100.0}}, this->exec_); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_THROW(abs_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed), + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()), gko::NotSupported); } } @@ -314,13 +334,15 @@ TYPED_TEST(ResidualNorm, RelativeSelfCalculatesThrowWithoutRhs) T initial_norm = 100.0; auto solution = gko::initialize({rhs_val - initial_norm}, this->exec_); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); - ASSERT_THROW(rel_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed), - gko::NotSupported); + ASSERT_THROW( + rel_criterion->update().solution(solution).check( + RelativeStoppingId, true, &stop_status, stop_indicators.get_data()), + gko::NotSupported); } @@ -347,72 +369,84 @@ TYPED_TEST(ResidualNorm, SelfCalculatesAndWaitsTillResidualGoal) auto rhs_norm = gko::initialize({100.0}, this->exec_); gko::as(rhs)->compute_norm2(rhs_norm); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(rhs_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); solution->at(0) = rhs_val - r::value * T{1.1} * rhs_norm->at(0); ASSERT_FALSE(rhs_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); solution->at(0) = rhs_val - r::value * T{0.5} * rhs_norm->at(0); ASSERT_TRUE(rhs_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } { T initial_norm = 100.0; auto solution = gko::initialize({rhs_val - initial_norm}, this->exec_); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(rel_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); solution->at(0) = rhs_val - r::value * T{1.1} * initial_norm; ASSERT_FALSE(rel_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); solution->at(0) = rhs_val - r::value * T{0.5} * initial_norm; ASSERT_TRUE(rel_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } { auto solution = gko::initialize({rhs_val - T{100.0}}, this->exec_); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(abs_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); // TODO FIXME: NVHPC calculates different result of rhs - r*1.2 from // rhs - tmp = rhs - (r * 1.2). https://godbolt.org/z/GrGE9PE67 solution->at(0) = rhs_val - r::value * T{1.4}; ASSERT_FALSE(abs_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); solution->at(0) = rhs_val - r::value * T{0.5}; ASSERT_TRUE(abs_criterion->update().solution(solution).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } } @@ -438,74 +472,86 @@ TYPED_TEST(ResidualNorm, WaitsTillResidualGoalMultipleRHS) auto rhs_norm = gko::initialize({I{100.0, 100.0}}, this->exec_); gko::as(rhs)->compute_norm2(rhs_norm); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 2); stop_status.get_data()[0].reset(); stop_status.get_data()[1].reset(); ASSERT_FALSE(rhs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); res_norm->at(0, 0) = r::value * 0.9 * rhs_norm->at(0, 0); ASSERT_FALSE(rhs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); res_norm->at(0, 1) = r::value * 0.9 * rhs_norm->at(0, 1); ASSERT_TRUE(rhs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[1].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } { auto res_norm = gko::initialize({I{100.0, 100.0}}, this->exec_); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 2); stop_status.get_data()[0].reset(); stop_status.get_data()[1].reset(); ASSERT_FALSE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); res_norm->at(0, 0) = r::value * 0.9 * res_norm->at(0, 0); ASSERT_FALSE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); res_norm->at(0, 1) = r::value * 0.9 * res_norm->at(0, 1); ASSERT_TRUE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[1].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } { auto res_norm = gko::initialize({I{100.0, 100.0}}, this->exec_); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 2); stop_status.get_data()[0].reset(); stop_status.get_data()[1].reset(); ASSERT_FALSE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); res_norm->at(0, 0) = r::value * 0.9; ASSERT_FALSE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); res_norm->at(0, 1) = r::value * 0.9; ASSERT_TRUE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[1].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } } @@ -557,25 +603,26 @@ TYPED_TEST(ResidualNormWithInitialResnorm, WaitsTillResidualGoal) auto init_res_val = res_norm->at(0, 0); auto criterion = this->factory_->generate(nullptr, rhs, nullptr, initial_res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); res_norm->at(0) = r::value * 1.1 * init_res_val; ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); res_norm->at(0) = r::value * 0.9 * init_res_val; ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } @@ -594,25 +641,26 @@ TYPED_TEST(ResidualNormWithInitialResnorm, std::shared_ptr mtx = gko::initialize({1.0}, this->exec_); auto criterion = this->factory_->generate(mtx, rhs, x.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(criterion->update().solution(x).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); x->at(0) = rhs_val - r::value * T{1.1} * initial_res; ASSERT_FALSE(criterion->update().solution(x).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); x->at(0) = rhs_val - r::value * T{0.5} * initial_res; ASSERT_TRUE(criterion->update().solution(x).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } @@ -628,26 +676,27 @@ TYPED_TEST(ResidualNormWithInitialResnorm, WaitsTillResidualGoalMultipleRHS) std::shared_ptr rhs = gko::initialize({I{10.0, 10.0}}, this->exec_); auto criterion = this->factory_->generate(nullptr, rhs, nullptr, res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 2); stop_status.get_data()[0].reset(); stop_status.get_data()[1].reset(); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); res_norm->at(0, 0) = r::value * 0.9 * res_norm->at(0, 0); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); res_norm->at(0, 1) = r::value * 0.9 * res_norm->at(0, 1); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[1].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } @@ -715,25 +764,26 @@ TYPED_TEST(ResidualNormWithRhsNorm, WaitsTillResidualGoal) auto res_norm = gko::initialize({100.0}, this->exec_); auto criterion = this->factory_->generate(nullptr, rhs, nullptr, initial_res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); res_norm->at(0) = r::value * 1.1 * rhs_norm->at(0); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); res_norm->at(0) = r::value * 0.9 * rhs_norm->at(0); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } @@ -752,26 +802,27 @@ TYPED_TEST(ResidualNormWithRhsNorm, WaitsTillResidualGoalMultipleRHS) gko::initialize({I{0.0, 0.0}}, this->exec_); gko::as(rhs)->compute_norm2(rhs_norm); auto criterion = this->factory_->generate(nullptr, rhs, nullptr, res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 2); stop_status.get_data()[0].reset(); stop_status.get_data()[1].reset(); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); res_norm->at(0, 0) = r::value * 0.9 * rhs_norm->at(0, 0); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); res_norm->at(0, 1) = r::value * 0.9 * rhs_norm->at(0, 1); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[1].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } @@ -847,16 +898,17 @@ TYPED_TEST(ImplicitResidualNorm, CheckIfResZeroConverges) .on(this->exec_) ->generate(mtx, rhs, x.get(), nullptr); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); - EXPECT_TRUE( - criterion->update() - .implicit_sq_residual_norm(implicit_sq_res_norm) - .check(RelativeStoppingId, true, &stop_status, &one_changed)); + EXPECT_TRUE(criterion->update() + .implicit_sq_residual_norm(implicit_sq_res_norm) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); EXPECT_TRUE(stop_status.get_data()[0].has_converged()); - EXPECT_TRUE(one_changed); + EXPECT_TRUE(stop_indicators.get_const_data()[0]); } } @@ -903,25 +955,26 @@ TYPED_TEST(ImplicitResidualNorm, WaitsTillResidualGoal) gko::as(rhs)->compute_norm2(rhs_norm); auto criterion = this->factory_->generate(nullptr, rhs, nullptr, initial_res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); res_norm->at(0) = std::pow(r::value * 1.1 * rhs_norm->at(0), 2); ASSERT_FALSE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); res_norm->at(0) = std::pow(r::value * 0.9 * rhs_norm->at(0), 2); ASSERT_TRUE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } @@ -939,28 +992,29 @@ TYPED_TEST(ImplicitResidualNorm, WaitsTillResidualGoalMultipleRHS) gko::initialize({I{0.0, 0.0}}, this->exec_); gko::as(rhs)->compute_norm2(rhs_norm); auto criterion = this->factory_->generate(nullptr, rhs, nullptr, res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 2); stop_status.get_data()[0].reset(); stop_status.get_data()[1].reset(); ASSERT_FALSE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); res_norm->at(0, 0) = std::pow(r::value * 0.9 * rhs_norm->at(0, 0), 2); ASSERT_FALSE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); res_norm->at(0, 1) = std::pow(r::value * 0.9 * rhs_norm->at(0, 1), 2); ASSERT_TRUE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[1].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } @@ -1024,25 +1078,26 @@ TYPED_TEST(ResidualNormWithAbsolute, WaitsTillResidualGoal) auto res_norm = gko::initialize({100.0}, this->exec_); auto criterion = this->factory_->generate(nullptr, rhs, nullptr, initial_res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 1); stop_status.get_data()[0].reset(); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); res_norm->at(0) = r::value * 1.1; ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), false); - ASSERT_EQ(one_changed, false); + ASSERT_EQ(stop_indicators.get_const_data()[0], false); res_norm->at(0) = r::value * 0.9; ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } @@ -1058,26 +1113,27 @@ TYPED_TEST(ResidualNormWithAbsolute, WaitsTillResidualGoalMultipleRHS) std::shared_ptr rhs = gko::initialize({I{10.0, 10.0}}, this->exec_); auto criterion = this->factory_->generate(nullptr, rhs, nullptr, res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->exec_, 2); stop_status.get_data()[0].reset(); stop_status.get_data()[1].reset(); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); res_norm->at(0, 0) = r::value * 0.9; ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[0].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); res_norm->at(0, 1) = r::value * 0.9; ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); ASSERT_EQ(stop_status.get_data()[1].has_converged(), true); - ASSERT_EQ(one_changed, true); + ASSERT_EQ(stop_indicators.get_const_data()[0], true); } diff --git a/reference/test/stop/time.cpp b/reference/test/stop/time.cpp index a5ea6107fbf..c41c43b500c 100644 --- a/reference/test/stop/time.cpp +++ b/reference/test/stop/time.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -64,15 +64,16 @@ TEST_F(Time, CanCreateCriterion) TEST_F(Time, WaitsTillTime) { auto criterion = factory_->generate(nullptr, nullptr, nullptr); - bool one_changed{}; + gko::array stop_indicators(exec_->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(exec_, 1); stop_status.get_data()[0].reset(); constexpr gko::uint8 RelativeStoppingId{1}; sleep_millisecond(test_ms); - ASSERT_TRUE(criterion->update().check(RelativeStoppingId, true, - &stop_status, &one_changed)); + ASSERT_TRUE(criterion->update().check( + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); } diff --git a/test/stop/criterion_kernels.cpp b/test/stop/criterion_kernels.cpp index 30280e848d8..34a614a69ef 100644 --- a/test/stop/criterion_kernels.cpp +++ b/test/stop/criterion_kernels.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -30,7 +30,8 @@ class Criterion : public CommonTestFixture { TEST_F(Criterion, SetsOneStopStatus) { - bool one_changed{}; + gko::array stop_indicators(ref, 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; auto criterion = factory->generate(nullptr, nullptr, nullptr); gko::array stop_status(ref, 1); @@ -39,7 +40,8 @@ TEST_F(Criterion, SetsOneStopStatus) stop_status.set_executor(exec); criterion->update() .num_iterations(test_iterations) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); stop_status.set_executor(ref); ASSERT_EQ(stop_status.get_data()[0].has_stopped(), true); @@ -48,7 +50,8 @@ TEST_F(Criterion, SetsOneStopStatus) TEST_F(Criterion, SetsMultipleStopStatuses) { - bool one_changed{}; + gko::array stop_indicators(ref, 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; auto criterion = factory->generate(nullptr, nullptr, nullptr); gko::array stop_status(ref, 3); @@ -59,7 +62,8 @@ TEST_F(Criterion, SetsMultipleStopStatuses) stop_status.set_executor(exec); criterion->update() .num_iterations(test_iterations) - .check(RelativeStoppingId, true, &stop_status, &one_changed); + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()); stop_status.set_executor(ref); ASSERT_EQ(stop_status.get_data()[0].has_stopped(), true); diff --git a/test/stop/residual_norm_kernels.cpp b/test/stop/residual_norm_kernels.cpp index 7be3e7fde48..e9f3aea543c 100644 --- a/test/stop/residual_norm_kernels.cpp +++ b/test/stop/residual_norm_kernels.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -70,16 +70,18 @@ TYPED_TEST(ResidualNorm, CanIgorneResidualNorm) auto criterion = this->factory->generate(nullptr, rhs, nullptr, initial_res.get()); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().ignore_residual_check(true).check( - RelativeStoppingId, true, &stop_status, &one_changed)); - ASSERT_THROW(criterion->update().check(RelativeStoppingId, true, - &stop_status, &one_changed), - gko::NotSupported); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); + ASSERT_THROW( + criterion->update().check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data()), + gko::NotSupported); } @@ -108,16 +110,18 @@ TYPED_TEST(ResidualNorm, CheckIfResZeroConverges) .on(this->exec) ->generate(mtx, rhs, x.get(), nullptr); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); EXPECT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); EXPECT_TRUE(stop_status.get_data()[0].has_converged()); - EXPECT_TRUE(one_changed); + EXPECT_TRUE(stop_indicators.get_const_data()[0]); } } @@ -139,82 +143,94 @@ TYPED_TEST(ResidualNorm, WaitsTillResidualGoal) auto rhs_norm = gko::initialize({100.0}, this->exec); gko::as(rhs)->compute_norm2(rhs_norm); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 1.1 * read(res_norm, 0, 0)); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_FALSE(stop_status.get_data()[0].has_converged()); - ASSERT_FALSE(one_changed); + ASSERT_FALSE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 0, r::value * 0.9 * read(res_norm, 0, 0)); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } { auto res_norm = gko::initialize({100.0}, this->exec); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); ASSERT_FALSE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 1.1 * read(res_norm, 0, 0)); ASSERT_FALSE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_FALSE(stop_status.get_data()[0].has_converged()); - ASSERT_FALSE(one_changed); + ASSERT_FALSE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 0, r::value * 0.9 * read(res_norm, 0, 0)); ASSERT_TRUE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } { auto res_norm = gko::initialize({100.0}, this->exec); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); ASSERT_FALSE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 1.1); ASSERT_FALSE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_FALSE(stop_status.get_data()[0].has_converged()); - ASSERT_FALSE(one_changed); + ASSERT_FALSE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 0, r::value * 0.9); ASSERT_TRUE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } } @@ -240,7 +256,8 @@ TYPED_TEST(ResidualNorm, WaitsTillResidualGoalMultipleRHS) auto rhs_norm = gko::initialize({I{100.0, 100.0}}, this->exec); gko::as(rhs)->compute_norm2(rhs_norm); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 2); stop_status.get_data()[0].reset(); @@ -248,28 +265,32 @@ TYPED_TEST(ResidualNorm, WaitsTillResidualGoalMultipleRHS) stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 0.9 * read(rhs_norm, 0, 0)); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 1, r::value * 0.9 * read(rhs_norm, 0, 1)); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[1].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } { auto res_norm = gko::initialize({I{100.0, 100.0}}, this->exec); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 2); stop_status.get_data()[0].reset(); @@ -277,28 +298,32 @@ TYPED_TEST(ResidualNorm, WaitsTillResidualGoalMultipleRHS) stop_status.set_executor(this->exec); ASSERT_FALSE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 0.9 * read(res_norm, 0, 0)); ASSERT_FALSE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 1, r::value * 0.9 * read(res_norm, 0, 1)); ASSERT_TRUE(rel_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[1].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } { auto res_norm = gko::initialize({I{100.0, 100.0}}, this->exec); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 2); stop_status.get_data()[0].reset(); @@ -306,22 +331,25 @@ TYPED_TEST(ResidualNorm, WaitsTillResidualGoalMultipleRHS) stop_status.set_executor(this->exec); ASSERT_FALSE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 0.9); ASSERT_FALSE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 1, r::value * 0.9); ASSERT_TRUE(abs_criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[1].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } } @@ -357,29 +385,30 @@ TYPED_TEST(ResidualNormWithInitialResnorm, WaitsTillResidualGoal) auto res_norm = gko::initialize({100.0}, this->exec); auto criterion = this->factory->generate(nullptr, rhs, nullptr, initial_res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 1.1 * read(res_norm, 0, 0)); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_FALSE(stop_status.get_data()[0].has_converged()); - ASSERT_FALSE(one_changed); + ASSERT_FALSE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 0, r::value * 0.9 * read(res_norm, 0, 0)); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } @@ -396,7 +425,8 @@ TYPED_TEST(ResidualNormWithInitialResnorm, WaitsTillResidualGoalMultipleRHS) std::shared_ptr rhs = gko::initialize({I{10.0, 10.0}}, this->exec); auto criterion = this->factory->generate(nullptr, rhs, nullptr, res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 2); stop_status.get_data()[0].reset(); @@ -404,22 +434,22 @@ TYPED_TEST(ResidualNormWithInitialResnorm, WaitsTillResidualGoalMultipleRHS) stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 0.9 * read(res_norm, 0, 0)); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 1, r::value * 0.9 * read(res_norm, 0, 1)); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[1].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } @@ -458,29 +488,30 @@ TYPED_TEST(ResidualNormWithRhsNorm, WaitsTillResidualGoal) auto res_norm = gko::initialize({100.0}, this->exec); auto criterion = this->factory->generate(nullptr, rhs, nullptr, initial_res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 1.1 * read(rhs_norm, 0, 0)); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_FALSE(stop_status.get_data()[0].has_converged()); - ASSERT_FALSE(one_changed); + ASSERT_FALSE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 0, r::value * 0.9 * read(rhs_norm, 0, 0)); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } @@ -500,7 +531,8 @@ TYPED_TEST(ResidualNormWithRhsNorm, WaitsTillResidualGoalMultipleRHS) gko::initialize({I{0.0, 0.0}}, this->exec); gko::as(rhs)->compute_norm2(rhs_norm); auto criterion = this->factory->generate(nullptr, rhs, nullptr, res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 2); stop_status.get_data()[0].reset(); @@ -508,22 +540,22 @@ TYPED_TEST(ResidualNormWithRhsNorm, WaitsTillResidualGoalMultipleRHS) stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 0.9 * read(rhs_norm, 0, 0)); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 1, r::value * 0.9 * read(rhs_norm, 0, 1)); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[1].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } @@ -574,18 +606,19 @@ TYPED_TEST(ImplicitResidualNorm, CheckIfResZeroConverges) .on(this->exec) ->generate(mtx, rhs, x.get(), nullptr); constexpr gko::uint8 RelativeStoppingId{1}; - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); - EXPECT_TRUE( - criterion->update() - .implicit_sq_residual_norm(implicit_sq_res_norm) - .check(RelativeStoppingId, true, &stop_status, &one_changed)); + EXPECT_TRUE(criterion->update() + .implicit_sq_residual_norm(implicit_sq_res_norm) + .check(RelativeStoppingId, true, &stop_status, + stop_indicators.get_data())); stop_status.set_executor(this->ref); EXPECT_TRUE(stop_status.get_data()[0].has_converged()); - EXPECT_TRUE(one_changed); + EXPECT_TRUE(stop_indicators.get_const_data()[0]); } } @@ -602,31 +635,32 @@ TYPED_TEST(ImplicitResidualNorm, WaitsTillResidualGoal) gko::as(rhs)->compute_norm2(rhs_norm); auto criterion = this->factory->generate(nullptr, rhs, nullptr, initial_res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); write(res_norm, 0, 0, std::pow(r::value * 1.1 * read(rhs_norm, 0, 0), 2)); ASSERT_FALSE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_FALSE(stop_status.get_data()[0].has_converged()); - ASSERT_FALSE(one_changed); + ASSERT_FALSE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 0, std::pow(r::value * 0.9 * read(rhs_norm, 0, 0), 2)); ASSERT_TRUE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } @@ -645,7 +679,8 @@ TYPED_TEST(ImplicitResidualNorm, WaitsTillResidualGoalMultipleRHS) gko::initialize({I{0.0, 0.0}}, this->exec); gko::as(rhs)->compute_norm2(rhs_norm); auto criterion = this->factory->generate(nullptr, rhs, nullptr, res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 2); stop_status.get_data()[0].reset(); @@ -653,24 +688,24 @@ TYPED_TEST(ImplicitResidualNorm, WaitsTillResidualGoalMultipleRHS) stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); write(res_norm, 0, 0, std::pow(r::value * 0.9 * read(rhs_norm, 0, 0), 2)); ASSERT_FALSE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 1, std::pow(r::value * 0.9 * read(rhs_norm, 0, 1), 2)); ASSERT_TRUE(criterion->update().implicit_sq_residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[1].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } @@ -705,29 +740,30 @@ TYPED_TEST(ResidualNormWithAbsolute, WaitsTillResidualGoal) auto res_norm = gko::initialize({100.0}, this->exec); auto criterion = this->factory->generate(nullptr, rhs, nullptr, initial_res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 1); stop_status.get_data()[0].reset(); stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 1.1); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_FALSE(stop_status.get_data()[0].has_converged()); - ASSERT_FALSE(one_changed); + ASSERT_FALSE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 0, r::value * 0.9); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); } @@ -744,7 +780,8 @@ TYPED_TEST(ResidualNormWithAbsolute, WaitsTillResidualGoalMultipleRHS) std::shared_ptr rhs = gko::initialize({I{10.0, 10.0}}, this->exec); auto criterion = this->factory->generate(nullptr, rhs, nullptr, res.get()); - bool one_changed{}; + gko::array stop_indicators(this->exec->get_master(), 2); + stop_indicators.get_data()[0] = false; constexpr gko::uint8 RelativeStoppingId{1}; gko::array stop_status(this->ref, 2); stop_status.get_data()[0].reset(); @@ -752,21 +789,21 @@ TYPED_TEST(ResidualNormWithAbsolute, WaitsTillResidualGoalMultipleRHS) stop_status.set_executor(this->exec); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); write(res_norm, 0, 0, r::value * 0.9); ASSERT_FALSE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[0].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); write(res_norm, 0, 1, r::value * 0.9); ASSERT_TRUE(criterion->update().residual_norm(res_norm).check( - RelativeStoppingId, true, &stop_status, &one_changed)); + RelativeStoppingId, true, &stop_status, stop_indicators.get_data())); stop_status.set_executor(this->ref); ASSERT_TRUE(stop_status.get_data()[1].has_converged()); - ASSERT_TRUE(one_changed); + ASSERT_TRUE(stop_indicators.get_const_data()[0]); stop_status.set_executor(this->exec); }