Merge (#1438): Add a batch::Bicgstab solver class, core, ref and omp …

…kernels

This PR adds a batch::Bicgstab solver and the reference (and OpenMP) kernels

In addition, some general solver, stopping critieria, logger and preconditioner framework is also added.

1. Batch stopping criteria
2. Simple batch logger
3. Some batch matrix generation utilities
4. A basic batch Identity matrix class and a corresponding Identity preconditioner to enable unpreconditioned solves.
5. The batch dispatch mechanism that selects the correct matrix, solver, preconditioner, stopping critieria at runtime and dispatches the correct kernel on the device.

Related PR: #1438

common/cuda_hip/log/batch_logger.hpp.inc

@@ -0,0 +1,56 @@
+/*******************************<GINKGO LICENSE>******************************
+Copyright (c) 2017-2023, the Ginkgo authors
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+1. Redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
+IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
+PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+******************************<GINKGO LICENSE>*******************************/
+
+/**
+ * @see reference/log/batch_logger.hpp
+ */
+template <typename RealType>
+class SimpleFinalLogger final {
+public:
+    using real_type = remove_complex<RealType>;
+
+    SimpleFinalLogger(real_type* const batch_residuals, int* const batch_iters)
+        : final_residuals_{batch_residuals}, final_iters_{batch_iters}
+    {}
+
+    __device__ __forceinline__ void log_iteration(const size_type batch_idx,
+                                                  const int iter,
+                                                  const real_type res_norm)
+    {
+        final_iters_[batch_idx] = iter;
+        final_residuals_[batch_idx] = res_norm;
+    }
+
+private:
+    real_type* const final_residuals_;
+    int* const final_iters_;
+};

common/cuda_hip/preconditioner/batch_identity.hpp.inc

@@ -0,0 +1,68 @@
+/*******************************<GINKGO LICENSE>******************************
+Copyright (c) 2017-2023, the Ginkgo authors
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+1. Redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
+IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
+PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+******************************<GINKGO LICENSE>*******************************/
+
+/**
+ * @see reference/preconditioner/batch_identity.hpp
+ */
+template <typename ValueType>
+class Identity final {
+public:
+    using value_type = ValueType;
+
+    static constexpr int work_size = 0;
+
+    __host__ __device__ static constexpr int dynamic_work_size(int, int)
+    {
+        return 0;
+    }
+
+    __device__ __forceinline__ void generate(
+        size_type,
+        const gko::batch::matrix::ell::batch_item<const ValueType, gko::int32>&,
+        ValueType*)
+    {}
+
+    __device__ __forceinline__ void generate(
+        size_type,
+        const gko::batch::matrix::dense::batch_item<const ValueType>&,
+        ValueType*)
+    {}
+
+    __device__ __forceinline__ void apply(const int num_rows,
+                                          const ValueType* const r,
+                                          ValueType* const z) const
+    {
+        for (int li = threadIdx.x; li < num_rows; li += blockDim.x) {
+            z[li] = r[li];
+        }
+    }
+};

common/cuda_hip/stop/batch_criteria.hpp.inc

@@ -0,0 +1,80 @@
+/*******************************<GINKGO LICENSE>******************************
+Copyright (c) 2017-2023, the Ginkgo authors
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+1. Redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
+IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
+PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+******************************<GINKGO LICENSE>*******************************/
+
+
+/**
+ * @see reference/stop/batch_criteria.hpp
+ */
+template <typename ValueType>
+class SimpleRelResidual {
+public:
+    using real_type = remove_complex<ValueType>;
+
+    __device__ __forceinline__ SimpleRelResidual(
+        const real_type rel_res_tol, const real_type* const rhs_b_norms)
+        : rel_tol_{rel_res_tol}, rhs_norms_{rhs_b_norms}
+    {}
+
+    __device__ __forceinline__ bool check_converged(
+        const real_type* const residual_norms) const
+    {
+        return residual_norms[0] <= (rel_tol_ * rhs_norms_[0]);
+    }
+
+private:
+    const real_type rel_tol_;
+    const real_type* const rhs_norms_;
+};
+
+
+/**
+ * @see reference/stop/batch_criteria.hpp
+ */
+template <typename ValueType>
+class SimpleAbsResidual {
+public:
+    using real_type = remove_complex<ValueType>;
+
+    __device__ __forceinline__ SimpleAbsResidual(const real_type tol,
+                                                 const real_type*)
+        : abs_tol_{tol}
+    {}
+
+    __device__ __forceinline__ bool check_converged(
+        const real_type* const residual_norms) const
+    {
+        return (residual_norms[0] <= abs_tol_);
+    }
+
+private:
+    const real_type abs_tol_;
+};

core/CMakeLists.txt

@@ -29,6 +29,7 @@ target_sources(ginkgo
     factorization/par_ilu.cpp
     factorization/par_ilut.cpp
     factorization/symbolic.cpp
+    log/batch_logger.cpp
     log/convergence.cpp
     log/logger.cpp
     log/performance_hint.cpp
@@ -41,6 +42,7 @@ target_sources(ginkgo
     log/stream.cpp
     matrix/batch_dense.cpp
     matrix/batch_ell.cpp
+    matrix/batch_identity.cpp
     matrix/coo.cpp
     matrix/csr.cpp
     matrix/dense.cpp
@@ -62,6 +64,7 @@ target_sources(ginkgo
     reorder/amd.cpp
     reorder/rcm.cpp
     reorder/scaled_reordered.cpp
+    solver/batch_bicgstab.cpp
     solver/bicg.cpp
     solver/bicgstab.cpp
     solver/cb_gmres.cpp

core/base/batch_struct.hpp

@@ -71,9 +71,9 @@ struct uniform_batch {
     int32 num_rows;
     int32 num_rhs;
 
-    size_type get_entry_storage() const
+    inline size_type get_single_item_num_nnz() const
     {
-        return num_rows * stride * sizeof(value_type);
+        return static_cast<size_type>(stride * num_rows);
     }
 };
 

core/base/batch_utilities.hpp

@@ -201,6 +201,8 @@ std::unique_ptr<OutputType> read(
                            std::forward<TArgs>(create_args)...);
 
     for (size_type b = 0; b < num_batch_items; ++b) {
+        if (data.at(b).size != data.at(0).size)
+            GKO_INVALID_STATE("Incorrect data passed in");
         tmp->create_view_for_item(b)->read(data[b]);
     }
 

core/device_hooks/common_kernels.inc.cpp

@@ -75,6 +75,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "core/preconditioner/isai_kernels.hpp"
 #include "core/preconditioner/jacobi_kernels.hpp"
 #include "core/reorder/rcm_kernels.hpp"
+#include "core/solver/batch_bicgstab_kernels.hpp"
 #include "core/solver/bicg_kernels.hpp"
 #include "core/solver/bicgstab_kernels.hpp"
 #include "core/solver/cb_gmres_kernels.hpp"
@@ -414,6 +415,15 @@ GKO_STUB_VALUE_AND_INDEX_TYPE(GKO_DECLARE_DIAGONAL_FILL_IN_MATRIX_DATA_KERNEL);
 }  // namespace diagonal
 
 
+namespace batch_bicgstab {
+
+
+GKO_STUB_VALUE_TYPE(GKO_DECLARE_BATCH_BICGSTAB_APPLY_KERNEL);
+
+
+}  // namespace batch_bicgstab
+
+
 namespace cg {
 
 

core/log/batch_logger.cpp

@@ -0,0 +1,70 @@
+/*******************************<GINKGO LICENSE>******************************
+Copyright (c) 2017-2023, the Ginkgo authors
+All rights reserved.
+
+Redistribution and use in source and binary forms, with or without
+modification, are permitted provided that the following conditions
+are met:
+
+1. Redistributions of source code must retain the above copyright
+notice, this list of conditions and the following disclaimer.
+
+2. Redistributions in binary form must reproduce the above copyright
+notice, this list of conditions and the following disclaimer in the
+documentation and/or other materials provided with the distribution.
+
+3. Neither the name of the copyright holder nor the names of its
+contributors may be used to endorse or promote products derived from
+this software without specific prior written permission.
+
+THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS
+IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED
+TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A
+PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
+HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
+SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
+LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
+DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
+THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+******************************<GINKGO LICENSE>*******************************/
+
+#include <ginkgo/core/log/batch_logger.hpp>
+
+
+#include <ginkgo/core/base/array.hpp>
+#include <ginkgo/core/base/batch_multi_vector.hpp>
+#include <ginkgo/core/base/math.hpp>
+
+
+namespace gko {
+namespace batch {
+namespace log {
+
+
+template <typename ValueType>
+void BatchConvergence<ValueType>::on_batch_solver_completed(
+    const array<int>& iteration_count,
+    const array<remove_complex<ValueType>>& residual_norm) const
+{
+    if (this->iteration_count_.get_num_elems() == 0) {
+        this->iteration_count_ = gko::array<int>(
+            iteration_count.get_executor(), iteration_count.get_num_elems());
+    }
+    if (this->residual_norm_.get_num_elems() == 0) {
+        this->residual_norm_ = gko::array<remove_complex<ValueType>>(
+            residual_norm.get_executor(), residual_norm.get_num_elems());
+    }
+    this->iteration_count_ = iteration_count;
+    this->residual_norm_ = residual_norm;
+}
+
+
+#define GKO_DECLARE_BATCH_CONVERGENCE(_type) class BatchConvergence<_type>
+GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BATCH_CONVERGENCE);
+
+
+}  // namespace log
+}  // namespace batch
+}  // namespace gko

core/log/logger.cpp

@@ -75,6 +75,8 @@ constexpr Logger::mask_type Logger::linop_factory_generate_completed_mask;
 constexpr Logger::mask_type Logger::criterion_check_started_mask;
 constexpr Logger::mask_type Logger::criterion_check_completed_mask;
 
+constexpr Logger::mask_type Logger::batch_solver_completed_mask;
+
 constexpr Logger::mask_type Logger::iteration_complete_mask;