Move more MatrixFree tests

diff --git a/tests/cuda/matrix_vector_mf.h b/tests/cuda/matrix_vector_mf.h
deleted file mode 100644 (file)
index c7ec9de..0000000
--- 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 <deal.II/lac/cuda_vector.h>
-#include <deal.II/lac/la_parallel_vector.h>
-
-#include <deal.II/matrix_free/cuda_fe_evaluation.h>
-#include <deal.II/matrix_free/cuda_matrix_free.h>
-
-#include "../tests.h"
-
-template <int dim, int fe_degree, typename Number, int n_q_points_1d>
-class HelmholtzOperatorQuad
-{
-public:
-  DEAL_II_HOST_DEVICE
-  HelmholtzOperatorQuad(
-    const typename CUDAWrappers::MatrixFree<dim, Number>::Data *gpu_data,
-    Number *                                                    coef,
-    int                                                         cell)
-    : gpu_data(gpu_data)
-    , coef(coef)
-    , cell(cell)
-  {}
-
-  DEAL_II_HOST_DEVICE void
-  operator()(
-    CUDAWrappers::FEEvaluation<dim, fe_degree, n_q_points_1d, 1, Number>
-      * 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<dim, Number>::Data *gpu_data;
-  Number *                                                    coef;
-  int                                                         cell;
-};
-
-
-
-template <int dim, int fe_degree, typename Number, int n_q_points_1d>
-DEAL_II_HOST_DEVICE void
-HelmholtzOperatorQuad<dim, fe_degree, Number, n_q_points_1d>::operator()(
-  CUDAWrappers::FEEvaluation<dim, fe_degree, n_q_points_1d, 1, Number> *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 <int dim, int fe_degree, typename Number, int n_q_points_1d>
-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<dim, Number>::Data *gpu_data,
-    CUDAWrappers::SharedData<dim, Number> *                     shared_data,
-    const Number *                                              src,
-    Number *                                                    dst) const;
-
-  Number *coef;
-};
-
-
-
-template <int dim, int fe_degree, typename Number, int n_q_points_1d>
-DEAL_II_HOST_DEVICE void
-HelmholtzOperator<dim, fe_degree, Number, n_q_points_1d>::operator()(
-  const unsigned int                                          cell,
-  const typename CUDAWrappers::MatrixFree<dim, Number>::Data *gpu_data,
-  CUDAWrappers::SharedData<dim, Number> *                     shared_data,
-  const Number *                                              src,
-  Number *                                                    dst) const
-{
-  CUDAWrappers::FEEvaluation<dim, fe_degree, n_q_points_1d, 1, Number> 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<dim, fe_degree, Number, n_q_points_1d>(gpu_data,
-                                                                 coef,
-                                                                 cell));
-  fe_eval.integrate(true, true);
-  fe_eval.distribute_local_to_global(dst);
-}
-
-
-
-template <int dim, int fe_degree, typename Number, int n_q_points_1d>
-class VaryingCoefficientFunctor
-{
-public:
-  VaryingCoefficientFunctor(Number *coefficient)
-    : coef(coefficient)
-  {}
-
-  DEAL_II_HOST_DEVICE void
-  operator()(
-    const typename CUDAWrappers::MatrixFree<dim, Number>::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 <int dim, int fe_degree, typename Number, int n_q_points_1d>
-DEAL_II_HOST_DEVICE void
-VaryingCoefficientFunctor<dim, fe_degree, Number, n_q_points_1d>::operator()(
-  const typename CUDAWrappers::MatrixFree<dim, Number>::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 dim,
-          int fe_degree,
-          typename Number,
-          typename VectorType = LinearAlgebra::CUDAWrappers::Vector<Number>,
-          int n_q_points_1d   = fe_degree + 1>
-class MatrixFreeTest : public Subscriptor
-{
-public:
-  MatrixFreeTest(const CUDAWrappers::MatrixFree<dim, Number> &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<dim, Number> &data;
-  LinearAlgebra::CUDAWrappers::Vector<Number>  coef;
-};
-
-template <int dim,
-          int fe_degree,
-          typename Number,
-          typename VectorType,
-          int n_q_points_1d>
-MatrixFreeTest<dim, fe_degree, Number, VectorType, n_q_points_1d>::
-  MatrixFreeTest(const CUDAWrappers::MatrixFree<dim, Number> &data_in,
-                 const unsigned int                           size,
-                 const bool                                   constant_coeff)
-  : data(data_in)
-{
-  coef.reinit(size);
-  if (constant_coeff)
-    {
-      coef.add(10.);
-    }
-  else
-    {
-      VaryingCoefficientFunctor<dim, fe_degree, Number, n_q_points_1d> functor(
-        coef.get_values());
-      data.evaluate_coefficients(functor);
-    }
-}
-
-template <int dim,
-          int fe_degree,
-          typename Number,
-          typename VectorType,
-          int n_q_points_1d>
-Number
-MatrixFreeTest<dim, fe_degree, Number, VectorType, n_q_points_1d>::el(
-  const unsigned int row,
-  const unsigned int col) const
-{
-  (void)col;
-  Assert(false, ExcNotImplemented());
-  return 0.;
-}
-
-template <int dim,
-          int fe_degree,
-          typename Number,
-          typename VectorType,
-          int n_q_points_1d>
-void
-MatrixFreeTest<dim, fe_degree, Number, VectorType, n_q_points_1d>::vmult(
-  VectorType &      dst,
-  const VectorType &src) const
-{
-  dst = static_cast<Number>(0.);
-  HelmholtzOperator<dim, fe_degree, Number, n_q_points_1d> 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 7469a928a4e97d92d2e76a766ae25417e0e29a36..9ae646ad34485acd0d28a0e73f27bba512eeb7e1 100644 (file)
--- a/tests/cuda/coefficient_eval.cc
+++ b/tests/matrix_free/coefficient_eval_device.cc
@@ -84,7 +84,8 @@ public:
   DummyMatrixFree(const CUDAWrappers::MatrixFree<dim, double> &data_in,
                   const unsigned int                           size);
   void
-  eval(LinearAlgebra::CUDAWrappers::Vector<double> &dst) const;
+  eval(LinearAlgebra::distributed::Vector<double, MemorySpace::Default> &dst)
+    const;
 
 private:
   const CUDAWrappers::MatrixFree<dim, double> &data;
@@ -101,10 +102,10 @@ DummyMatrixFree<dim, fe_degree>::DummyMatrixFree(
 template <int dim, int fe_degree>
 void
 DummyMatrixFree<dim, fe_degree>::eval(
-  LinearAlgebra::CUDAWrappers::Vector<double> &dst) const
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> &dst) const
 {
-  LinearAlgebra::CUDAWrappers::Vector<double> src(dst);
-  DummyOperator<dim, fe_degree>               dummy_operator;
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> src(dst);
+  DummyOperator<dim, fe_degree> 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<double>      coef(size);
-  LinearAlgebra::CUDAWrappers::Vector<double> coef_device(size);
+  LinearAlgebra::ReadWriteVector<double>                           coef(size);
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> 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 cf2ce3d863d9e3d023e34a2e8e3f3f0cb43a5416..05455b25710c101deee0e10ba52834534508e6f3 100644 (file)
--- 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<dim,
-                 fe_degree,
-                 Number,
-                 LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>
+  MatrixFreeTest<
+    dim,
+    fe_degree,
+    Number,
+    LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>
     mf(mf_data, coef_size);
   mf.internal_m = owned_set.size();
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev(
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev(
     owned_set, MPI_COMM_WORLD);
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev(
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev(
     owned_set, MPI_COMM_WORLD);
 
   LinearAlgebra::ReadWriteVector<Number> rw_in(owned_set);
@@ -223,13 +224,13 @@ test()
       dim,
       fe_degree,
       Number,
-      LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>,
-    LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>;
+      LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>,
+    LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>;
   DevicePreconditionerType precondition_chebyshev_device;
   typename DevicePreconditionerType::AdditionalData device_preconditioner_data;
   device_preconditioner_data.preconditioner = std::make_shared<DiagonalMatrix<
-    LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>>(
-    LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>(
+    LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>>(
+    LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>(
       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