Merge (#1671): Unify batch functionality: solvers

Unify and simplify batch functionality: solvers, preconditioners, logger, stopping criteria Related PR: #1671
ginkgo-project · Aug 26, 2024 · 55e6bc2 · 55e6bc2
2 parents c09529f + 9f74b1b
commit 55e6bc2
Show file tree

Hide file tree

Showing 56 changed files with 882 additions and 609 deletions.
diff --git a/common/cuda_hip/base/batch_multi_vector_kernels.hpp b/common/cuda_hip/base/batch_multi_vector_kernels.hpp
@@ -2,8 +2,9 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
-#include <thrust/functional.h>
-#include <thrust/transform.h>
+#ifndef GKO_COMMON_CUDA_HIP_BASE_BATCH_MULTI_VECTOR_KERNELS_HPP_
+#define GKO_COMMON_CUDA_HIP_BASE_BATCH_MULTI_VECTOR_KERNELS_HPP_
+
 
 #include <ginkgo/core/base/batch_multi_vector.hpp>
 #include <ginkgo/core/base/exception_helpers.hpp>
@@ -14,12 +15,9 @@
 #include "common/cuda_hip/base/config.hpp"
 #include "common/cuda_hip/base/math.hpp"
 #include "common/cuda_hip/base/runtime.hpp"
-#include "common/cuda_hip/base/thrust.hpp"
 #include "common/cuda_hip/base/types.hpp"
 #include "common/cuda_hip/components/cooperative_groups.hpp"
-#include "common/cuda_hip/components/format_conversion.hpp"
 #include "common/cuda_hip/components/reduction.hpp"
-#include "common/cuda_hip/components/segment_scan.hpp"
 #include "common/cuda_hip/components/thread_ids.hpp"
 #include "common/cuda_hip/components/warp_blas.hpp"
 
@@ -315,3 +313,6 @@ __global__ __launch_bounds__(default_block_size) void copy_kernel(
 }  // namespace GKO_DEVICE_NAMESPACE
 }  // namespace kernels
 }  // namespace gko
+
+
+#endif
diff --git a/common/cuda_hip/log/batch_logger.hpp.inc → common/cuda_hip/log/batch_logger.hpp b/common/cuda_hip/log/batch_logger.hpp.inc → common/cuda_hip/log/batch_logger.hpp
@@ -2,6 +2,18 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
+#ifndef GKO_COMMON_CUDA_HIP_LOG_BATCH_LOGGER_HPP_
+#define GKO_COMMON_CUDA_HIP_LOG_BATCH_LOGGER_HPP_
+
+
+#include <ginkgo/core/base/types.hpp>
+
+
+namespace gko {
+namespace kernels {
+namespace GKO_DEVICE_NAMESPACE {
+namespace batch_log {
+
 /**
  * @see reference/log/batch_logger.hpp
  */
@@ -28,3 +40,12 @@ class SimpleFinalLogger final {
     real_type* const final_residuals_;
     idx_type* const final_iters_;
 };
+
+
+}  // namespace batch_log
+}  // namespace GKO_DEVICE_NAMESPACE
+}  // namespace kernels
+}  // namespace gko
+
+
+#endif
diff --git a/common/cuda_hip/matrix/batch_csr_kernels.cpp b/common/cuda_hip/matrix/batch_csr_kernels.cpp
@@ -4,9 +4,6 @@
 
 #include "common/cuda_hip/matrix/batch_csr_kernels.hpp"
 
-#include <thrust/functional.h>
-#include <thrust/transform.h>
-
 #include <ginkgo/core/base/batch_multi_vector.hpp>
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/math.hpp>

diff --git a/common/cuda_hip/matrix/batch_csr_kernels.hpp b/common/cuda_hip/matrix/batch_csr_kernels.hpp
@@ -2,8 +2,9 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
-#include <thrust/functional.h>
-#include <thrust/transform.h>
+#ifndef GKO_COMMON_CUDA_HIP_MATRIX_BATCH_CSR_KERNELS_HPP_
+#define GKO_COMMON_CUDA_HIP_MATRIX_BATCH_CSR_KERNELS_HPP_
+
 
 #include <ginkgo/core/base/batch_multi_vector.hpp>
 #include <ginkgo/core/base/exception_helpers.hpp>
