From 943f110dc76b730eec1968ce0b423e0dd13239a0 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Tue, 27 Aug 2019 17:54:05 -0400 Subject: [PATCH] Allow using PreconditionChebyshev for parallel::distributed::Vector with MemorySpace::CUDA --- include/deal.II/lac/affine_constraints.h | 35 +++ .../lac/affine_constraints.templates.h | 15 +- include/deal.II/lac/precondition.h | 110 ++++++- tests/cuda/matrix_vector_mf.h | 32 +- tests/cuda/precondition_03.cu | 286 ++++++++++++++++++ ...h_mpi=true.with_p4est=true.mpirun=1.output | 10 + ...h_mpi=true.with_p4est=true.mpirun=2.output | 10 + 7 files changed, 468 insertions(+), 30 deletions(-) create mode 100644 tests/cuda/precondition_03.cu create mode 100644 tests/cuda/precondition_03.with_trilinos=true.with_mpi=true.with_p4est=true.mpirun=1.output create mode 100644 tests/cuda/precondition_03.with_trilinos=true.with_mpi=true.with_p4est=true.mpirun=2.output diff --git a/include/deal.II/lac/affine_constraints.h b/include/deal.II/lac/affine_constraints.h index 7a710868cf..993f59b072 100644 --- a/include/deal.II/lac/affine_constraints.h +++ b/include/deal.II/lac/affine_constraints.h @@ -54,6 +54,27 @@ namespace internals class GlobalRowsFromLocal; } +namespace internal +{ + namespace AffineConstraintsImplementation + { + template + void + set_zero_all(const std::vector &cm, + VectorType & vec); + + template + void + set_zero_all(const std::vector &cm, + dealii::Vector & vec); + + template + void + set_zero_all(const std::vector &cm, + dealii::BlockVector & vec); + } // namespace AffineConstraintsImplementation +} // namespace internal + template class AffineConstraints; @@ -1675,6 +1696,20 @@ AffineConstraints::set_inhomogeneity(const size_type line_n, line_ptr->inhomogeneity = value; } +template +template +inline void +AffineConstraints::set_zero(VectorType &vec) const +{ + // since lines is a private member, we cannot pass it to the functions + // above. therefore, copy the content which is cheap + std::vector constrained_lines(lines.size()); + for (unsigned int i = 0; i < lines.size(); ++i) + constrained_lines[i] = lines[i].index; + internal::AffineConstraintsImplementation::set_zero_all(constrained_lines, + vec); +} + template inline types::global_dof_index AffineConstraints::n_constraints() const diff --git a/include/deal.II/lac/affine_constraints.templates.h b/include/deal.II/lac/affine_constraints.templates.h index 4670b160ae..ebe4835f64 100644 --- a/include/deal.II/lac/affine_constraints.templates.h +++ b/include/deal.II/lac/affine_constraints.templates.h @@ -1904,6 +1904,7 @@ namespace internal size_type shift = 0) { Assert(shift == 0, ExcNotImplemented()); + (void)shift; std::vector constrained_local_dofs_host; constrained_local_dofs_host.reserve(cm.size()); @@ -1993,20 +1994,6 @@ namespace internal } // namespace AffineConstraintsImplementation } // namespace internal -template -template -void -AffineConstraints::set_zero(VectorType &vec) const -{ - // since lines is a private member, we cannot pass it to the functions - // above. therefore, copy the content which is cheap - std::vector constrained_lines(lines.size()); - for (unsigned int i = 0; i < lines.size(); ++i) - constrained_lines[i] = lines[i].index; - internal::AffineConstraintsImplementation::set_zero_all(constrained_lines, - vec); -} - template template void diff --git a/include/deal.II/lac/precondition.h b/include/deal.II/lac/precondition.h index 2f887a3826..01cbd8001b 100644 --- a/include/deal.II/lac/precondition.h +++ b/include/deal.II/lac/precondition.h @@ -20,6 +20,7 @@ #include +#include #include #include #include @@ -27,6 +28,7 @@ #include #include +#include #include #include #include @@ -1028,6 +1030,12 @@ public: const double eig_cg_residual = 1e-2, const double max_eigenvalue = 1); + /** + * Copy assignment operator. + */ + AdditionalData & + operator=(const AdditionalData &other_data); + /** * This determines the degree of the Chebyshev polynomial. The degree of * the polynomial gives the number of matrix-vector products to be @@ -1075,6 +1083,12 @@ public: */ double max_eigenvalue; + /** + * Constraints to be used for the operator given. This variable is used to + * zero out the correct entries when creating an initial guess. + */ + AffineConstraints constraints; + /** * Stores the preconditioner object that the Chebyshev is wrapped around. */ @@ -1947,19 +1961,21 @@ namespace internal } // selection for diagonal matrix around parallel deal.II vector - template + template inline void vector_updates( - const LinearAlgebra::distributed::Vector &rhs, + const LinearAlgebra::distributed::Vector &rhs, const DiagonalMatrix< - LinearAlgebra::distributed::Vector> &jacobi, - const unsigned int iteration_index, - const double factor1, - const double factor2, - LinearAlgebra::distributed::Vector &solution_old, - LinearAlgebra::distributed::Vector &temp_vector1, - LinearAlgebra::distributed::Vector &, - LinearAlgebra::distributed::Vector &solution) + LinearAlgebra::distributed::Vector> &jacobi, + const unsigned int iteration_index, + const double factor1, + const double factor2, + LinearAlgebra::distributed::Vector + &solution_old, + LinearAlgebra::distributed::Vector + &temp_vector1, + LinearAlgebra::distributed::Vector &, + LinearAlgebra::distributed::Vector &solution) { VectorUpdater upd(rhs.begin(), jacobi.get_vector().begin(), @@ -2047,10 +2063,11 @@ namespace internal vector.add(-mean_value); } - template + template void set_initial_guess( - ::dealii::LinearAlgebra::distributed::Vector &vector) + ::dealii::LinearAlgebra::distributed::Vector + &vector) { // Choose a high-frequency mode consisting of numbers between 0 and 1 // that is cheap to compute (cheaper than random numbers) but avoids @@ -2068,6 +2085,54 @@ namespace internal vector.add(-mean_value); } + +# ifdef DEAL_II_COMPILER_CUDA_AWARE + template + __global__ void + set_initial_guess_kernel(const types::global_dof_index offset, + const unsigned int local_size, + Number * values) + + { + const unsigned int index = threadIdx.x + blockDim.x * blockIdx.x; + if (index < local_size) + values[index] = (index + offset) % 11; + } + + template + void + set_initial_guess( + ::dealii::LinearAlgebra::distributed::Vector + &vector) + { + // Choose a high-frequency mode consisting of numbers between 0 and 1 + // that is cheap to compute (cheaper than random numbers) but avoids + // obviously re-occurring numbers in multi-component systems by choosing + // a period of 11. + // Make initial guess robust with respect to number of processors + // by operating on the global index. + types::global_dof_index first_local_range = 0; + if (!vector.locally_owned_elements().is_empty()) + first_local_range = vector.locally_owned_elements().nth_index_in_set(0); + + const auto n_local_elements = vector.local_size(); + const int n_blocks = + 1 + (n_local_elements - 1) / CUDAWrappers::block_size; + set_initial_guess_kernel<<>>( + first_local_range, n_local_elements, vector.get_values()); + +# ifdef DEBUG + // Check that the kernel was launched correctly + AssertCuda(cudaGetLastError()); + // Check that there was no problem during the execution of the kernel + AssertCuda(cudaDeviceSynchronize()); +# endif + + const Number mean_value = vector.mean_value(); + vector.add(-mean_value); + } +# endif // DEAL_II_COMPILER_CUDA_AWARE + struct EigenvalueTracker { public: @@ -2100,6 +2165,26 @@ inline PreconditionChebyshev:: +template +inline typename PreconditionChebyshev::AdditionalData & + PreconditionChebyshev:: + AdditionalData::operator=(const AdditionalData &other_data) +{ + degree = other_data.degree; + smoothing_range = other_data.smoothing_range; + eig_cg_n_iterations = other_data.eig_cg_n_iterations; + eig_cg_residual = other_data.eig_cg_residual; + max_eigenvalue = other_data.max_eigenvalue; + preconditioner = other_data.preconditioner; + constraints.copy_from(other_data.constraints); + + return *this; +} + + + template inline PreconditionChebyshev:: PreconditionChebyshev() @@ -2191,6 +2276,7 @@ PreconditionChebyshev:: // one entry is different to trigger high frequencies internal::PreconditionChebyshevImplementation::set_initial_guess( temp_vector1); + data.constraints.set_zero(temp_vector1); try { diff --git a/tests/cuda/matrix_vector_mf.h b/tests/cuda/matrix_vector_mf.h index 0fddfdcdd6..ef4a0c3d4f 100644 --- a/tests/cuda/matrix_vector_mf.h +++ b/tests/cuda/matrix_vector_mf.h @@ -161,7 +161,7 @@ template , int n_q_points_1d = fe_degree + 1> -class MatrixFreeTest +class MatrixFreeTest : public Subscriptor { public: MatrixFreeTest(const CUDAWrappers::MatrixFree &data_in, @@ -169,7 +169,18 @@ public: const bool constant_coeff = true); void - vmult(VectorType &dst, const VectorType &src); + vmult(VectorType &dst, const VectorType &src) const; + + Number + el(const unsigned int row, const unsigned int col) const; + + types::global_dof_index + m() const + { + return internal_m; + } + + types::global_dof_index internal_m; private: const CUDAWrappers::MatrixFree &data; @@ -200,7 +211,20 @@ MatrixFreeTest:: } } - +template +Number +MatrixFreeTest::el( + const unsigned int row, + const unsigned int col) const +{ + (void)col; + Assert(false, ExcNotImplemented()); + return 0.; +} template ::vmult( VectorType & dst, - const VectorType &src) + const VectorType &src) const { dst = static_cast(0.); HelmholtzOperator helmholtz_operator( diff --git a/tests/cuda/precondition_03.cu b/tests/cuda/precondition_03.cu new file mode 100644 index 0000000000..eba67855af --- /dev/null +++ b/tests/cuda/precondition_03.cu @@ -0,0 +1,286 @@ +// --------------------------------------------------------------------- +// +// Copyright (C) 2018 by the deal.II authors +// +// This file is part of the deal.II library. +// +// The deal.II library is free software; you can use it, redistribute +// it, and/or modify it under the terms of the GNU Lesser General +// Public License as published by the Free Software Foundation; either +// version 2.1 of the License, or (at your option) any later version. +// The full text of the license can be found in the file LICENSE.md at +// the top level directory of deal.II. +// +// --------------------------------------------------------------------- + + + +// this tests the correctness of MPI-parallel matrix free matrix-vector +// products by comparing the result with a Trilinos sparse matrix assembled in +// the usual way. The mesh is distributed among processors (hypercube) and has +// both hanging nodes (by randomly refining some cells, so the mesh is going +// to be different when run with different numbers of processors) and +// Dirichlet boundary conditions + +#include +#include + +#include + +#include +#include + +#include +#include + +#include +#include + +#include +#include +#include +#include + +#include + +#include + +#include "../tests.h" +#include "matrix_vector_mf.h" + + + +template +void +test() +{ + typedef double Number; + + parallel::distributed::Triangulation tria(MPI_COMM_WORLD); + GridGenerator::hyper_cube(tria); + tria.refine_global(1); + typename Triangulation::active_cell_iterator cell = tria.begin_active(), + endc = tria.end(); + for (; cell != endc; ++cell) + if (cell->is_locally_owned()) + if (cell->center().norm() < 0.2) + cell->set_refine_flag(); + tria.execute_coarsening_and_refinement(); + if (dim < 3 && fe_degree < 2) + tria.refine_global(2); + else + tria.refine_global(1); + if (tria.begin(tria.n_levels() - 1)->is_locally_owned()) + tria.begin(tria.n_levels() - 1)->set_refine_flag(); + if (tria.last()->is_locally_owned()) + tria.last()->set_refine_flag(); + tria.execute_coarsening_and_refinement(); + cell = tria.begin_active(); + for (unsigned int i = 0; i < 10 - 3 * dim; ++i) + { + cell = tria.begin_active(); + unsigned int counter = 0; + for (; cell != endc; ++cell, ++counter) + if (cell->is_locally_owned()) + if (counter % (7 - i) == 0) + cell->set_refine_flag(); + tria.execute_coarsening_and_refinement(); + } + + FE_Q fe(fe_degree); + DoFHandler dof(tria); + dof.distribute_dofs(fe); + + IndexSet owned_set = dof.locally_owned_dofs(); + IndexSet relevant_set; + DoFTools::extract_locally_relevant_dofs(dof, relevant_set); + + AffineConstraints constraints(relevant_set); + DoFTools::make_hanging_node_constraints(dof, constraints); + VectorTools::interpolate_boundary_values(dof, + 0, + Functions::ZeroFunction(), + constraints); + constraints.close(); + + deallog << "Testing " << dof.get_fe().get_name() << std::endl; + + MappingQGeneric mapping(fe_degree); + CUDAWrappers::MatrixFree mf_data; + const QGauss<1> quad(fe_degree + 1); + typename CUDAWrappers::MatrixFree::AdditionalData + additional_data; + additional_data.mapping_update_flags = update_values | update_gradients | + update_JxW_values | + update_quadrature_points; + mf_data.reinit(mapping, dof, constraints, quad, additional_data); + + const unsigned int coef_size = + tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim); + MatrixFreeTest> + mf(mf_data, coef_size); + mf.internal_m = owned_set.size(); + LinearAlgebra::distributed::Vector in_dev( + owned_set, MPI_COMM_WORLD); + LinearAlgebra::distributed::Vector out_dev( + owned_set, MPI_COMM_WORLD); + + LinearAlgebra::ReadWriteVector rw_in(owned_set); + for (unsigned int i = 0; i < in_dev.local_size(); ++i) + { + const unsigned int glob_index = owned_set.nth_index_in_set(i); + if (constraints.is_constrained(glob_index)) + continue; + rw_in.local_element(i) = random_value(); + } + in_dev.import(rw_in, VectorOperation::insert); + + // assemble trilinos sparse matrix with + // (\nabla v, \nabla u) + (v, 10 * u) for + // reference + LinearAlgebra::distributed::Vector in_host( + owned_set, MPI_COMM_WORLD); + in_host.import(rw_in, VectorOperation::insert); + LinearAlgebra::distributed::Vector ref( + owned_set, MPI_COMM_WORLD); + TrilinosWrappers::SparseMatrix sparse_matrix; + { + TrilinosWrappers::SparsityPattern csp(owned_set, MPI_COMM_WORLD); + DoFTools::make_sparsity_pattern(dof, + csp, + constraints, + true, + Utilities::MPI::this_mpi_process( + MPI_COMM_WORLD)); + csp.compress(); + sparse_matrix.reinit(csp); + } + { + QGauss quadrature_formula(fe_degree + 1); + + FEValues fe_values(dof.get_fe(), + quadrature_formula, + update_values | update_gradients | + update_JxW_values); + + const unsigned int dofs_per_cell = dof.get_fe().dofs_per_cell; + const unsigned int n_q_points = quadrature_formula.size(); + + FullMatrix cell_matrix(dofs_per_cell, dofs_per_cell); + std::vector local_dof_indices(dofs_per_cell); + + typename DoFHandler::active_cell_iterator cell = dof.begin_active(), + endc = dof.end(); + for (; cell != endc; ++cell) + if (cell->is_locally_owned()) + { + cell_matrix = 0; + fe_values.reinit(cell); + + for (unsigned int q_point = 0; q_point < n_q_points; ++q_point) + for (unsigned int i = 0; i < dofs_per_cell; ++i) + { + for (unsigned int j = 0; j < dofs_per_cell; ++j) + cell_matrix(i, j) += + ((fe_values.shape_grad(i, q_point) * + fe_values.shape_grad(j, q_point) + + 10. * fe_values.shape_value(i, q_point) * + fe_values.shape_value(j, q_point)) * + fe_values.JxW(q_point)); + } + + cell->get_dof_indices(local_dof_indices); + constraints.distribute_local_to_global(cell_matrix, + local_dof_indices, + sparse_matrix); + } + } + sparse_matrix.compress(VectorOperation::add); + + LinearAlgebra::distributed::Vector matrix_diagonal( + ref.get_partitioner()); + for (const auto index : matrix_diagonal.locally_owned_elements()) + matrix_diagonal[index] = sparse_matrix(index, index); + + using HostPreconditionerType = PreconditionChebyshev< + TrilinosWrappers::SparseMatrix, + LinearAlgebra::distributed::Vector>; + HostPreconditionerType precondition_chebyshev_host; + typename HostPreconditionerType::AdditionalData host_preconditioner_data; + host_preconditioner_data.preconditioner = std::make_shared>>( + matrix_diagonal); + host_preconditioner_data.constraints.copy_from(constraints); + precondition_chebyshev_host.initialize(sparse_matrix, + host_preconditioner_data); + + using DevicePreconditionerType = PreconditionChebyshev< + MatrixFreeTest< + dim, + fe_degree, + Number, + LinearAlgebra::distributed::Vector>, + LinearAlgebra::distributed::Vector>; + DevicePreconditionerType precondition_chebyshev_device; + typename DevicePreconditionerType::AdditionalData device_preconditioner_data; + device_preconditioner_data.preconditioner = std::make_shared>>( + LinearAlgebra::distributed::Vector( + ref.get_partitioner())); + device_preconditioner_data.preconditioner->get_vector().import( + matrix_diagonal, VectorOperation::insert); + device_preconditioner_data.constraints.copy_from(constraints); + precondition_chebyshev_device.initialize(mf, device_preconditioner_data); + + precondition_chebyshev_device.vmult(out_dev, in_dev); + + LinearAlgebra::distributed::Vector out_host( + owned_set, MPI_COMM_WORLD); + LinearAlgebra::ReadWriteVector rw_out(owned_set); + rw_out.import(out_dev, VectorOperation::insert); + out_host.import(rw_out, VectorOperation::insert); + + precondition_chebyshev_host.vmult(ref, in_host); + out_host -= ref; + + const double diff_norm = out_host.linfty_norm(); + + deallog << "Norm of difference: " << diff_norm << std::endl << std::endl; +} + + +int +main(int argc, char **argv) +{ + Utilities::MPI::MPI_InitFinalize mpi_initialization( + argc, argv, testing_max_num_threads()); + + unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD); + deallog.push(Utilities::int_to_string(myid)); + + init_cuda(true); + + if (myid == 0) + { + initlog(); + deallog << std::setprecision(4); + + deallog.push("2d"); + test<2, 1>(); + deallog.pop(); + + deallog.push("3d"); + test<3, 1>(); + test<3, 2>(); + deallog.pop(); + } + else + { + test<2, 1>(); + test<3, 1>(); + test<3, 2>(); + } +} diff --git a/tests/cuda/precondition_03.with_trilinos=true.with_mpi=true.with_p4est=true.mpirun=1.output b/tests/cuda/precondition_03.with_trilinos=true.with_mpi=true.with_p4est=true.mpirun=1.output new file mode 100644 index 0000000000..3b90478aad --- /dev/null +++ b/tests/cuda/precondition_03.with_trilinos=true.with_mpi=true.with_p4est=true.mpirun=1.output @@ -0,0 +1,10 @@ + +DEAL:0:2d::Testing FE_Q<2>(1) +DEAL:0:2d::Norm of difference: 0 +DEAL:0:2d:: +DEAL:0:3d::Testing FE_Q<3>(1) +DEAL:0:3d::Norm of difference: 0 +DEAL:0:3d:: +DEAL:0:3d::Testing FE_Q<3>(2) +DEAL:0:3d::Norm of difference: 0 +DEAL:0:3d:: diff --git a/tests/cuda/precondition_03.with_trilinos=true.with_mpi=true.with_p4est=true.mpirun=2.output b/tests/cuda/precondition_03.with_trilinos=true.with_mpi=true.with_p4est=true.mpirun=2.output new file mode 100644 index 0000000000..3b90478aad --- /dev/null +++ b/tests/cuda/precondition_03.with_trilinos=true.with_mpi=true.with_p4est=true.mpirun=2.output @@ -0,0 +1,10 @@ + +DEAL:0:2d::Testing FE_Q<2>(1) +DEAL:0:2d::Norm of difference: 0 +DEAL:0:2d:: +DEAL:0:3d::Testing FE_Q<3>(1) +DEAL:0:3d::Norm of difference: 0 +DEAL:0:3d:: +DEAL:0:3d::Testing FE_Q<3>(2) +DEAL:0:3d::Norm of difference: 0 +DEAL:0:3d:: -- 2.39.5