From: Daniel Arndt <arndtd@ornl.gov>
Date: Wed, 24 May 2023 02:13:23 +0000 (-0400)
Subject: Test CUDAMatrixFree not only for DEAL_II_WITH_CUDA
X-Git-Tag: v9.5.0-rc1~177^2~10
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=138ae122dba4c5052edb5b2730f954b7b748f54c;p=dealii.git

Test CUDAMatrixFree not only for DEAL_II_WITH_CUDA
---

diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h
index b106f850b4..ae1b6ed9ff 100644
--- a/include/deal.II/matrix_free/cuda_fe_evaluation.h
+++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h
@@ -18,21 +18,19 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_WITH_CUDA
+#include <deal.II/base/memory_space.h>
+#include <deal.II/base/tensor.h>
+#include <deal.II/base/utilities.h>
 
-#  include <deal.II/base/memory_space.h>
-#  include <deal.II/base/tensor.h>
-#  include <deal.II/base/utilities.h>
+#include <deal.II/lac/cuda_atomic.h>
+#include <deal.II/lac/cuda_vector.h>
 
-#  include <deal.II/lac/cuda_atomic.h>
-#  include <deal.II/lac/cuda_vector.h>
+#include <deal.II/matrix_free/cuda_hanging_nodes_internal.h>
+#include <deal.II/matrix_free/cuda_matrix_free.h>
+#include <deal.II/matrix_free/cuda_matrix_free.templates.h>
+#include <deal.II/matrix_free/cuda_tensor_product_kernels.h>
 
-#  include <deal.II/matrix_free/cuda_hanging_nodes_internal.h>
-#  include <deal.II/matrix_free/cuda_matrix_free.h>
-#  include <deal.II/matrix_free/cuda_matrix_free.templates.h>
-#  include <deal.II/matrix_free/cuda_tensor_product_kernels.h>
-
-#  include <Kokkos_Core.hpp>
+#include <Kokkos_Core.hpp>
 
 DEAL_II_NAMESPACE_OPEN
 
@@ -526,5 +524,3 @@ namespace CUDAWrappers
 DEAL_II_NAMESPACE_CLOSE
 
 #endif
-
-#endif
diff --git a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
index 90942f92fe..743b06f015 100644
--- a/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
+++ b/include/deal.II/matrix_free/cuda_hanging_nodes_internal.h
@@ -18,13 +18,11 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_WITH_CUDA
+#include <deal.II/base/cuda_size.h>
 
-#  include <deal.II/base/cuda_size.h>
+#include <deal.II/matrix_free/hanging_nodes_internal.h>
 
-#  include <deal.II/matrix_free/hanging_nodes_internal.h>
-
-#  include <Kokkos_Macros.hpp>
+#include <Kokkos_Macros.hpp>
 
 DEAL_II_NAMESPACE_OPEN
 namespace CUDAWrappers
@@ -464,5 +462,3 @@ namespace CUDAWrappers
 
 DEAL_II_NAMESPACE_CLOSE
 #endif
-
-#endif
diff --git a/include/deal.II/matrix_free/cuda_matrix_free.h b/include/deal.II/matrix_free/cuda_matrix_free.h
index 3666d7e853..a30d57b260 100644
--- a/include/deal.II/matrix_free/cuda_matrix_free.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.h
@@ -19,28 +19,26 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_WITH_CUDA
-
-#  include <deal.II/base/cuda_size.h>
-#  include <deal.II/base/memory_space.h>
-#  include <deal.II/base/mpi_stub.h>
-#  include <deal.II/base/partitioner.h>
-#  include <deal.II/base/quadrature.h>
-#  include <deal.II/base/tensor.h>
-#  include <deal.II/base/utilities.h>
+#include <deal.II/base/cuda_size.h>
+#include <deal.II/base/memory_space.h>
+#include <deal.II/base/mpi_stub.h>
+#include <deal.II/base/partitioner.h>
+#include <deal.II/base/quadrature.h>
+#include <deal.II/base/tensor.h>
+#include <deal.II/base/utilities.h>
 
-#  include <deal.II/dofs/dof_handler.h>
+#include <deal.II/dofs/dof_handler.h>
 
-#  include <deal.II/fe/fe_update_flags.h>
-#  include <deal.II/fe/mapping.h>
+#include <deal.II/fe/fe_update_flags.h>
+#include <deal.II/fe/mapping.h>
 
-#  include <deal.II/grid/filtered_iterator.h>
+#include <deal.II/grid/filtered_iterator.h>
 
-#  include <deal.II/lac/affine_constraints.h>
-#  include <deal.II/lac/cuda_vector.h>
-#  include <deal.II/lac/la_parallel_vector.h>
+#include <deal.II/lac/affine_constraints.h>
+#include <deal.II/lac/cuda_vector.h>
+#include <deal.II/lac/la_parallel_vector.h>
 
-#  include <Kokkos_Core.hpp>
+#include <Kokkos_Core.hpp>
 
 
 