@@ -18,11 +19,7 @@
 #include "common/cuda_hip/base/thrust.hpp"
 #include "common/cuda_hip/base/types.hpp"
 #include "common/cuda_hip/components/cooperative_groups.hpp"
-#include "common/cuda_hip/components/format_conversion.hpp"
-#include "common/cuda_hip/components/reduction.hpp"
-#include "common/cuda_hip/components/segment_scan.hpp"
 #include "common/cuda_hip/components/thread_ids.hpp"
-#include "common/cuda_hip/components/warp_blas.hpp"
 #include "common/cuda_hip/matrix/batch_struct.hpp"
 
 
@@ -200,3 +197,6 @@ __global__ void add_scaled_identity_kernel(
 }  // namespace GKO_DEVICE_NAMESPACE
 }  // namespace kernels
 }  // namespace gko
+
+
+#endif
diff --git a/common/cuda_hip/matrix/batch_dense_kernels.cpp b/common/cuda_hip/matrix/batch_dense_kernels.cpp
@@ -4,9 +4,6 @@
 
 #include "common/cuda_hip/matrix/batch_dense_kernels.hpp"
 
-#include <thrust/functional.h>
-#include <thrust/transform.h>
-
 #include <ginkgo/core/base/batch_multi_vector.hpp>
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/math.hpp>

diff --git a/common/cuda_hip/matrix/batch_dense_kernels.hpp b/common/cuda_hip/matrix/batch_dense_kernels.hpp
@@ -2,8 +2,9 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
-#include <thrust/functional.h>
-#include <thrust/transform.h>
+#ifndef GKO_COMMON_CUDA_HIP_MATRIX_BATCH_DENSE_KERNELS_HPP_
+#define GKO_COMMON_CUDA_HIP_MATRIX_BATCH_DENSE_KERNELS_HPP_
+
 
 #include <ginkgo/core/base/batch_multi_vector.hpp>
 #include <ginkgo/core/base/exception_helpers.hpp>
@@ -15,12 +16,8 @@
 #include "common/cuda_hip/base/config.hpp"
 #include "common/cuda_hip/base/math.hpp"
 #include "common/cuda_hip/base/runtime.hpp"
-#include "common/cuda_hip/base/thrust.hpp"
 #include "common/cuda_hip/base/types.hpp"
 #include "common/cuda_hip/components/cooperative_groups.hpp"
-#include "common/cuda_hip/components/format_conversion.hpp"
-#include "common/cuda_hip/components/reduction.hpp"
-#include "common/cuda_hip/components/segment_scan.hpp"
 #include "common/cuda_hip/components/thread_ids.hpp"
 #include "common/cuda_hip/components/warp_blas.hpp"
 #include "common/cuda_hip/matrix/batch_struct.hpp"
@@ -247,3 +244,6 @@ __global__ void add_scaled_identity_kernel(
 }  // namespace GKO_DEVICE_NAMESPACE
 }  // namespace kernels
 }  // namespace gko
+
+
+#endif
diff --git a/common/cuda_hip/matrix/batch_ell_kernels.cpp b/common/cuda_hip/matrix/batch_ell_kernels.cpp
@@ -4,9 +4,6 @@
 
 #include "common/cuda_hip/matrix/batch_ell_kernels.hpp"
 
-#include <thrust/functional.h>
-#include <thrust/transform.h>
-
 #include <ginkgo/core/base/batch_multi_vector.hpp>
 #include <ginkgo/core/base/exception_helpers.hpp>
 #include <ginkgo/core/base/math.hpp>

diff --git a/common/cuda_hip/matrix/batch_ell_kernels.hpp b/common/cuda_hip/matrix/batch_ell_kernels.hpp
@@ -2,8 +2,9 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
-#include <thrust/functional.h>
-#include <thrust/transform.h>
+#ifndef GKO_COMMON_CUDA_HIP_MATRIX_BATCH_ELL_KERNELS_HPP_
+#define GKO_COMMON_CUDA_HIP_MATRIX_BATCH_ELL_KERNELS_HPP_
+
 
 #include <ginkgo/core/base/batch_multi_vector.hpp>
 #include <ginkgo/core/base/exception_helpers.hpp>
