From 4077f97fce0aa5ac149f59ce79e94f6f559006d0 Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Fri, 26 May 2023 08:56:17 -0400 Subject: [PATCH] Move more MatrixFree tests --- tests/cuda/matrix_vector_mf.h | 257 ------------------ .../coefficient_eval_device.cc} | 16 +- .../coefficient_eval_device.output} | 0 ...rix_free_device_precondition_chebyshev.cc} | 25 +- ...est=true.with_lapack=true.mpirun=1.output} | 0 ...est=true.with_lapack=true.mpirun=2.output} | 0 6 files changed, 21 insertions(+), 277 deletions(-) delete mode 100644 tests/cuda/matrix_vector_mf.h rename tests/{cuda/coefficient_eval.cc => matrix_free/coefficient_eval_device.cc} (92%) rename tests/{cuda/coefficient_eval.output => matrix_free/coefficient_eval_device.output} (100%) rename tests/{cuda/precondition_03.cc => matrix_free/matrix_free_device_precondition_chebyshev.cc} (94%) rename tests/{cuda/precondition_03.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=1.output => matrix_free/matrix_free_device_precondition_chebyshev.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=1.output} (100%) rename tests/{cuda/precondition_03.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=2.output => matrix_free/matrix_free_device_precondition_chebyshev.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=2.output} (100%) diff --git a/tests/cuda/matrix_vector_mf.h b/tests/cuda/matrix_vector_mf.h deleted file mode 100644 index c7ec9de966..0000000000 --- a/tests/cuda/matrix_vector_mf.h +++ /dev/null @@ -1,257 +0,0 @@ -// --------------------------------------------------------------------- -// -// Copyright (C) 2017 - 2021 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 is a template for matrix-vector products with the Helmholtz equation -// (zero and first derivatives) on different kinds of meshes (Cartesian, -// general, with and without hanging nodes). - -#include -#include - -#include -#include - -#include "../tests.h" - -template -class HelmholtzOperatorQuad -{ -public: - DEAL_II_HOST_DEVICE - HelmholtzOperatorQuad( - const typename CUDAWrappers::MatrixFree::Data *gpu_data, - Number * coef, - int cell) - : gpu_data(gpu_data) - , coef(coef) - , cell(cell) - {} - - DEAL_II_HOST_DEVICE void - operator()( - CUDAWrappers::FEEvaluation - * fe_eval, - int q_point) const; - - static const unsigned int n_q_points = - dealii::Utilities::pow(n_q_points_1d, dim); - -private: - const typename CUDAWrappers::MatrixFree::Data *gpu_data; - Number * coef; - int cell; -}; - - - -template -DEAL_II_HOST_DEVICE void -HelmholtzOperatorQuad::operator()( - CUDAWrappers::FEEvaluation *fe_eval, - int q_point) const -{ - unsigned int pos = gpu_data->local_q_point_id(cell, n_q_points, q_point); - fe_eval->submit_value(coef[pos] * fe_eval->get_value(q_point), q_point); - fe_eval->submit_gradient(fe_eval->get_gradient(q_point), q_point); -} - - - -template -class HelmholtzOperator -{ -public: - static const unsigned int n_dofs_1d = fe_degree + 1; - static const unsigned int n_local_dofs = - dealii::Utilities::pow(fe_degree + 1, dim); - static const unsigned int n_q_points = - dealii::Utilities::pow(n_q_points_1d, dim); - - HelmholtzOperator(Number *coefficient) - : coef(coefficient) - {} - - DEAL_II_HOST_DEVICE void - operator()( - const unsigned int cell, - const typename CUDAWrappers::MatrixFree::Data *gpu_data, - CUDAWrappers::SharedData * shared_data, - const Number * src, - Number * dst) const; - - Number *coef; -}; - - - -template -DEAL_II_HOST_DEVICE void -HelmholtzOperator::operator()( - const unsigned int cell, - const typename CUDAWrappers::MatrixFree::Data *gpu_data, - CUDAWrappers::SharedData * shared_data, - const Number * src, - Number * dst) const -{ - CUDAWrappers::FEEvaluation fe_eval( - gpu_data, shared_data); - fe_eval.read_dof_values(src); - fe_eval.evaluate(true, true); - fe_eval.apply_for_each_quad_point( - HelmholtzOperatorQuad(gpu_data, - coef, - cell)); - fe_eval.integrate(true, true); - fe_eval.distribute_local_to_global(dst); -} - - - -template -class VaryingCoefficientFunctor -{ -public: - VaryingCoefficientFunctor(Number *coefficient) - : coef(coefficient) - {} - - DEAL_II_HOST_DEVICE void - operator()( - const typename CUDAWrappers::MatrixFree::Data *gpu_data, - const unsigned int cell, - const unsigned int q) const; - - static const unsigned int n_dofs_1d = fe_degree + 1; - static const unsigned int n_local_dofs = - dealii::Utilities::pow(fe_degree + 1, dim); - static const unsigned int n_q_points = - dealii::Utilities::pow(n_q_points_1d, dim); - -private: - Number *coef; -}; - - - -template -DEAL_II_HOST_DEVICE void -VaryingCoefficientFunctor::operator()( - const typename CUDAWrappers::MatrixFree::Data *gpu_data, - const unsigned int cell, - const unsigned int q) const -{ - const unsigned int pos = gpu_data->local_q_point_id(cell, n_q_points, q); - const auto q_point = gpu_data->get_quadrature_point(cell, q); - - - - Number p_square = 0.; - for (unsigned int i = 0; i < dim; ++i) - { - Number coord = q_point[i]; - p_square += coord * coord; - } - coef[pos] = 10. / (0.05 + 2. * p_square); -} - - - -template , - int n_q_points_1d = fe_degree + 1> -class MatrixFreeTest : public Subscriptor -{ -public: - MatrixFreeTest(const CUDAWrappers::MatrixFree &data_in, - const unsigned int size, - const bool constant_coeff = true); - - void - 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; - LinearAlgebra::CUDAWrappers::Vector coef; -}; - -template -MatrixFreeTest:: - MatrixFreeTest(const CUDAWrappers::MatrixFree &data_in, - const unsigned int size, - const bool constant_coeff) - : data(data_in) -{ - coef.reinit(size); - if (constant_coeff) - { - coef.add(10.); - } - else - { - VaryingCoefficientFunctor functor( - coef.get_values()); - data.evaluate_coefficients(functor); - } -} - -template -Number -MatrixFreeTest::el( - const unsigned int row, - const unsigned int col) const -{ - (void)col; - Assert(false, ExcNotImplemented()); - return 0.; -} - -template -void -MatrixFreeTest::vmult( - VectorType & dst, - const VectorType &src) const -{ - dst = static_cast(0.); - HelmholtzOperator helmholtz_operator( - coef.get_values()); - data.cell_loop(helmholtz_operator, src, dst); - data.copy_constrained_values(src, dst); -} diff --git a/tests/cuda/coefficient_eval.cc b/tests/matrix_free/coefficient_eval_device.cc similarity index 92% rename from tests/cuda/coefficient_eval.cc rename to tests/matrix_free/coefficient_eval_device.cc index 7469a928a4..9ae646ad34 100644 --- a/tests/cuda/coefficient_eval.cc +++ b/tests/matrix_free/coefficient_eval_device.cc @@ -84,7 +84,8 @@ public: DummyMatrixFree(const CUDAWrappers::MatrixFree &data_in, const unsigned int size); void - eval(LinearAlgebra::CUDAWrappers::Vector &dst) const; + eval(LinearAlgebra::distributed::Vector &dst) + const; private: const CUDAWrappers::MatrixFree &data; @@ -101,10 +102,10 @@ DummyMatrixFree::DummyMatrixFree( template void DummyMatrixFree::eval( - LinearAlgebra::CUDAWrappers::Vector &dst) const + LinearAlgebra::distributed::Vector &dst) const { - LinearAlgebra::CUDAWrappers::Vector src(dst); - DummyOperator dummy_operator; + LinearAlgebra::distributed::Vector src(dst); + DummyOperator dummy_operator; data.cell_loop(dummy_operator, src, dst); } @@ -140,11 +141,12 @@ test() tria.n_active_cells() * n_q_points_per_cell); const unsigned int size = tria.n_active_cells() * n_q_points_per_cell; - LinearAlgebra::ReadWriteVector coef(size); - LinearAlgebra::CUDAWrappers::Vector coef_device(size); + LinearAlgebra::ReadWriteVector coef(size); + LinearAlgebra::distributed::Vector coef_device( + size); mf.eval(coef_device); - cudaDeviceSynchronize(); + Kokkos::fence(); coef.import_elements(coef_device, VectorOperation::insert); // Computation the host diff --git a/tests/cuda/coefficient_eval.output b/tests/matrix_free/coefficient_eval_device.output similarity index 100% rename from tests/cuda/coefficient_eval.output rename to tests/matrix_free/coefficient_eval_device.output diff --git a/tests/cuda/precondition_03.cc b/tests/matrix_free/matrix_free_device_precondition_chebyshev.cc similarity index 94% rename from tests/cuda/precondition_03.cc rename to tests/matrix_free/matrix_free_device_precondition_chebyshev.cc index cf2ce3d863..05455b2571 100644 --- a/tests/cuda/precondition_03.cc +++ b/tests/matrix_free/matrix_free_device_precondition_chebyshev.cc @@ -47,7 +47,7 @@ #include "../tests.h" -#include "matrix_vector_mf.h" +#include "matrix_vector_device_mf.h" @@ -118,15 +118,16 @@ test() const unsigned int coef_size = tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim); - MatrixFreeTest> + MatrixFreeTest< + dim, + fe_degree, + Number, + LinearAlgebra::distributed::Vector> mf(mf_data, coef_size); mf.internal_m = owned_set.size(); - LinearAlgebra::distributed::Vector in_dev( + LinearAlgebra::distributed::Vector in_dev( owned_set, MPI_COMM_WORLD); - LinearAlgebra::distributed::Vector out_dev( + LinearAlgebra::distributed::Vector out_dev( owned_set, MPI_COMM_WORLD); LinearAlgebra::ReadWriteVector rw_in(owned_set); @@ -223,13 +224,13 @@ test() dim, fe_degree, Number, - LinearAlgebra::distributed::Vector>, - LinearAlgebra::distributed::Vector>; + 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( + LinearAlgebra::distributed::Vector>>( + LinearAlgebra::distributed::Vector( ref.get_partitioner())); device_preconditioner_data.preconditioner->get_vector().import_elements( matrix_diagonal, VectorOperation::insert); @@ -262,8 +263,6 @@ main(int argc, char **argv) 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(); diff --git a/tests/cuda/precondition_03.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_precondition_chebyshev.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=1.output similarity index 100% rename from tests/cuda/precondition_03.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=1.output rename to tests/matrix_free/matrix_free_device_precondition_chebyshev.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=1.output diff --git a/tests/cuda/precondition_03.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_precondition_chebyshev.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=2.output similarity index 100% rename from tests/cuda/precondition_03.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=2.output rename to tests/matrix_free/matrix_free_device_precondition_chebyshev.with_trilinos=true.with_p4est=true.with_lapack=true.mpirun=2.output -- 2.39.5