@@ -58,13 +56,13 @@ namespace internal
 namespace CUDAWrappers
 {
   // forward declaration
-#  ifndef DOXYGEN
+#ifndef DOXYGEN
   namespace internal
   {
     template <int dim, typename Number>
     class ReinitHelper;
   }
-#  endif
+#endif
 
   /**
    * This class collects all the data that is stored for the matrix free
@@ -118,12 +116,12 @@ namespace CUDAWrappers
         , use_coloring(use_coloring)
         , overlap_communication_computation(overlap_communication_computation)
       {
-#  ifndef DEAL_II_MPI_WITH_DEVICE_SUPPORT
+#ifndef DEAL_II_MPI_WITH_DEVICE_SUPPORT
         AssertThrow(
           overlap_communication_computation == false,
           ExcMessage(
             "Overlapping communication and computation requires CUDA-aware MPI."));
-#  endif
+#endif
         if (overlap_communication_computation == true)
           AssertThrow(
             use_coloring == false || overlap_communication_computation == false,
@@ -377,6 +375,7 @@ namespace CUDAWrappers
     void
     set_constrained_values(const Number val, VectorType &dst) const;
 
+#ifdef DEAL_II_WITH_CUDA
     /**
      * Initialize a serial vector. The size corresponds to the number of degrees
      * of freedom in the DoFHandler object.
@@ -384,6 +383,7 @@ namespace CUDAWrappers
     void
     initialize_dof_vector(
       LinearAlgebra::CUDAWrappers::Vector<Number> &vec) const;
+#endif
 
     /**
      * Initialize a distributed vector. The local elements correspond to the
@@ -392,7 +392,8 @@ namespace CUDAWrappers
      */
     void
     initialize_dof_vector(
-      LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &vec) const;
+      LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &vec)
+      const;
 
     /**
      * Return the colored graph of locally owned active cells.
@@ -456,10 +457,13 @@ namespace CUDAWrappers
     template <typename Functor>
     void
     distributed_cell_loop(
-      const Functor &                                                      func,
-      const LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &src,
-      LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &dst) const;
+      const Functor &func,
+      const LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>
+        &                                                               src,
+      LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &dst)
+      const;
 
+#ifdef DEAL_II_WITH_CUDA
     /**
      * This function should never be called. Calling it results in an internal
      * error. This function exists only because cell_loop needs
@@ -471,6 +475,7 @@ namespace CUDAWrappers
       const Functor &                                    func,
       const LinearAlgebra::CUDAWrappers::Vector<Number> &src,
       LinearAlgebra::CUDAWrappers::Vector<Number> &      dst) const;
+#endif
 
     /**
      * Unique ID associated with the object.
@@ -591,18 +596,6 @@ namespace CUDAWrappers
     Kokkos::View<Number *, MemorySpace::Default::kokkos_space>
       constraint_weights;
 
-    /**
-     * Grid dimensions associated to the different colors. The grid dimensions
-     * are used to launch the CUDA kernels.
-     */
-    std::vector<dim3> grid_dim;
-
-    /**
-     * Block dimensions associated to the different colors. The block dimensions
-     * are used to launch the CUDA kernels.
-     */
-    std::vector<dim3> block_dim;
-
     /**
      * Shared pointer to a Partitioner for distributed Vectors used in
      * cell_loop. When MPI is not used the pointer is null.
@@ -788,8 +781,6 @@ namespace CUDAWrappers
     data_host.row_start      = data.row_start;
     data_host.use_coloring   = data.use_coloring;
 
-    const unsigned int n_elements =
-      data_host.n_cells * data_host.padding_length;
     if (update_flags & update_quadrature_points)
       {
         data_host.q_points = Kokkos::create_mirror(data.q_points);
@@ -820,7 +811,7 @@ namespace CUDAWrappers
 
   /*----------------------- Inline functions ---------------------------------*/
 
-#  ifndef DOXYGEN
+#ifndef DOXYGEN
 
   template <int dim, typename Number>
   inline const std::vector<std::vector<
@@ -850,11 +841,10 @@ namespace CUDAWrappers
     return *dof_handler;
   }
 
-#  endif
+#endif
 
 } // namespace CUDAWrappers
 
 DEAL_II_NAMESPACE_CLOSE
 
 #endif
-#endif
diff --git a/include/deal.II/matrix_free/cuda_matrix_free.templates.h b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
index 6884cf5681..981e525d44 100644
--- a/include/deal.II/matrix_free/cuda_matrix_free.templates.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
@@ -19,29 +19,29 @@
 
 #include <deal.II/base/config.h>
 
-#ifdef DEAL_II_WITH_CUDA
-
-#  include <deal.II/base/cuda.h>
-#  include <deal.II/base/cuda_size.h>
-#  include <deal.II/base/graph_coloring.h>
-#  include <deal.II/base/memory_space.h>
+#include <deal.II/base/cuda.h>
+#include <deal.II/base/cuda_size.h>
+#include <deal.II/base/graph_coloring.h>
+#include <deal.II/base/memory_space.h>
 
-#  include <deal.II/dofs/dof_tools.h>
+#include <deal.II/dofs/dof_tools.h>
 