@@ -15,14 +16,9 @@
 #include "common/cuda_hip/base/config.hpp"
 #include "common/cuda_hip/base/math.hpp"
 #include "common/cuda_hip/base/runtime.hpp"
-#include "common/cuda_hip/base/thrust.hpp"
 #include "common/cuda_hip/base/types.hpp"
 #include "common/cuda_hip/components/cooperative_groups.hpp"
-#include "common/cuda_hip/components/format_conversion.hpp"
-#include "common/cuda_hip/components/reduction.hpp"
-#include "common/cuda_hip/components/segment_scan.hpp"
 #include "common/cuda_hip/components/thread_ids.hpp"
-#include "common/cuda_hip/components/warp_blas.hpp"
 #include "common/cuda_hip/matrix/batch_struct.hpp"
 
 
@@ -210,3 +206,6 @@ __global__ void add_scaled_identity_kernel(
 }  // namespace GKO_DEVICE_NAMESPACE
 }  // namespace kernels
 }  // namespace gko
+
+
+#endif
diff --git a/...preconditioner/batch_block_jacobi.hpp.inc → ...hip/preconditioner/batch_block_jacobi.hpp b/...preconditioner/batch_block_jacobi.hpp.inc → ...hip/preconditioner/batch_block_jacobi.hpp
@@ -2,6 +2,30 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
+#ifndef GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_BLOCK_JACOBI_HPP_
+#define GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_BLOCK_JACOBI_HPP_
+
+
+#include <ginkgo/core/base/exception_helpers.hpp>
+#include <ginkgo/core/base/math.hpp>
+#include <ginkgo/core/base/types.hpp>
+#include <ginkgo/core/preconditioner/batch_jacobi.hpp>
+
+#include "common/cuda_hip/base/batch_struct.hpp"
+#include "common/cuda_hip/base/config.hpp"
+#include "common/cuda_hip/base/math.hpp"
+#include "common/cuda_hip/base/types.hpp"
+#include "common/cuda_hip/components/cooperative_groups.hpp"
+#include "common/cuda_hip/matrix/batch_struct.hpp"
+#include "core/preconditioner/batch_jacobi_helpers.hpp"
+
+
+namespace gko {
+namespace kernels {
+namespace GKO_DEVICE_NAMESPACE {
+namespace batch_preconditioner {
+
+
 /**
  * BlockJacobi preconditioner for batch solvers.
  */
@@ -173,3 +197,12 @@ class BlockJacobi final {
     const int* __restrict__ const block_ptrs_arr_;
     const int* __restrict__ const row_block_map_;
 };
+
+
+}  // namespace batch_preconditioner
+}  // namespace GKO_DEVICE_NAMESPACE
+}  // namespace kernels
+}  // namespace gko
+
+
+#endif
diff --git a/...hip/preconditioner/batch_identity.hpp.inc → ...uda_hip/preconditioner/batch_identity.hpp b/...hip/preconditioner/batch_identity.hpp.inc → ...uda_hip/preconditioner/batch_identity.hpp
@@ -2,6 +2,25 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
+#ifndef GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_IDENTITY_HPP_
+#define GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_IDENTITY_HPP_
+
+
+#include <ginkgo/core/base/exception_helpers.hpp>
+#include <ginkgo/core/base/math.hpp>
+#include <ginkgo/core/base/types.hpp>
+
+#include "common/cuda_hip/base/batch_struct.hpp"
+#include "common/cuda_hip/base/config.hpp"
+#include "common/cuda_hip/matrix/batch_struct.hpp"
+
+
+namespace gko {
+namespace kernels {
+namespace GKO_DEVICE_NAMESPACE {
+namespace batch_preconditioner {
+
+
 /**
  * @see reference/preconditioner/batch_identity.hpp
  */
@@ -31,3 +50,12 @@ class Identity final {
         }
     }
 };
+
+
+}  // namespace batch_preconditioner
+}  // namespace GKO_DEVICE_NAMESPACE
+}  // namespace kernels
+}  // namespace gko
+
+
+#endif
diff --git a/...econditioner/batch_jacobi_kernels.hpp.inc → ...p/preconditioner/batch_jacobi_kernels.hpp b/...econditioner/batch_jacobi_kernels.hpp.inc → ...p/preconditioner/batch_jacobi_kernels.hpp
@@ -2,6 +2,34 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
+#ifndef GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_JACOBI_KERNELS_HPP_
+#define GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_JACOBI_KERNELS_HPP_
+
+
+#include <ginkgo/core/base/exception_helpers.hpp>
+#include <ginkgo/core/base/math.hpp>
+#include <ginkgo/core/base/types.hpp>
+
+#include "common/cuda_hip/base/batch_multi_vector_kernels.hpp"
+#include "common/cuda_hip/base/batch_struct.hpp"
+#include "common/cuda_hip/base/config.hpp"
+#include "common/cuda_hip/base/math.hpp"
+#include "common/cuda_hip/base/runtime.hpp"
+#include "common/cuda_hip/base/types.hpp"
+#include "common/cuda_hip/components/cooperative_groups.hpp"
+#include "common/cuda_hip/components/thread_ids.hpp"
+#include "common/cuda_hip/matrix/batch_csr_kernels.hpp"
+#include "common/cuda_hip/matrix/batch_dense_kernels.hpp"
+#include "common/cuda_hip/matrix/batch_ell_kernels.hpp"
+#include "common/cuda_hip/matrix/batch_struct.hpp"
+
+
+namespace gko {
+namespace kernels {
+namespace GKO_DEVICE_NAMESPACE {
+namespace batch_single_kernels {
+
+
 __global__ void compute_block_storage_kernel(
     const gko::size_type num_blocks,
     const int* const __restrict__ block_pointers,
@@ -243,3 +271,12 @@ __launch_bounds__(default_block_size) void compute_block_jacobi_kernel(
         }
     }
 }
+
+
+}  // namespace batch_single_kernels
+}  // namespace GKO_DEVICE_NAMESPACE
+}  // namespace kernels
+}  // namespace gko
+
+
+#endif
diff --git a/common/cuda_hip/preconditioner/batch_preconditioners.hpp b/common/cuda_hip/preconditioner/batch_preconditioners.hpp
@@ -0,0 +1,17 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_PRECONDITIONERS_HPP_
+#define GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_PRECONDITIONERS_HPP_
+
+
+#include "common/cuda_hip/preconditioner/batch_block_jacobi.hpp"
+#include "common/cuda_hip/preconditioner/batch_identity.hpp"
+#include "common/cuda_hip/preconditioner/batch_scalar_jacobi.hpp"
+#include "core/base/batch_struct.hpp"
+#include "core/matrix/batch_struct.hpp"
+#include "core/preconditioner/batch_jacobi_helpers.hpp"
+
+
+#endif  // GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_PRECONDITIONERS_HPP_
diff --git a/...reconditioner/batch_scalar_jacobi.hpp.inc → ...ip/preconditioner/batch_scalar_jacobi.hpp b/...reconditioner/batch_scalar_jacobi.hpp.inc → ...ip/preconditioner/batch_scalar_jacobi.hpp
@@ -2,6 +2,29 @@
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
+#ifndef GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_SCALAR_JACOBI_HPP_
+#define GKO_COMMON_CUDA_HIP_PRECONDITIONER_BATCH_SCALAR_JACOBI_HPP_
+
+
+#include <ginkgo/core/base/exception_helpers.hpp>
+#include <ginkgo/core/base/math.hpp>
+#include <ginkgo/core/base/types.hpp>
+#include <ginkgo/core/preconditioner/batch_jacobi.hpp>
+
+#include "common/cuda_hip/base/batch_struct.hpp"
+#include "common/cuda_hip/base/config.hpp"
+#include "common/cuda_hip/base/math.hpp"
+#include "common/cuda_hip/base/types.hpp"
+#include "common/cuda_hip/components/cooperative_groups.hpp"
+#include "common/cuda_hip/matrix/batch_struct.hpp"
+
+
+namespace gko {
+namespace kernels {
+namespace GKO_DEVICE_NAMESPACE {
+namespace batch_preconditioner {
+
+
 /**
  * (Scalar) Jacobi preconditioner for batch solvers.
  */
@@ -132,3 +155,12 @@ class ScalarJacobi final {
 private:
     value_type* __restrict__ work_;
 };
+
+
+}  // namespace batch_preconditioner
+}  // namespace GKO_DEVICE_NAMESPACE
+}  // namespace kernels
+}  // namespace gko
+
+
+#endif