-#  include <deal.II/fe/fe_dgq.h>
-#  include <deal.II/fe/fe_values.h>
-#  include <deal.II/fe/mapping_q1.h>
+#include <deal.II/fe/fe_dgq.h>
+#include <deal.II/fe/fe_values.h>
+#include <deal.II/fe/mapping_q1.h>
 
-#  include <deal.II/matrix_free/cuda_hanging_nodes_internal.h>
-#  include <deal.II/matrix_free/cuda_matrix_free.h>
-#  include <deal.II/matrix_free/shape_info.h>
+#include <deal.II/matrix_free/cuda_hanging_nodes_internal.h>
+#include <deal.II/matrix_free/cuda_matrix_free.h>
+#include <deal.II/matrix_free/shape_info.h>
 
-#  include <Kokkos_Core.hpp>
+#include <Kokkos_Core.hpp>
+#ifdef DEAL_II_WITH_CUDA
 #  include <cuda_runtime_api.h>
+#endif
 
-#  include <cmath>
-#  include <functional>
-#  include <string>
+#include <cmath>
+#include <functional>
+#include <string>
 
 
 DEAL_II_NAMESPACE_OPEN
@@ -132,8 +132,6 @@ namespace CUDAWrappers
       // We need at least three colors when we are using CUDA-aware MPI and
       // overlapping the communication
       data->n_cells.resize(std::max(n_colors, 3U), 0);
-      data->grid_dim.resize(n_colors);
-      data->block_dim.resize(n_colors);
       data->local_to_global.resize(n_colors);
       data->constraint_mask.resize(n_colors);
 
@@ -309,6 +307,7 @@ namespace CUDAWrappers
       }
     };
 
+#ifdef DEAL_II_WITH_CUDA
     template <>
     struct VectorLocalSize<LinearAlgebra::CUDAWrappers::Vector<double>>
     {
@@ -328,6 +327,7 @@ namespace CUDAWrappers
         return vec.size();
       }
     };
+#endif
 
 
 
@@ -557,7 +557,7 @@ namespace CUDAWrappers
   }
 
 
-
+#ifdef DEAL_II_WITH_CUDA
   template <int dim, typename Number>
   void
   MatrixFree<dim, Number>::initialize_dof_vector(
@@ -565,13 +565,14 @@ namespace CUDAWrappers
   {
     vec.reinit(n_dofs);
   }
+#endif
 
 
 
   template <int dim, typename Number>
   void
   MatrixFree<dim, Number>::initialize_dof_vector(
-    LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &vec) const
+    LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &vec) const
   {
     if (partitioner)
       vec.reinit(partitioner);
@@ -635,7 +636,6 @@ namespace CUDAWrappers
     // First compute the size of n_cells, row_starts, kernel launch parameters,
     // and constrained_dofs
     std::size_t bytes = n_cells.size() * sizeof(unsigned int) * 2 +
-                        2 * n_colors * sizeof(dim3) +
                         n_constrained_dofs * sizeof(unsigned int);
 
     // For each color, add local_to_global, inv_jacobian, JxW, and q_points.
@@ -668,9 +668,6 @@ namespace CUDAWrappers
   {
     dof_handler = &dof_handler_;
 
-    if (typeid(Number) == typeid(double))
-      cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
-
     UpdateFlags update_flags = additional_data.mapping_update_flags;
     if (update_flags & update_gradients)
       update_flags |= update_JxW_values;
@@ -937,9 +934,9 @@ namespace CUDAWrappers
   template <typename Functor>
   void
   MatrixFree<dim, Number>::distributed_cell_loop(
-    const Functor &                                                      func,
-    const LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &src,
-    LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &dst) const
+    const Functor &func,
+    const LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &src,
+    LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &dst) const
   {
     MemorySpace::Default::kokkos_space::execution_space exec;
 
@@ -1046,12 +1043,13 @@ namespace CUDAWrappers
     else
       {
         // Create the ghosted source and the ghosted destination
-        LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>
+        LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>
           ghosted_src(partitioner);
-        LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>
+        LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>
           ghosted_dst(ghosted_src);
         ghosted_src = src;
         ghosted_dst = dst;
+        ghosted_dst.zero_out_ghost_values();
 
         // Execute the loop on the cells
         for (unsigned int i = 0; i < n_colors; ++i)
@@ -1062,7 +1060,10 @@ namespace CUDAWrappers
                 team_policy(exec, n_cells[i], Kokkos::AUTO);
 
               internal::ApplyKernel<dim, Number, Functor> apply_kernel(
-                func, get_data(i), src.get_values(), dst.get_values());
+                func,
+                get_data(i),
+                ghosted_src.get_values(),
+                ghosted_dst.get_values());
 
               Kokkos::parallel_for(
                 "dealii::MatrixFree::distributed_cell_loop_" +
@@ -1077,8 +1078,7 @@ namespace CUDAWrappers
       }
   }
 
-
-
+#ifdef DEAL_II_WITH_CUDA
   template <int dim, typename Number>
   template <typename Functor>
   void
@@ -1089,10 +1089,9 @@ namespace CUDAWrappers
   {
     Assert(false, ExcInternalError());
   }
+#endif
 } // namespace CUDAWrappers
 
 DEAL_II_NAMESPACE_CLOSE
 
 #endif
-
-#endif
diff --git a/include/deal.II/matrix_free/cuda_tensor_product_kernels.h b/include/deal.II/matrix_free/cuda_tensor_product_kernels.h
index b191a0cf5b..ae39a56f51 100644
--- a/include/deal.II/matrix_free/cuda_tensor_product_kernels.h
+++ b/include/deal.II/matrix_free/cuda_tensor_product_kernels.h
@@ -23,8 +23,6 @@
 
 #include <deal.II/matrix_free/cuda_matrix_free.templates.h>
 
-#ifdef DEAL_II_WITH_CUDA
-
 DEAL_II_NAMESPACE_OPEN
 
 
@@ -756,5 +754,3 @@ namespace CUDAWrappers
 DEAL_II_NAMESPACE_CLOSE
 
 #endif
-
-#endif
diff --git a/source/matrix_free/CMakeLists.txt b/source/matrix_free/CMakeLists.txt
index 995f85a391..19212afd94 100644
--- a/source/matrix_free/CMakeLists.txt
+++ b/source/matrix_free/CMakeLists.txt
@@ -14,6 +14,7 @@
 ## ---------------------------------------------------------------------
 
 set(_src
+  cuda_matrix_free.cc
   dof_info.cc
   evaluation_template_factory.cc
   evaluation_template_factory_inst2.cc
@@ -52,15 +53,5 @@ file(GLOB _header
   ${CMAKE_SOURCE_DIR}/include/deal.II/matrix_free/*.h
   )
 
-if(DEAL_II_WITH_CUDA)
-  set(_src
-    cuda_matrix_free.cc
-    ${_src}
-    )
-  set(_inst
-    ${_inst}
-    )
-endif()
-
 define_object_library(object_matrix_free OBJECT ${_src} ${_header} ${_inst})
 expand_instantiations(object_matrix_free "${_inst}")
diff --git a/source/matrix_free/cuda_matrix_free.cc b/source/matrix_free/cuda_matrix_free.cc
index 2ff0a5968f..22f89f3b11 100644
--- a/source/matrix_free/cuda_matrix_free.cc
+++ b/source/matrix_free/cuda_matrix_free.cc
@@ -15,8 +15,6 @@
 
 #include <deal.II/matrix_free/cuda_matrix_free.templates.h>
 
-#ifdef DEAL_II_WITH_CUDA
-
 DEAL_II_NAMESPACE_OPEN
 
 
@@ -31,5 +29,3 @@ namespace CUDAWrappers
 } // namespace CUDAWrappers
 
 DEAL_II_NAMESPACE_CLOSE
-
-#endif
diff --git a/tests/cuda/matrix_free_initialize_vector.cc b/tests/matrix_free/matrix_free_device_initialize_vector.cc
similarity index 97%
rename from tests/cuda/matrix_free_initialize_vector.cc
rename to tests/matrix_free/matrix_free_device_initialize_vector.cc
index 7cdefebe5f..0c53f24828 100644
--- a/tests/cuda/matrix_free_initialize_vector.cc
+++ b/tests/matrix_free/matrix_free_device_initialize_vector.cc
@@ -50,6 +50,7 @@ check(const LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>
          ExcInternalError());
 }
 
+#ifdef DEAL_II_WITH_CUDA
 template <typename Number>
 void
 check(const LinearAlgebra::CUDAWrappers::Vector<Number> &vector,
@@ -57,6 +58,7 @@ check(const LinearAlgebra::CUDAWrappers::Vector<Number> &vector,
 {
   AssertDimension(vector.size(), reference_partitioner.size());
 }
+#endif
 
 template <int dim, int fe_degree, typename VectorType>
 void
@@ -113,14 +115,16 @@ 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);
+  // init_cuda(true);
   MPILogInitAll mpi_inilog;
 
   test<2,
        1,
        LinearAlgebra::distributed::Vector<double, MemorySpace::Default>>();
+
+#ifdef DEAL_II_WITH_CUDA
   if (Utilities::MPI::n_mpi_processes(MPI_COMM_WORLD) == 1)
     test<2, 1, LinearAlgebra::CUDAWrappers::Vector<double>>();
-
+#endif
   return 0;
 }
diff --git a/tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output
similarity index 100%
rename from tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output
rename to tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.mpirun=2.output
diff --git a/tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=1.output b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=off.mpirun=1.output
similarity index 50%
rename from tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=1.output
rename to tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=off.mpirun=1.output
index a0d4178229..274361dac4 100644
--- a/tests/cuda/matrix_free_initialize_vector.with_mpi=on.with_p4est=on.mpirun=1.output
+++ b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=off.mpirun=1.output
@@ -4,8 +4,3 @@ DEAL:0:0::locally owned dofs :
 DEAL:0:0::locally relevant dofs :
 {[0,24]}
 DEAL:0:0::OK
-DEAL:0:0::locally owned dofs :
-{[0,24]}
-DEAL:0:0::locally relevant dofs :
-{[0,24]}
-DEAL:0:0::OK
diff --git a/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=on.mpirun=1.output b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=on.mpirun=1.output
new file mode 100644
index 0000000000..7e4be316a7
--- /dev/null
+++ b/tests/matrix_free/matrix_free_device_initialize_vector.with_mpi=on.with_p4est=on.with_cuda=on.mpirun=1.output
@@ -0,0 +1,13 @@
+
+
+-
+-DEAL:0:0::locally owned dofs :
+-{[0,24]}
+-DEAL:0:0::locally relevant dofs :
+-{[0,24]}
+-DEAL:0:0::OK
+-DEAL:0:0::locally owned dofs :
+-{[0,24]}
+-DEAL:0:0::locally relevant dofs :
+-{[0,24]}
+-DEAL:0:0::OK
diff --git a/tests/cuda/matrix_free_matrix_vector_01.cc b/tests/matrix_free/matrix_free_device_matrix_vector_01.cc
similarity index 91%
rename from tests/cuda/matrix_free_matrix_vector_01.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_01.cc
index c6a79083fc..e626ae812c 100644
--- a/tests/cuda/matrix_free_matrix_vector_01.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_01.cc
@@ -22,7 +22,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 template <int dim, int fe_degree, typename Number>
 void
@@ -41,6 +41,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::CUDAWrappers::Vector<Number>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells());
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_01.output b/tests/matrix_free/matrix_free_device_matrix_vector_01.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_01.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_01.output
diff --git a/tests/cuda/matrix_free_matrix_vector_02.cc b/tests/matrix_free/matrix_free_device_matrix_vector_02.cc
similarity index 93%
rename from tests/cuda/matrix_free_matrix_vector_02.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_02.cc
index ceaa73a31a..f50cb6c522 100644
--- a/tests/cuda/matrix_free_matrix_vector_02.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_02.cc
@@ -22,7 +22,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 template <int dim, int fe_degree, typename Number>
 void
@@ -45,6 +45,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::CUDAWrappers::Vector<Number>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells());
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_02.output b/tests/matrix_free/matrix_free_device_matrix_vector_02.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_02.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_02.output
diff --git a/tests/cuda/matrix_free_matrix_vector_03.cc b/tests/matrix_free/matrix_free_device_matrix_vector_03.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_03.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_03.cc
index cd581a99c5..2894340c29 100644
--- a/tests/cuda/matrix_free_matrix_vector_03.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_03.cc
@@ -24,7 +24,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 template <int dim, int fe_degree, typename Number>
 void
@@ -66,6 +66,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::CUDAWrappers::Vector<Number>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells());
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_03.output b/tests/matrix_free/matrix_free_device_matrix_vector_03.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_03.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_03.output
diff --git a/tests/cuda/matrix_free_matrix_vector_03b.cc b/tests/matrix_free/matrix_free_device_matrix_vector_03b.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_03b.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_03b.cc
index 464aadc676..38e924ea14 100644
--- a/tests/cuda/matrix_free_matrix_vector_03b.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_03b.cc
@@ -25,7 +25,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 template <int dim, int fe_degree, typename Number>
 void
@@ -68,6 +68,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::CUDAWrappers::Vector<Number>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells(), true, true);
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_03b.output b/tests/matrix_free/matrix_free_device_matrix_vector_03b.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_03b.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_03b.output
diff --git a/tests/cuda/matrix_free_matrix_vector_04.cc b/tests/matrix_free/matrix_free_device_matrix_vector_04.cc
similarity index 92%
rename from tests/cuda/matrix_free_matrix_vector_04.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_04.cc
index c27adc0907..f68d81e5fb 100644
--- a/tests/cuda/matrix_free_matrix_vector_04.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_04.cc
@@ -21,7 +21,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 
 template <int dim, int fe_degree, typename Number>
@@ -45,6 +45,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::CUDAWrappers::Vector<Number>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells());
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_04.output b/tests/matrix_free/matrix_free_device_matrix_vector_04.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_04.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_04.output
diff --git a/tests/cuda/matrix_free_matrix_vector_05.cc b/tests/matrix_free/matrix_free_device_matrix_vector_05.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_05.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_05.cc
index cffa01115f..b8daf713cc 100644
--- a/tests/cuda/matrix_free_matrix_vector_05.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_05.cc
@@ -22,7 +22,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 
 template <int dim, int fe_degree, typename Number>
@@ -77,6 +77,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::CUDAWrappers::Vector<Number>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells());
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_05.output b/tests/matrix_free/matrix_free_device_matrix_vector_05.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_05.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_05.output
diff --git a/tests/cuda/matrix_free_matrix_vector_06.cc b/tests/matrix_free/matrix_free_device_matrix_vector_06.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_06.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_06.cc
index 27fdf46f73..14776e9dc8 100644
--- a/tests/cuda/matrix_free_matrix_vector_06.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_06.cc
@@ -26,7 +26,7 @@
 #include "../tests.h"
 
 #include "create_mesh.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 template <int dim, int fe_degree, typename Number>
 void
@@ -75,6 +75,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::CUDAWrappers::Vector<Number>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells());
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_06.output b/tests/matrix_free/matrix_free_device_matrix_vector_06.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_06.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_06.output
diff --git a/tests/cuda/matrix_free_matrix_vector_06a.cc b/tests/matrix_free/matrix_free_device_matrix_vector_06a.cc
similarity index 97%
rename from tests/cuda/matrix_free_matrix_vector_06a.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_06a.cc
index e20b65caaa..41ea3e611b 100644
--- a/tests/cuda/matrix_free_matrix_vector_06a.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_06a.cc
@@ -22,7 +22,7 @@
 #include "../tests.h"
 
 #include "create_mesh.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 template <int dim, int fe_degree, typename Number>
 void
@@ -71,6 +71,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells());
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_06a.output b/tests/matrix_free/matrix_free_device_matrix_vector_06a.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_06a.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_06a.output
diff --git a/tests/cuda/matrix_free_matrix_vector_06b.cc b/tests/matrix_free/matrix_free_device_matrix_vector_06b.cc
similarity index 97%
rename from tests/cuda/matrix_free_matrix_vector_06b.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_06b.cc
index 473297c9b3..5d7d72718a 100644
--- a/tests/cuda/matrix_free_matrix_vector_06b.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_06b.cc
@@ -21,7 +21,7 @@
 #include "../tests.h"
 
 #include "create_mesh.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 template <int dim, int fe_degree, typename Number>
 void
@@ -70,6 +70,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells(), false);
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_06b.output b/tests/matrix_free/matrix_free_device_matrix_vector_06b.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_06b.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_06b.output
diff --git a/tests/cuda/matrix_free_matrix_vector_09.cc b/tests/matrix_free/matrix_free_device_matrix_vector_09.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_09.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_09.cc
index 13244c274d..833319faef 100644
--- a/tests/cuda/matrix_free_matrix_vector_09.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_09.cc
@@ -24,7 +24,7 @@
 #include "../tests.h"
 
 #include "create_mesh.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
 
 template <int dim, int fe_degree, typename Number>
 void
@@ -73,6 +73,6 @@ test()
   do_test<dim,
           fe_degree,
           Number,
-          LinearAlgebra::CUDAWrappers::Vector<Number>,
+          LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
           fe_degree + 1>(dof, constraints, tria.n_active_cells());
 }
diff --git a/tests/cuda/matrix_free_matrix_vector_09.output b/tests/matrix_free/matrix_free_device_matrix_vector_09.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_09.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_09.output
diff --git a/tests/cuda/matrix_free_matrix_vector_10.cc b/tests/matrix_free/matrix_free_device_matrix_vector_10.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_10.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_10.cc
index 6ea858759a..92b50ded61 100644
--- a/tests/cuda/matrix_free_matrix_vector_10.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_10.cc
@@ -46,7 +46,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
 
 
 
@@ -117,13 +117,14 @@ 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);
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev;
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev;
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev;
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev;
   mf_data.initialize_dof_vector(in_dev);
   mf_data.initialize_dof_vector(out_dev);
 
@@ -224,8 +225,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/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=1.output
diff --git a/tests/cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_10.with_trilinos=true.with_p4est=true.mpirun=2.output
diff --git a/tests/cuda/matrix_free_matrix_vector_10a.cc b/tests/matrix_free/matrix_free_device_matrix_vector_10a.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_10a.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_10a.cc
index 7c434dd7bb..513c8bdb82 100644
--- a/tests/cuda/matrix_free_matrix_vector_10a.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_10a.cc
@@ -41,7 +41,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
 
 
 template <int dim>
@@ -119,14 +119,15 @@ 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, false);
 
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev;
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev;
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev;
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev;
   mf_data.initialize_dof_vector(in_dev);
   mf_data.initialize_dof_vector(out_dev);
   LinearAlgebra::ReadWriteVector<Number> rw_in(owned_set);
@@ -230,8 +231,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/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=1.output
diff --git a/tests/cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_10a.with_trilinos=true.with_p4est=true.mpirun=2.output
diff --git a/tests/cuda/matrix_free_matrix_vector_19.cc b/tests/matrix_free/matrix_free_device_matrix_vector_19.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_19.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_19.cc
index 39fba64516..f71c3c5f33 100644
--- a/tests/cuda/matrix_free_matrix_vector_19.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_19.cc
@@ -42,7 +42,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
 
 
 
@@ -113,14 +113,15 @@ 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);
-  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);
@@ -220,8 +221,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/matrix_free_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_19.with_trilinos=true.with_trilinos_with_zoltan=true.mpirun=2.output
diff --git a/tests/cuda/matrix_free_matrix_vector_25.cc b/tests/matrix_free/matrix_free_device_matrix_vector_25.cc
similarity index 95%
rename from tests/cuda/matrix_free_matrix_vector_25.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_25.cc
index a80c6ac3b4..dcb5464c43 100644
--- a/tests/cuda/matrix_free_matrix_vector_25.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_25.cc
@@ -40,7 +40,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
 
 
 
@@ -120,13 +120,14 @@ 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);
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev;
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev;
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev;
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev;
   mf_data.initialize_dof_vector(in_dev);
   mf_data.initialize_dof_vector(out_dev);
 
@@ -227,8 +228,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/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=1.output
diff --git a/tests/cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_25.with_trilinos=true.with_p4est=true.mpirun=2.output
diff --git a/tests/cuda/matrix_free_matrix_vector_26.cc b/tests/matrix_free/matrix_free_device_matrix_vector_26.cc
similarity index 96%
rename from tests/cuda/matrix_free_matrix_vector_26.cc
rename to tests/matrix_free/matrix_free_device_matrix_vector_26.cc
index 8c1eeaf803..32b73b449b 100644
--- a/tests/cuda/matrix_free_matrix_vector_26.cc
+++ b/tests/matrix_free/matrix_free_device_matrix_vector_26.cc
@@ -42,7 +42,7 @@
 
 #include "../tests.h"
 
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
 
 
 
@@ -126,13 +126,14 @@ 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);
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev;
-  LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev;
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev;
+  LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev;
   mf_data.initialize_dof_vector(in_dev);
   mf_data.initialize_dof_vector(out_dev);
 
@@ -233,8 +234,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/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output b/tests/matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=1.output
diff --git a/tests/cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output b/tests/matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output
similarity index 100%
rename from tests/cuda/matrix_free_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output
rename to tests/matrix_free/matrix_free_device_matrix_vector_26.with_trilinos=true.with_p4est=true.mpirun=2.output
diff --git a/tests/cuda/matrix_free_multiple_objects.cc b/tests/matrix_free/matrix_free_device_multiple_objects.cc
similarity index 87%
rename from tests/cuda/matrix_free_multiple_objects.cc
rename to tests/matrix_free/matrix_free_device_multiple_objects.cc
index e1bf9da3b2..04d44ba029 100644
--- a/tests/cuda/matrix_free_multiple_objects.cc
+++ b/tests/matrix_free/matrix_free_device_multiple_objects.cc
@@ -44,24 +44,27 @@
 #include "../tests.h"
 
 #include "Kokkos_Core.hpp"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
 
 template <int fe_degree, int n_q_points_1d>
 void
-do_test(const DoFHandler<2> &            dof,
-        MatrixFreeTest<2,
-                       fe_degree,
-                       double,
-                       LinearAlgebra::CUDAWrappers::Vector<double>,
-                       n_q_points_1d> &  mf,
+do_test(const DoFHandler<2> &dof,
+        MatrixFreeTest<
+          2,
+          fe_degree,
+          double,
+          LinearAlgebra::distributed::Vector<double, MemorySpace::Default>,
+          n_q_points_1d> &               mf,
         unsigned int                     n_dofs,
         MappingQ<2> &                    mapping,
         const AffineConstraints<double> &constraints)
 {
-  Vector<double>                              in_host(n_dofs), out_host(n_dofs);
-  LinearAlgebra::ReadWriteVector<double>      in(n_dofs), out(n_dofs);
-  LinearAlgebra::CUDAWrappers::Vector<double> in_device(n_dofs);
-  LinearAlgebra::CUDAWrappers::Vector<double> out_device(n_dofs);
+  Vector<double>                         in_host(n_dofs), out_host(n_dofs);
+  LinearAlgebra::ReadWriteVector<double> in(n_dofs), out(n_dofs);
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> in_device(
+    n_dofs);
+  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> out_device(
+    n_dofs);
 
   for (unsigned int i = 0; i < n_dofs; ++i)
     {
@@ -73,7 +76,7 @@ do_test(const DoFHandler<2> &            dof,
 
   in_device.import_elements(in, VectorOperation::insert);
   mf.vmult(out_device, in_device);
-  cudaDeviceSynchronize();
+  Kokkos::fence();
   out.import_elements(out_device, VectorOperation::insert);
 
   // assemble sparse matrix with (\nabla v, \nabla u) + (v, 10 * u)
@@ -149,7 +152,6 @@ main()
   deallog << std::setprecision(3);
 
   Kokkos::initialize();
-  init_cuda();
 
   {
     Triangulation<2> tria;
@@ -174,11 +176,12 @@ main()
     const QGauss<1> quad_1(n_q_points_1d_1);
     mf_data_1.reinit(mapping_1, dof_1, constraints, quad_1, additional_data_1);
     const unsigned int n_dofs_1 = dof_1.n_dofs();
-    MatrixFreeTest<2,
-                   fe_degree_1,
-                   double,
-                   LinearAlgebra::CUDAWrappers::Vector<double>,
-                   n_q_points_1d_1>
+    MatrixFreeTest<
+      2,
+      fe_degree_1,
+      double,
+      LinearAlgebra::distributed::Vector<double, MemorySpace::Default>,
+      n_q_points_1d_1>
       mf_1(mf_data_1,
            n_dofs_1 * std::pow(n_q_points_1d_1, 2),
            constant_coefficient);
@@ -198,11 +201,12 @@ main()
     const QGauss<1> quad_2(n_q_points_1d_2);
     mf_data_2.reinit(mapping_2, dof_2, constraints, quad_2, additional_data_2);
     const unsigned int n_dofs_2 = dof_2.n_dofs();
-    MatrixFreeTest<2,
-                   fe_degree_2,
-                   double,
-                   LinearAlgebra::CUDAWrappers::Vector<double>,
-                   n_q_points_1d_2>
+    MatrixFreeTest<
+      2,
+      fe_degree_2,
+      double,
+      LinearAlgebra::distributed::Vector<double, MemorySpace::Default>,
+      n_q_points_1d_2>
       mf_2(mf_data_2,
            n_dofs_2 * std::pow(n_q_points_1d_2, 2),
            constant_coefficient);
diff --git a/tests/cuda/matrix_free_multiple_objects.output b/tests/matrix_free/matrix_free_device_multiple_objects.output
similarity index 100%
rename from tests/cuda/matrix_free_multiple_objects.output
rename to tests/matrix_free/matrix_free_device_multiple_objects.output
diff --git a/tests/cuda/matrix_free_no_index_initialize.cc b/tests/matrix_free/matrix_free_device_no_index_initialize.cc
similarity index 92%
rename from tests/cuda/matrix_free_no_index_initialize.cc
rename to tests/matrix_free/matrix_free_device_no_index_initialize.cc
index bda055ddc6..9ab2222c4f 100644
--- a/tests/cuda/matrix_free_no_index_initialize.cc
+++ b/tests/matrix_free/matrix_free_device_no_index_initialize.cc
@@ -83,15 +83,12 @@ public:
   void
   test() const
   {
-    LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> dst_dummy;
-    LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> src_dummy;
+    LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> dst_dummy;
+    LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> src_dummy;
 
     data.cell_loop(*this, src_dummy, dst_dummy);
 
-    // Check that the kernel was launched correctly
-    AssertCuda(cudaPeekAtLastError());
-    // Check that there was no problem during the execution of the kernel
-    AssertCuda(cudaDeviceSynchronize());
+    Kokkos::fence();
 
     deallog << "OK" << std::endl;
   };
@@ -160,7 +157,6 @@ main()
   initlog();
 
   Kokkos::initialize();
-  init_cuda();
 
   test<2, 1>();
 
diff --git a/tests/cuda/matrix_free_no_index_initialize.output b/tests/matrix_free/matrix_free_device_no_index_initialize.output
similarity index 100%
rename from tests/cuda/matrix_free_no_index_initialize.output
rename to tests/matrix_free/matrix_free_device_no_index_initialize.output
diff --git a/tests/cuda/matrix_free_reinit_01.cc b/tests/matrix_free/matrix_free_device_reinit_01.cc
similarity index 99%
rename from tests/cuda/matrix_free_reinit_01.cc
rename to tests/matrix_free/matrix_free_device_reinit_01.cc
index 0d072baba3..6067986f16 100644
--- a/tests/cuda/matrix_free_reinit_01.cc
+++ b/tests/matrix_free/matrix_free_device_reinit_01.cc
@@ -91,7 +91,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);
   MPILogInitAll mpi_inilog;
 
   test<2, 1>();
diff --git a/tests/cuda/matrix_free_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output b/tests/matrix_free/matrix_free_device_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output
similarity index 100%
rename from tests/cuda/matrix_free_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output
rename to tests/matrix_free/matrix_free_device_reinit_01.with_mpi=on.with_p4est=on.mpirun=2.output
diff --git a/tests/cuda/matrix_vector_common.h b/tests/matrix_free/matrix_vector_device_common.h
similarity index 98%
rename from tests/cuda/matrix_vector_common.h
rename to tests/matrix_free/matrix_vector_device_common.h
index 7420cf5032..de11174444 100644
--- a/tests/cuda/matrix_vector_common.h
+++ b/tests/matrix_free/matrix_vector_device_common.h
@@ -47,7 +47,7 @@
 #include "../tests.h"
 
 #include "Kokkos_Core.hpp"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
 
 
 // forward declare this function. will be implemented in .cc files
@@ -116,7 +116,7 @@ do_test(const DoFHandler<dim> &          dof,
 
   in_device.import_elements(in, VectorOperation::insert);
   mf.vmult(out_device, in_device);
-  cudaDeviceSynchronize();
+  Kokkos::fence();
   out.import_elements(out_device, VectorOperation::insert);
 
   // assemble sparse matrix with (\nabla v, \nabla u) + (v, 10 * u)
@@ -197,8 +197,6 @@ main()
 
   Kokkos::initialize();
 
-  init_cuda();
-
   {
     deallog.push("double");
     deallog.push("2d");