From: Daniel Arndt <arndtd@ornl.gov>
Date: Tue, 20 Dec 2022 00:31:34 +0000 (-0500)
Subject: Merge remote-tracking branch 'upstream/master' into kokkos_la_d_vector_kokkos
X-Git-Tag: v9.5.0-rc1~697^2~14
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=46e9cd44216fe4fda53c20848de49484c87d73f9;p=dealii.git

Merge remote-tracking branch 'upstream/master' into kokkos_la_d_vector_kokkos
---

46e9cd44216fe4fda53c20848de49484c87d73f9
diff --cc include/deal.II/base/memory_space_data.h
index 564d7afccb,19c615a628..f30bbb32a1
--- a/include/deal.II/base/memory_space_data.h
+++ b/include/deal.II/base/memory_space_data.h
@@@ -38,9 -38,10 +38,10 @@@ namespace MemorySpac
  {
    /**
     * Structure which stores data on the host or the device depending on the
-    * template parameter @p MemorySpace. The data is copied into the structure
-    * which then owns the data and will release the memory when the destructor is
-    * called.
+    * template parameter @p MemorySpace. Valid choices are MemorySpace::Host,
 -   * MemorySpace::Device, and MemorySpace::CUDA (if CUDA was enabled in
++   * MemorySpace::Default, and MemorySpace::CUDA (if CUDA was enabled in
+    * deal.II). The data is copied into the structure which then owns the data
+    * and will release the memory when the destructor is called.
     */
    template <typename T, typename MemorySpace>
    struct MemorySpaceData
diff --cc include/deal.II/lac/la_parallel_vector.templates.h
index e310f207ec,21b101d18b..f0060403e5
--- a/include/deal.II/lac/la_parallel_vector.templates.h
+++ b/include/deal.II/lac/la_parallel_vector.templates.h
@@@ -86,10 -86,9 +86,10 @@@ namespace LinearAlgebr
        template <typename Number, typename MemorySpaceType>
        struct la_parallel_vector_templates_functions
        {
 -        static_assert(std::is_same<MemorySpaceType, MemorySpace::Host>::value ||
 -                        std::is_same<MemorySpaceType, MemorySpace::CUDA>::value,
 -                      "MemorySpace should be Host or CUDA");
 +        static_assert(
 +          std::is_same<MemorySpaceType, MemorySpace::Host>::value ||
-             std::is_same<MemorySpaceType, MemorySpace::Device>::value,
-           "MemorySpace should be Host or Device");
++            std::is_same<MemorySpaceType, MemorySpace::Default>::value,
++          "MemorySpace should be Host or Default");
  
          static void
          resize_val(
@@@ -317,28 -316,26 +317,28 @@@
          }
        };
  
 -#ifdef DEAL_II_COMPILER_CUDA_AWARE
        template <typename Number>
 -      struct la_parallel_vector_templates_functions<Number,
 -                                                    ::dealii::MemorySpace::CUDA>
 +      struct la_parallel_vector_templates_functions<
 +        Number,
-         ::dealii::MemorySpace::Device>
++        ::dealii::MemorySpace::Default>
        {
          using size_type = types::global_dof_index;
  
          static void
 -        resize_val(const types::global_dof_index new_alloc_size,
 -                   types::global_dof_index &     allocated_size,
 -                   ::dealii::MemorySpace::
 -                     MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data,
 -                   const MPI_Comm &comm_sm)
 +        resize_val(
 +          const types::global_dof_index new_alloc_size,
 +          types::global_dof_index &     allocated_size,
 +          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                  ::dealii::MemorySpace::Device>
++                                                 ::dealii::MemorySpace::Default>
 +            &             data,
 +          const MPI_Comm &comm_sm)
          {
            (void)comm_sm;
  
            static_assert(
              std::is_same<Number, float>::value ||
                std::is_same<Number, double>::value,
-             "Number should be float or double for Device memory space");
 -            "Number should be float or double for CUDA memory space");
++            "Number should be float or double for Default memory space");
  
            if (new_alloc_size > allocated_size)
              {
@@@ -362,10 -359,10 +362,10 @@@
            const ReadWriteVector<Number> &V,
            VectorOperation::values        operation,
            std::shared_ptr<const Utilities::MPI::Partitioner>
 -                          communication_pattern,
 +            &             communication_pattern,
            const IndexSet &locally_owned_elem,
            ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                  ::dealii::MemorySpace::Device>
 -                                                 ::dealii::MemorySpace::CUDA>
++                                                 ::dealii::MemorySpace::Default>
              &data)
          {
            Assert(
@@@ -375,123 -372,108 +375,124 @@@
                "Only VectorOperation::add and VectorOperation::insert are allowed"));
  
            ::dealii::LinearAlgebra::distributed::
-             Vector<Number, ::dealii::MemorySpace::Device>
 -            Vector<Number, ::dealii::MemorySpace::CUDA>
++            Vector<Number, ::dealii::MemorySpace::Default>
                tmp_vector(communication_pattern);
  
            // fill entries from ReadWriteVector into the distributed vector,
            // including ghost entries. this is not really efficient right now
            // because indices are translated twice, once by nth_index_in_set(i)
            // and once for operator() of tmp_vector
 -          const IndexSet &       v_stored   = V.get_stored_elements();
 -          const size_type        n_elements = v_stored.n_elements();
 -          std::vector<size_type> indices(n_elements);
 -          for (size_type i = 0; i < n_elements; ++i)
 -            indices[i] = communication_pattern->global_to_local(
 -              v_stored.nth_index_in_set(i));
 +          const IndexSet &                  v_stored = V.get_stored_elements();
 +          const size_type                   n_elements = v_stored.n_elements();
 +          Kokkos::DefaultHostExecutionSpace host_exec;
 +          Kokkos::View<size_type *, Kokkos::HostSpace> indices("indices",
 +                                                               n_elements);
 +          Kokkos::parallel_for(
 +            Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec,
 +                                                                   0,
 +                                                                   n_elements),
 +            KOKKOS_LAMBDA(size_type i) {
 +              indices[i] = communication_pattern->global_to_local(
 +                v_stored.nth_index_in_set(i));
 +            });
 +          host_exec.fence();
 +
            // Move the indices to the device
-           ::dealii::MemorySpace::Device::kokkos_space::execution_space exec;
-           Kokkos::View<size_type *, ::dealii::MemorySpace::Device::kokkos_space>
 -          size_type *indices_dev;
 -          ::dealii::Utilities::CUDA::malloc(indices_dev, n_elements);
 -          ::dealii::Utilities::CUDA::copy_to_dev(indices, indices_dev);
++          ::dealii::MemorySpace::Default::kokkos_space::execution_space exec;
++          Kokkos::View<size_type *,
++                       ::dealii::MemorySpace::Default::kokkos_space>
 +            indices_dev("indices_dev", n_elements);
 +          Kokkos::deep_copy(exec, indices_dev, indices);
 +          exec.fence();
 +
            // Move the data to the device
-           Kokkos::View<Number *, ::dealii::MemorySpace::Device::kokkos_space>
 -          Number *V_dev;
 -          ::dealii::Utilities::CUDA::malloc(V_dev, n_elements);
 -          cudaError_t cuda_error_code = cudaMemcpy(V_dev,
 -                                                   V.begin(),
 -                                                   n_elements * sizeof(Number),
 -                                                   cudaMemcpyHostToDevice);
 -          AssertCuda(cuda_error_code);
++          Kokkos::View<Number *, ::dealii::MemorySpace::Default::kokkos_space>
 +            V_dev("V_dev", n_elements);
 +          Kokkos::deep_copy(
 +            exec,
 +            V_dev,
 +            Kokkos::View<Number *, Kokkos::HostSpace>(V.begin(), n_elements));
 +          exec.fence();
  
            // Set the values in tmp_vector
 -          const int n_blocks =
 -            1 + n_elements / (::dealii::CUDAWrappers::chunk_size *
 -                              ::dealii::CUDAWrappers::block_size);
 -          ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<Number>
 -            <<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
 -              indices_dev, tmp_vector.begin(), V_dev, n_elements);
 +          Kokkos::parallel_for(
 +            Kokkos::RangePolicy<
-               ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
++              ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
 +              exec, 0, n_elements),
 +            KOKKOS_LAMBDA(size_type i) {
 +              tmp_vector(indices_dev(i)) = V_dev(i);
 +            });
 +          exec.fence();
  
            tmp_vector.compress(operation);
  
            // Copy the local elements of tmp_vector to the right place in val
            IndexSet        tmp_index_set  = tmp_vector.locally_owned_elements();
            const size_type tmp_n_elements = tmp_index_set.n_elements();
 -          indices.resize(tmp_n_elements);
 -          for (size_type i = 0; i < tmp_n_elements; ++i)
 -            indices[i] = locally_owned_elem.index_within_set(
 -              tmp_index_set.nth_index_in_set(i));
 -          ::dealii::Utilities::CUDA::free(indices_dev);
 -          ::dealii::Utilities::CUDA::malloc(indices_dev, tmp_n_elements);
 -          ::dealii::Utilities::CUDA::copy_to_dev(indices, indices_dev);
 +          Kokkos::realloc(indices, tmp_n_elements);
 +          Kokkos::parallel_for(
 +            Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec,
 +                                                                   0,
 +                                                                   n_elements),
 +            KOKKOS_LAMBDA(size_type i) {
 +              indices[i] = locally_owned_elem.index_within_set(
 +                tmp_index_set.nth_index_in_set(i));
 +            });
 +          host_exec.fence();
 +          Kokkos::realloc(indices_dev, tmp_n_elements);
 +          Kokkos::deep_copy(indices_dev,
 +                            Kokkos::View<size_type *>(indices.data(),
 +                                                      tmp_n_elements));
  
            if (operation == VectorOperation::add)
 -            ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_permutated<
 -              Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
 -              indices_dev,
 -              data.values.data(),
 -              tmp_vector.begin(),
 -              tmp_n_elements);
 +            Kokkos::parallel_for(
 +              Kokkos::RangePolicy<
-                 ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
++                ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
 +                exec, 0, n_elements),
 +              KOKKOS_LAMBDA(size_type i) {
 +                data.values(indices_dev(i)) += tmp_vector(i);
 +              });
            else
 -            ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<
 -              Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
 -              indices_dev,
 -              data.values.data(),
 -              tmp_vector.begin(),
 -              tmp_n_elements);
 -
 -          ::dealii::Utilities::CUDA::free(indices_dev);
 -          ::dealii::Utilities::CUDA::free(V_dev);
 +            Kokkos::parallel_for(
 +              Kokkos::RangePolicy<
-                 ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
++                ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
 +                exec, 0, n_elements),
 +              KOKKOS_LAMBDA(size_type i) {
 +                data.values(indices_dev(i)) = tmp_vector(i);
 +              });
 +          exec.fence();
          }
  
          template <typename RealType>
          static void
          linfty_norm_local(const ::dealii::MemorySpace::MemorySpaceData<
                              Number,
-                             ::dealii::MemorySpace::Device> &data,
-                           const unsigned int                size,
-                           RealType &                        result)
 -                            ::dealii::MemorySpace::CUDA> &data,
 -                          const unsigned int              size,
 -                          RealType &                      result)
++                            ::dealii::MemorySpace::Default> &data,
++                          const unsigned int                 size,
++                          RealType &                         result)
          {
            static_assert(std::is_same<Number, RealType>::value,
                          "RealType should be the same type as Number");
  
-           typename ::dealii::MemorySpace::Device::kokkos_space::execution_space
 -          Number *    result_device;
 -          cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
 -          AssertCuda(error_code);
 -          error_code = cudaMemset(result_device, 0, sizeof(Number));
 -
 -          const int n_blocks = 1 + size / (::dealii::CUDAWrappers::chunk_size *
 -                                           ::dealii::CUDAWrappers::block_size);
 -          ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
 -            Number,
 -            ::dealii::LinearAlgebra::CUDAWrappers::kernel::LInfty<Number>>
 -            <<<dim3(n_blocks, 1), dim3(::dealii::CUDAWrappers::block_size)>>>(
 -              result_device, data.values.data(), size);
 -
 -          // Copy the result back to the host
 -          error_code = cudaMemcpy(&result,
 -                                  result_device,
 -                                  sizeof(Number),
 -                                  cudaMemcpyDeviceToHost);
 -          AssertCuda(error_code);
 -          // Free the memory on the device
 -          error_code = cudaFree(result_device);
 -          AssertCuda(error_code);
++          typename ::dealii::MemorySpace::Default::kokkos_space::execution_space
 +            exec;
 +          Kokkos::parallel_reduce(
 +            Kokkos::RangePolicy<
-               ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
++              ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
 +              exec, 0, size),
 +            KOKKOS_LAMBDA(size_type i, RealType & update) {
 +#if KOKKOS_VERSION < 30400
 +              update = fmax(update, fabs(data.values(i)));
 +#elif KOKKOS_VERSION < 30700
 +              update = Kokkos::Experimental::fmax(
 +                update, Kokkos::Experimental::fabs(data.values(i)));
 +#else
 +              update = Kokkos::fmax(update, Kokkos::abs(data.values(i)));
 +#endif
 +            },
 +            Kokkos::Max<RealType, Kokkos::HostSpace>(result));
          }
        };
 -#endif
      } // namespace internal
  
  
@@@ -978,8 -974,9 +979,9 @@@
        // allocate import_data in case it is not set up yet
        if (partitioner->n_import_indices() > 0)
          {
 -#  if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -          if (std::is_same<MemorySpaceType, dealii::MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-           if (std::is_same<MemorySpaceType, dealii::MemorySpace::Device>::value)
++          if (std::is_same<MemorySpaceType,
++                           dealii::MemorySpace::Default>::value)
              {
                if (import_data.values_host_buffer.size() == 0)
                  Kokkos::resize(import_data.values_host_buffer,
@@@ -994,8 -997,9 +996,8 @@@
              }
          }
  
 -#  if defined DEAL_II_COMPILER_CUDA_AWARE && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -      if (std::is_same<MemorySpaceType, dealii::MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-       if (std::is_same<MemorySpaceType, dealii::MemorySpace::Device>::value)
++      if (std::is_same<MemorySpaceType, dealii::MemorySpace::Default>::value)
          {
            // Move the data to the host and then move it back to the
            // device. We use values to store the elements because the function
@@@ -1051,8 -1055,9 +1053,8 @@@
  
        // make this function thread safe
        std::lock_guard<std::mutex> lock(mutex);
 -#  if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -      if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-       if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
++      if (std::is_same<MemorySpaceType, MemorySpace::Default>::value)
          {
            Assert(partitioner->n_import_indices() == 0 ||
                     import_data.values_host_buffer.size() != 0,
@@@ -1096,7 -1106,7 +1098,6 @@@
                  partitioner->n_ghost_indices()),
                compress_requests);
          }
--
  #else
        (void)operation;
  #endif
@@@ -1122,8 -1132,9 +1123,8 @@@
        // allocate import_data in case it is not set up yet
        if (partitioner->n_import_indices() > 0)
          {
 -#  if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -          if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-           if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
++          if (std::is_same<MemorySpaceType, MemorySpace::Default>::value)
              {
                if (import_data.values_host_buffer.size() == 0)
                  Kokkos::resize(import_data.values_host_buffer,
@@@ -1143,8 -1154,9 +1144,8 @@@
              }
          }
  
 -#  if defined DEAL_II_COMPILER_CUDA_AWARE && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -      if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-       if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
++      if (std::is_same<MemorySpaceType, MemorySpace::Default>::value)
          {
            // Move the data to the host and then move it back to the
            // device. We use values to store the elements because the function
@@@ -1205,8 -1217,9 +1206,8 @@@
            // make this function thread safe
            std::lock_guard<std::mutex> lock(mutex);
  
 -#  if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
 -    !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
 -          if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
 +#  if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
-           if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
++          if (std::is_same<MemorySpaceType, MemorySpace::Default>::value)
              {
                partitioner->export_to_ghosted_array_finish(
                  ArrayView<Number, MemorySpace::Host>(
diff --cc include/deal.II/lac/read_write_vector.templates.h
index 08135c2d10,d1b9d6c070..c9a8f33470
--- a/include/deal.II/lac/read_write_vector.templates.h
+++ b/include/deal.II/lac/read_write_vector.templates.h
@@@ -109,8 -109,8 +109,8 @@@ namespace LinearAlgebr
  
          static_assert(
            std::is_same<MemorySpace, ::dealii::MemorySpace::Host>::value ||
-             std::is_same<MemorySpace, ::dealii::MemorySpace::Device>::value,
-           "MemorySpace should be Host or Device");
 -            std::is_same<MemorySpace, ::dealii::MemorySpace::CUDA>::value,
 -          "MemorySpace should be Host or CUDA");
++            std::is_same<MemorySpace, ::dealii::MemorySpace::Default>::value,
++          "MemorySpace should be Host or Default");
        }
      };
  
@@@ -160,8 -160,9 +160,8 @@@
  
  
  
 -#ifdef DEAL_II_COMPILER_CUDA_AWARE
      template <typename Number>
-     struct read_write_vector_functions<Number, ::dealii::MemorySpace::Device>
 -    struct read_write_vector_functions<Number, ::dealii::MemorySpace::CUDA>
++    struct read_write_vector_functions<Number, ::dealii::MemorySpace::Default>
      {
        using size_type = types::global_dof_index;
  
@@@ -177,12 -178,11 +177,12 @@@
  
          const unsigned int n_elements =
            communication_pattern->locally_owned_size();
 -        cudaError_t cuda_error_code = cudaMemcpy(tmp_vector.begin(),
 -                                                 values,
 -                                                 n_elements * sizeof(Number),
 -                                                 cudaMemcpyDeviceToHost);
 -        AssertCuda(cuda_error_code);
 +        Kokkos::deep_copy(
 +          Kokkos::View<Number *, Kokkos::HostSpace>(tmp_vector.begin(),
 +                                                    n_elements),
 +          Kokkos::View<const Number *,
-                        ::dealii::MemorySpace::Device::kokkos_space>(
++                       ::dealii::MemorySpace::Default::kokkos_space>(
 +            values, n_elements));
          tmp_vector.update_ghost_values();
  
          const IndexSet &stored = rw_vector.get_stored_elements();
diff --cc include/deal.II/lac/vector_operations_internal.h
index 4f2a2b475a,21d08f4413..92300d124e
--- a/include/deal.II/lac/vector_operations_internal.h
+++ b/include/deal.II/lac/vector_operations_internal.h
@@@ -1484,9 -1484,9 +1484,9 @@@ namespace interna
          ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace> & /*data*/)
        {
          static_assert(
-           std::is_same<MemorySpace, ::dealii::MemorySpace::Device>::value &&
 -          std::is_same<MemorySpace, ::dealii::MemorySpace::CUDA>::value &&
++          std::is_same<MemorySpace, ::dealii::MemorySpace::Default>::value &&
              std::is_same<Number, Number2>::value,
-           "For the Device MemorySpace Number and Number2 should be the same type");
 -          "For the CUDA MemorySpace Number and Number2 should be the same type");
++          "For the Default MemorySpace Number and Number2 should be the same type");
        }
  
        static void
@@@ -2110,7 -2111,7 +2110,7 @@@
                                                 ::dealii::MemorySpace::Host>
            &data,
          std::enable_if_t<
-           std::is_same<MemorySpace2, ::dealii::MemorySpace::Device>::value,
 -          std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::value,
++          std::is_same<MemorySpace2, ::dealii::MemorySpace::Default>::value,
            int> = 0)
        {
          if (operation == VectorOperation::insert)
@@@ -2130,24 -2132,30 +2130,24 @@@
  
  
  
 -#ifdef DEAL_II_COMPILER_CUDA_AWARE
      template <typename Number>
-     struct functions<Number, Number, ::dealii::MemorySpace::Device>
 -    struct functions<Number, Number, ::dealii::MemorySpace::CUDA>
++    struct functions<Number, Number, ::dealii::MemorySpace::Default>
      {
 -      static const int block_size =
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::block_size;
 -      static const int chunk_size =
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::chunk_size;
 -
        static void
        copy(
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
 -        cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
 -                                                 v_data.values.data(),
 -                                                 size * sizeof(Number),
 -                                                 cudaMemcpyDeviceToDevice);
 -        AssertCuda(cuda_error_code);
 +        Kokkos::deep_copy(
 +          Kokkos::subview(data.values,
 +                          Kokkos::pair<size_type, size_type>(0, size)),
 +          Kokkos::subview(v_data.values,
 +                          Kokkos::pair<size_type, size_type>(0, size)));
        }
  
        static void
@@@ -2155,13 -2163,13 +2155,13 @@@
            const size_type size,
            const Number    s,
            ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                  ::dealii::MemorySpace::Device>
 -                                                 ::dealii::MemorySpace::CUDA>
++                                                 ::dealii::MemorySpace::Default>
              &data)
        {
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::set<Number>
 -          <<<n_blocks, block_size>>>(data.values.data(), s, size);
 -        AssertCudaKernel();
 +        Kokkos::deep_copy(
 +          Kokkos::subview(data.values,
 +                          Kokkos::pair<size_type, size_type>(0, size)),
 +          s);
        }
  
        static void
@@@ -2169,20 -2177,18 +2169,19 @@@
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
 -          <<<n_blocks, block_size>>>(data.values.data(),
 -                                     1.,
 -                                     v_data.values.data(),
 -                                     size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(int i) { data.values(i) += v_data.values(i); });
 +        exec.fence();
        }
  
        static void
@@@ -2190,20 -2196,18 +2189,19 @@@
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
 -          <<<n_blocks, block_size>>>(data.values.data(),
 -                                     -1.,
 -                                     v_data.values.data(),
 -                                     size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) { data.values(i) -= v_data.values(i); });
 +        exec.fence();
        }
  
        static void
@@@ -2212,18 -2216,13 +2210,17 @@@
          const size_type size,
          Number          a,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_add<Number>
 -          <<<n_blocks, block_size>>>(data.values.data(), a, size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) { data.values(i) += a; });
 +        exec.fence();
        }
  
        static void
@@@ -2232,22 -2231,18 +2229,21 @@@
          const size_type size,
          const Number    a,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
 -          <<<n_blocks, block_size>>>(data.values.data(),
 -                                     a,
 -                                     v_data.values.data(),
 -                                     size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) += a * v_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static void
@@@ -2257,24 -2252,22 +2253,23 @@@
          const Number    a,
          const Number    b,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &w_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aVbW<Number>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
 -                                                    a,
 -                                                    v_data.values.data(),
 -                                                    b,
 -                                                    w_data.values.data(),
 -                                                    size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) += a * v_data.values(i) + b * w_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static void
@@@ -2283,22 -2276,16 +2278,21 @@@
          const size_type size,
          const Number    x,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(
 -            x, data.values.data(), 1., v_data.values.data(), size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) = x * data.values(i) + v_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static void
@@@ -2308,22 -2295,16 +2302,21 @@@
          const Number    x,
          const Number    a,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(
 -            x, data.values.data(), a, v_data.values.data(), size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) = x * data.values(i) + a * v_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static void
@@@ -2334,25 -2315,23 +2327,24 @@@
          const Number    a,
          const Number    b,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &w_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(x,
 -                                                    data.values.data(),
 -                                                    a,
 -                                                    v_data.values.data(),
 -                                                    b,
 -                                                    w_data.values.data(),
 -                                                    size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) =
 +              x * data.values(i) + a * v_data.values(i) + b * w_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static void
@@@ -2361,18 -2340,13 +2353,17 @@@
          const size_type size,
          const Number    factor,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_scale<Number>
 -          <<<n_blocks, block_size>>>(data.values.data(), factor, size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) { data.values(i) *= factor; });
 +        exec.fence();
        }
  
        static void
@@@ -2380,20 -2354,17 +2371,19 @@@
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::scale<Number>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
 -                                                    v_data.values.data(),
 -                                                    size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) { data.values(i) *= v_data.values(i); });
 +        exec.fence();
        }
  
        static void
@@@ -2402,22 -2373,18 +2392,21 @@@
          const size_type size,
          const Number    a,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::equ<Number>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
 -                                                    a,
 -                                                    v_data.values.data(),
 -                                                    size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) = a * v_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static void
@@@ -2427,50 -2394,63 +2416,48 @@@
          const Number    a,
          const Number    b,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &w_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::equ<Number>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
 -                                                    a,
 -                                                    v_data.values.data(),
 -                                                    b,
 -                                                    w_data.values.data(),
 -                                                    size);
 -        AssertCudaKernel();
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_for(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i) {
 +            data.values(i) = a * v_data.values(i) + b * w_data.values(i);
 +          });
 +        exec.fence();
        }
  
        static Number
        dot(const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
            const size_type size,
            const ::dealii::MemorySpace::
-             MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -            MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++            MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
            ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                  ::dealii::MemorySpace::Device>
 -                                                 ::dealii::MemorySpace::CUDA>
++                                                 ::dealii::MemorySpace::Default>
              &data)
        {
 -        Number *    result_device;
 -        cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
 -        AssertCuda(error_code);
 -        error_code = cudaMemset(result_device, 0, sizeof(Number));
 -        AssertCuda(error_code);
 -
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::double_vector_reduction<
 -          Number,
 -          ::dealii::LinearAlgebra::CUDAWrappers::kernel::DotProduct<Number>>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
 -                                                    data.values.data(),
 -                                                    v_data.values.data(),
 -                                                    static_cast<unsigned int>(
 -                                                      size));
 -        AssertCudaKernel();
 -
 -        // Copy the result back to the host
          Number result;
 -        error_code = cudaMemcpy(&result,
 -                                result_device,
 -                                sizeof(Number),
 -                                cudaMemcpyDeviceToHost);
 -        AssertCuda(error_code);
 -        // Free the memory on the device
 -        error_code = cudaFree(result_device);
 -        AssertCuda(error_code);
  
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        AssertIsFinite(result);
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i, Number & update) {
 +            update += data.values(i) * v_data.values(i);
 +          },
 +          result);
  
 +        AssertIsFinite(result);
          return result;
        }
  
@@@ -2480,8 -2460,9 +2467,8 @@@
                 &             thread_loop_partitioner,
               const size_type size,
               real_type &     sum,
 -             ::dealii::MemorySpace::MemorySpaceData<Number,
 -                                                    ::dealii::MemorySpace::CUDA>
 -               &data)
 +             ::dealii::MemorySpace::
-                MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data)
++               MemorySpaceData<Number, ::dealii::MemorySpace::Default> &data)
        {
          sum = dot(thread_loop_partitioner, size, data, data);
        }
@@@ -2491,23 -2472,32 +2478,22 @@@
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
          const size_type size,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data)
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
 -      {
 -        Number *    result_device;
 -        cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
 -        AssertCuda(error_code);
 -        error_code = cudaMemset(result_device, 0, sizeof(Number));
 -
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
 -          Number,
 -          ::dealii::LinearAlgebra::CUDAWrappers::kernel::ElemSum<Number>>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
 -                                                    data.values.data(),
 -                                                    size);
 -
 -        // Copy the result back to the host
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &data)
 +      {
          Number result;
 -        error_code = cudaMemcpy(&result,
 -                                result_device,
 -                                sizeof(Number),
 -                                cudaMemcpyDeviceToHost);
 -        AssertCuda(error_code);
 -        // Free the memory on the device
 -        error_code = cudaFree(result_device);
 -        AssertCuda(error_code);
  
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i, Number & update) {
 +            update += data.values(i);
 +          },
 +          result);
 +
 +        AssertIsFinite(result);
          return result;
        }
  
@@@ -2518,57 -2508,44 +2504,55 @@@
          const size_type size,
          real_type &     sum,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        Number *    result_device;
 -        cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
 -        AssertCuda(error_code);
 -        error_code = cudaMemset(result_device, 0, sizeof(Number));
 -
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
 -          Number,
 -          ::dealii::LinearAlgebra::CUDAWrappers::kernel::L1Norm<Number>>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
 -                                                    data.values.data(),
 -                                                    size);
 -
 -        // Copy the result back to the host
 -        error_code = cudaMemcpy(&sum,
 -                                result_device,
 -                                sizeof(Number),
 -                                cudaMemcpyDeviceToHost);
 -        AssertCuda(error_code);
 -        // Free the memory on the device
 -        error_code = cudaFree(result_device);
 -        AssertCuda(error_code);
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i, Number & update) {
 +#if KOKKOS_VERSION < 30400
 +            update += fabs(data.values(i));
 +#elif KOKKOS_VERSION < 30700
 +            update += Kokkos::Experimental::fabs(data.values(i));
 +#else
 +            update += Kokkos::fabs(data.values(i));
 +#endif
 +          },
 +          sum);
        }
  
        template <typename real_type>
        static void
        norm_p(
          const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
 -        const size_type,
 -        real_type &,
 -        real_type,
 +        const size_type size,
 +        real_type &     sum,
 +        real_type       exp,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA> &)
++                                               ::dealii::MemorySpace::Default>
 +          &data)
        {
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
 -        Assert(false, ExcNotImplemented());
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i, Number & update) {
 +#if KOKKOS_VERSION < 30400
 +            update += pow(fabs(data.values(i)), exp);
 +#elif KOKKOS_VERSION < 30700
 +            update += Kokkos::Experimental::pow(
 +              Kokkos::Experimental::fabs(data.values(i)), exp);
 +#else
 +            update += Kokkos::pow(Kokkos::fabs(data.values(i)), exp);
 +#endif
 +          },
 +          sum);
        }
  
        static Number
@@@ -2577,29 -2554,33 +2561,28 @@@
          const size_type size,
          const Number    a,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &v_data,
          const ::dealii::MemorySpace::
-           MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
 -          MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
++          MemorySpaceData<Number, ::dealii::MemorySpace::Default> &w_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data)
        {
 -        Number *    res_d;
 -        cudaError_t error_code = cudaMalloc(&res_d, sizeof(Number));
 -        AssertCuda(error_code);
 -        error_code = cudaMemset(res_d, 0, sizeof(Number));
 -        AssertCuda(error_code);
 -
 -        const int n_blocks = 1 + size / (chunk_size * block_size);
 -        ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_and_dot<Number>
 -          <<<dim3(n_blocks, 1), dim3(block_size)>>>(res_d,
 -                                                    data.values.data(),
 -                                                    v_data.values.data(),
 -                                                    w_data.values.data(),
 -                                                    a,
 -                                                    size);
 -
          Number res;
 -        error_code =
 -          cudaMemcpy(&res, res_d, sizeof(Number), cudaMemcpyDeviceToHost);
 -        AssertCuda(error_code);
 -        error_code = cudaFree(res_d);
 +
-         auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
++        auto exec = typename ::dealii::MemorySpace::Default::kokkos_space::
 +          execution_space{};
 +        Kokkos::parallel_reduce(
 +          Kokkos::RangePolicy<
-             ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
-                                                                           0,
-                                                                           size),
++            ::dealii::MemorySpace::Default::kokkos_space::execution_space>(
++            exec, 0, size),
 +          KOKKOS_LAMBDA(size_type i, Number & update) {
 +            data.values(i) += a * v_data.values(i);
 +            update +=
 +              data.values(i) * Number(numbers::NumberTraits<Number>::conjugate(
 +                                 w_data.values(i)));
 +          },
 +          res);
  
          return res;
        }
@@@ -2614,10 -2595,10 +2597,10 @@@
          const ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace2>
            &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data,
          std::enable_if_t<
-           std::is_same<MemorySpace2, ::dealii::MemorySpace::Device>::value,
 -          std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::value,
++          std::is_same<MemorySpace2, ::dealii::MemorySpace::Default>::value,
            int> = 0)
        {
          if (operation == VectorOperation::insert)
@@@ -2644,7 -2625,7 +2627,7 @@@
          const ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace2>
            &v_data,
          ::dealii::MemorySpace::MemorySpaceData<Number,
-                                                ::dealii::MemorySpace::Device>
 -                                               ::dealii::MemorySpace::CUDA>
++                                               ::dealii::MemorySpace::Default>
            &data,
          std::enable_if_t<
            std::is_same<MemorySpace2, ::dealii::MemorySpace::Host>::value,
diff --cc source/lac/la_parallel_vector_cuda.cc
index 76dc6c01eb,e4ebd6c6f5..91fc69a625
--- a/source/lac/la_parallel_vector_cuda.cc
+++ b/source/lac/la_parallel_vector_cuda.cc
@@@ -24,50 -24,40 +24,50 @@@ namespace LinearAlgebr
  {
    namespace distributed
    {
-     template class Vector<float, ::dealii::MemorySpace::Device>;
-     template class Vector<double, ::dealii::MemorySpace::Device>;
 -    template class Vector<float, ::dealii::MemorySpace::CUDA>;
 -    template class Vector<double, ::dealii::MemorySpace::CUDA>;
++    template class Vector<float, ::dealii::MemorySpace::Default>;
++    template class Vector<double, ::dealii::MemorySpace::Default>;
 +
      template void
      Vector<float, ::dealii::MemorySpace::Host>::import<
-       ::dealii::MemorySpace::Device>(
-       const Vector<float, ::dealii::MemorySpace::Device> &,
 -      ::dealii::MemorySpace::CUDA>(
 -      const Vector<float, ::dealii::MemorySpace::CUDA> &,
++      ::dealii::MemorySpace::Default>(
++      const Vector<float, ::dealii::MemorySpace::Default> &,
        VectorOperation::values);
      template void
      Vector<double, ::dealii::MemorySpace::Host>::import<
-       ::dealii::MemorySpace::Device>(
-       const Vector<double, ::dealii::MemorySpace::Device> &,
 -      ::dealii::MemorySpace::CUDA>(
 -      const Vector<double, ::dealii::MemorySpace::CUDA> &,
++      ::dealii::MemorySpace::Default>(
++      const Vector<double, ::dealii::MemorySpace::Default> &,
        VectorOperation::values);
  
      template void
-     Vector<float, ::dealii::MemorySpace::Device>::import<
 -    Vector<float, ::dealii::MemorySpace::CUDA>::import<
++    Vector<float, ::dealii::MemorySpace::Default>::import<
        ::dealii::MemorySpace::Host>(
        const Vector<float, ::dealii::MemorySpace::Host> &,
        VectorOperation::values);
      template void
-     Vector<double, ::dealii::MemorySpace::Device>::import<
 -    Vector<double, ::dealii::MemorySpace::CUDA>::import<
++    Vector<double, ::dealii::MemorySpace::Default>::import<
        ::dealii::MemorySpace::Host>(
        const Vector<double, ::dealii::MemorySpace::Host> &,
        VectorOperation::values);
  
      template void
-     Vector<float, ::dealii::MemorySpace::Device>::import<
-       ::dealii::MemorySpace::Device>(
-       const Vector<float, ::dealii::MemorySpace::Device> &,
 -    Vector<float, ::dealii::MemorySpace::CUDA>::import<
 -      ::dealii::MemorySpace::CUDA>(
 -      const Vector<float, ::dealii::MemorySpace::CUDA> &,
++    Vector<float, ::dealii::MemorySpace::Default>::import<
++      ::dealii::MemorySpace::Default>(
++      const Vector<float, ::dealii::MemorySpace::Default> &,
        VectorOperation::values);
      template void
-     Vector<double, ::dealii::MemorySpace::Device>::import<
-       ::dealii::MemorySpace::Device>(
-       const Vector<double, ::dealii::MemorySpace::Device> &,
 -    Vector<double, ::dealii::MemorySpace::CUDA>::import<
 -      ::dealii::MemorySpace::CUDA>(
 -      const Vector<double, ::dealii::MemorySpace::CUDA> &,
++    Vector<double, ::dealii::MemorySpace::Default>::import<
++      ::dealii::MemorySpace::Default>(
++      const Vector<double, ::dealii::MemorySpace::Default> &,
        VectorOperation::values);
 +
 +    template void
-     Vector<float, ::dealii::MemorySpace::Device>::reinit<float>(
-       const Vector<float, ::dealii::MemorySpace::Device> &,
++    Vector<float, ::dealii::MemorySpace::Default>::reinit<float>(
++      const Vector<float, ::dealii::MemorySpace::Default> &,
 +      const bool);
 +    template void
-     Vector<double, ::dealii::MemorySpace::Device>::reinit<double>(
-       const Vector<double, ::dealii::MemorySpace::Device> &,
++    Vector<double, ::dealii::MemorySpace::Default>::reinit<double>(
++      const Vector<double, ::dealii::MemorySpace::Default> &,
 +      const bool);
    } // namespace distributed
  } // namespace LinearAlgebra
  
diff --cc source/lac/read_write_vector.cc
index a6dc6febd2,f1ca4fee7c..23ced9c5bf
--- a/source/lac/read_write_vector.cc
+++ b/source/lac/read_write_vector.cc
@@@ -46,22 -46,12 +46,22 @@@ namespace LinearAlgebr
      const distributed::Vector<float, ::dealii::MemorySpace::Host> &,
      VectorOperation::values,
      const std::shared_ptr<const Utilities::MPI::CommunicationPatternBase> &);
 +  template void
 +  ReadWriteVector<float>::import(
-     const distributed::Vector<float, ::dealii::MemorySpace::Device> &,
++    const distributed::Vector<float, ::dealii::MemorySpace::Default> &,
 +    VectorOperation::values,
 +    const std::shared_ptr<const Utilities::MPI::CommunicationPatternBase> &);
  
    template void
    ReadWriteVector<double>::import(
      const distributed::Vector<double, ::dealii::MemorySpace::Host> &,
      VectorOperation::values,
      const std::shared_ptr<const Utilities::MPI::CommunicationPatternBase> &);
 +  template void
 +  ReadWriteVector<double>::import(
-     const distributed::Vector<double, ::dealii::MemorySpace::Device> &,
++    const distributed::Vector<double, ::dealii::MemorySpace::Default> &,
 +    VectorOperation::values,
 +    const std::shared_ptr<const Utilities::MPI::CommunicationPatternBase> &);
  #  ifdef DEAL_II_WITH_COMPLEX_VALUES
    template void
    ReadWriteVector<std::complex<float>>::import(
diff --cc source/lac/vector_memory.cc
index db9ea9dfa2,461c1cf50c..5ab7a4c8d6
--- a/source/lac/vector_memory.cc
+++ b/source/lac/vector_memory.cc
@@@ -30,14 -30,6 +30,14 @@@
  DEAL_II_NAMESPACE_OPEN
  
  #include "vector_memory.inst"
 +template class VectorMemory<
-   LinearAlgebra::distributed::Vector<float, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<float, MemorySpace::Default>>;
 +template class VectorMemory<
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default>>;
 +template class GrowingVectorMemory<
-   LinearAlgebra::distributed::Vector<float, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<float, MemorySpace::Default>>;
 +template class GrowingVectorMemory<
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default>>;
  
  namespace internal
  {
@@@ -52,12 -44,6 +52,12 @@@
      release_all_unused_memory()
      {
  #include "vector_memory_release.inst"
 +      dealii::GrowingVectorMemory<dealii::LinearAlgebra::distributed::Vector<
 +        float,
-         MemorySpace::Device>>::release_unused_memory();
++        MemorySpace::Default>>::release_unused_memory();
 +      dealii::GrowingVectorMemory<dealii::LinearAlgebra::distributed::Vector<
 +        double,
-         MemorySpace::Device>>::release_unused_memory();
++        MemorySpace::Default>>::release_unused_memory();
  #ifdef DEAL_II_WITH_CUDA
        release_all_unused_cuda_memory();
  #endif
diff --cc tests/lac/parallel_vector_01.cc
index 60dfbc9f2b,5021d1aa16..cad97db078
--- a/tests/lac/parallel_vector_01.cc
+++ b/tests/lac/parallel_vector_01.cc
@@@ -44,7 -44,7 +44,7 @@@ test(
    IndexSet local_owned(numproc * 2);
    local_owned.add_range(myid * 2, myid * 2 + 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_owned, MPI_COMM_WORLD);
  
    // set local values
diff --cc tests/lac/parallel_vector_02.cc
index 891f0ee834,e7664f4e1c..79e9922cee
--- a/tests/lac/parallel_vector_02.cc
+++ b/tests/lac/parallel_vector_02.cc
@@@ -47,7 -47,7 +47,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
diff --cc tests/lac/parallel_vector_03.cc
index 5d686847c8,0e34c7ed54..33145afb4e
--- a/tests/lac/parallel_vector_03.cc
+++ b/tests/lac/parallel_vector_03.cc
@@@ -47,7 -47,7 +47,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
diff --cc tests/lac/parallel_vector_03a.cc
index a7ff1fc723,47f510ab7e..190dc37fc1
--- a/tests/lac/parallel_vector_03a.cc
+++ b/tests/lac/parallel_vector_03a.cc
@@@ -80,7 -80,7 +80,7 @@@ test(
        local_relevant.add_range(4, 7);
      }
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
    v = 0.;
  
diff --cc tests/lac/parallel_vector_04.cc
index a3cb093854,e928ed32b1..08f2e81776
--- a/tests/lac/parallel_vector_04.cc
+++ b/tests/lac/parallel_vector_04.cc
@@@ -48,7 -54,7 +48,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
@@@ -70,16 -76,12 +70,17 @@@
      {
        unsigned int local_index = partitioner->global_to_local(1);
        double *     values_dev  = v.get_values();
-       Kokkos::deep_copy(Kokkos::View<double, MemorySpace::Device::kokkos_space>(
-                           values_dev + local_index),
-                         7);
 -      set_value<<<1, 1>>>(values_dev, local_index, 7);
++      Kokkos::deep_copy(
++        Kokkos::View<double, MemorySpace::Default::kokkos_space>(values_dev +
++                                                                 local_index),
++        7);
      }
  
 -  unsigned int        allocated_size = local_relevant.n_elements();
 -  std::vector<double> v_host(allocated_size);
 -  Utilities::CUDA::copy_to_host(v.get_values(), v_host);
 +  unsigned int allocated_size = local_relevant.n_elements();
-   Kokkos::View<double *, MemorySpace::Device::kokkos_space> v_device(
++  Kokkos::View<double *, MemorySpace::Default::kokkos_space> v_device(
 +    v.get_values(), allocated_size);
 +  Kokkos::View<double *, Kokkos::HostSpace> v_host("v_host", allocated_size);
 +  Kokkos::deep_copy(v_host, v_device);
  
    AssertThrow(v_host[partitioner->global_to_local(myid * 2)] == myid * 4.0,
                ExcInternalError());
diff --cc tests/lac/parallel_vector_05.cc
index 9cf4631e2f,1519bd97e6..005e4e1c17
--- a/tests/lac/parallel_vector_05.cc
+++ b/tests/lac/parallel_vector_05.cc
@@@ -47,7 -47,7 +47,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
diff --cc tests/lac/parallel_vector_06.cc
index 01b1432a56,bfe06fffae..804248706c
--- a/tests/lac/parallel_vector_06.cc
+++ b/tests/lac/parallel_vector_06.cc
@@@ -48,7 -48,7 +48,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_owned, MPI_COMM_WORLD);
  
    // set local values
@@@ -102,7 -102,7 +102,7 @@@
    {
      const double norm_sqr = v.l2_norm() * v.l2_norm();
      AssertThrow(std::fabs(v * v - norm_sqr) < 1e-15, ExcInternalError());
-     LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
 -    LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
++    LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v2;
      v2 = v;
      AssertThrow(std::fabs(v2 * v - norm_sqr) < 1e-15, ExcInternalError());
  
@@@ -116,7 -116,7 +116,7 @@@
      bool allzero = v.all_zero();
      if (myid == 0)
        deallog << " v==0 ? " << allzero << std::endl;
-     LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
 -    LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
++    LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v2;
      v2.reinit(v);
      allzero = v2.all_zero();
      if (myid == 0)
diff --cc tests/lac/parallel_vector_08.cc
index c33c5dadf0,f0c4716232..b9c4b6cd97
--- a/tests/lac/parallel_vector_08.cc
+++ b/tests/lac/parallel_vector_08.cc
@@@ -66,9 -66,9 +66,9 @@@ test(
    // v has ghosts, w has none. set some entries
    // on w, copy into v and check if they are
    // there
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w(
      local_owned, local_owned, MPI_COMM_WORLD);
  
    // set a few of the local elements
diff --cc tests/lac/parallel_vector_10.cc
index 0e13b1b21f,6b250c3292..715a646fbd
--- a/tests/lac/parallel_vector_10.cc
+++ b/tests/lac/parallel_vector_10.cc
@@@ -47,9 -47,9 +47,9 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v);
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v);
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w(v);
  
    // set local values and check them
    LinearAlgebra::ReadWriteVector<double> rw_vector(local_owned);
diff --cc tests/lac/parallel_vector_11.cc
index 2185db2980,a46a73426e..df0285c862
--- a/tests/lac/parallel_vector_11.cc
+++ b/tests/lac/parallel_vector_11.cc
@@@ -55,10 -55,10 +55,10 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_index(2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
    AssertDimension(static_cast<unsigned int>(actual_local_size), v.local_size());
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v), x(v),
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v), x(v),
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w(v), x(v),
      y(v);
  
    // set local elements
diff --cc tests/lac/parallel_vector_12.cc
index 0cdea76c30,9103d9ef77..97f83361f8
--- a/tests/lac/parallel_vector_12.cc
+++ b/tests/lac/parallel_vector_12.cc
@@@ -61,7 -61,7 +61,7 @@@ test(
    if (numproc > 2)
      local_relevant0.add_index(8);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v0(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v0(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v0(
      local_owned0, local_relevant0, MPI_COMM_WORLD);
  
    // vector1: local size 4
@@@ -83,7 -83,7 +83,7 @@@
        local_relevant1.add_index(10);
      }
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v1(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v1(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v1(
      local_owned1, local_relevant1, MPI_COMM_WORLD);
  
    v0 = 1;
@@@ -179,7 -179,7 +179,7 @@@
      deallog << "Ghost values after re-set OK" << std::endl;
  
    // swap with an empty vector
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v2;
    v2.swap(v0);
    AssertDimension(v0.size(), 0);
    AssertDimension(v2.size(), global_size1);
diff --cc tests/lac/parallel_vector_13.cc
index eafc5552b5,c157244982..5249faa317
--- a/tests/lac/parallel_vector_13.cc
+++ b/tests/lac/parallel_vector_13.cc
@@@ -49,7 -49,7 +49,7 @@@ test(
    local_relevant = local_owned;
    local_relevant.add_range(1, 2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values and check them
diff --cc tests/lac/parallel_vector_14.cc
index 4ff6f1bbec,c3fab28bad..2d568e8cc1
--- a/tests/lac/parallel_vector_14.cc
+++ b/tests/lac/parallel_vector_14.cc
@@@ -55,7 -55,7 +55,7 @@@ test(
    // and once where they have not
    for (unsigned int run = 0; run < 2; ++run)
      {
-       LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++      LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
          local_owned, local_relevant, MPI_COMM_WORLD);
  
        // set local values
@@@ -68,8 -68,7 +68,8 @@@
  
        v.import(rw_vector, VectorOperation::insert);
  
-       LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v),
 -      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v), u(v);
++      LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w(v),
 +        u(v);
        u = 0;
  
        v *= 2.0;
@@@ -91,7 -90,7 +91,7 @@@
          }
  
        // copy vector content to non-ghosted vectors, manually created.
-       LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v_dist(
 -      LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v_dist(
++      LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v_dist(
          local_owned, MPI_COMM_WORLD),
          w_dist(v_dist), u_dist(v_dist);
  
diff --cc tests/lac/parallel_vector_15.cc
index 2246d321a2,050c320fed..4233a9d88f
--- a/tests/lac/parallel_vector_15.cc
+++ b/tests/lac/parallel_vector_15.cc
@@@ -51,7 -51,7 +51,7 @@@ test(
    if (numproc > 1)
      local_relevant.add_range(3, 4);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
  
    // set local values
@@@ -67,7 -67,7 +67,7 @@@
    if (myid == 0)
      deallog << "v has ghost elements: " << v.has_ghost_elements() << std::endl;
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w, x;
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w, x;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> w, x;
    w = v;
    if (myid == 0)
      deallog << "w has ghost elements: " << w.has_ghost_elements() << std::endl;
diff --cc tests/lac/parallel_vector_16.cc
index 21e0f5cc23,68e71da766..23670af241
--- a/tests/lac/parallel_vector_16.cc
+++ b/tests/lac/parallel_vector_16.cc
@@@ -52,14 -58,11 +52,14 @@@ test(
    local_relevant.add_range(min_index + 38, min_index + 40);
    local_relevant.add_range(min_index + 41, min_index + 43);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
-   Kokkos::View<double *, MemorySpace::Device::kokkos_space> v_view(
++  Kokkos::View<double *, MemorySpace::Default::kokkos_space> v_view(
 +    v.get_values(), local_relevant.n_elements());
  
 -  deallog << "Local range of proc 0: " << v.local_range().first << " "
 -          << v.local_range().second << std::endl;
 +  deallog << "Local range of proc 0: "
 +          << v.get_partitioner()->local_range().first << " "
 +          << v.get_partitioner()->local_range().second << std::endl;
  
    // set local values
    for (types::global_dof_index i = 0; i < local_size; ++i)
diff --cc tests/lac/parallel_vector_21.cc
index ac37704758,99f867327d..ddedd449b9
--- a/tests/lac/parallel_vector_21.cc
+++ b/tests/lac/parallel_vector_21.cc
@@@ -41,23 -47,17 +41,23 @@@ test(
    ghost_set.add_index(0);
    ghost_set.add_index(2);
  
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      locally_owned, ghost_set, MPI_COMM_WORLD);
  
    // create vector without actually setting the entries since they will be
    // overwritten soon anyway
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v2;
    v2.reinit(v, true);
  
    // set locally owned range of v2 manually
-   Kokkos::View<double *, MemorySpace::Device::kokkos_space> v2_view(
 -  set<<<1, v2.local_size()>>>(v2.get_values());
 -
++  Kokkos::View<double *, MemorySpace::Default::kokkos_space> v2_view(
 +    v2.get_values(), v2.local_size());
 +  Kokkos::deep_copy(v2_view, 1.);
 +
 +  Kokkos::parallel_for(
 +    v2.local_size(), KOKKOS_LAMBDA(int i) {
 +      KOKKOS_IMPL_DO_NOT_USE_PRINTF("%d: %f\n", i, v2_view(i));
 +    });
    // add entries to ghost values
    // Because of limitation in import, the IndexSet of the ReadWriteVector needs
    // to have the local elements.
diff --cc tests/lac/parallel_vector_22.cc
index 1697623d11,6456af6e9b..cfde38daf6
--- a/tests/lac/parallel_vector_22.cc
+++ b/tests/lac/parallel_vector_22.cc
@@@ -57,10 -64,8 +57,10 @@@ test(
    local_relevant.add_range(1, 2);
  
    // create vector
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
-   Kokkos::View<double *, MemorySpace::Device::kokkos_space> v_device(
++  Kokkos::View<double *, MemorySpace::Default::kokkos_space> v_device(
 +    v.get_values(), local_relevant.n_elements());
    const auto &partitioner = v.get_partitioner();
  
    // set local values
diff --cc tests/lac/parallel_vector_23.cc
index 16bcfa90e6,a9bb9c5d9c..064e50e712
--- a/tests/lac/parallel_vector_23.cc
+++ b/tests/lac/parallel_vector_23.cc
@@@ -58,7 -65,7 +58,7 @@@ test(
    local_relevant.add_range(1, 2);
  
    // create vector
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> v(
      local_owned, local_relevant, MPI_COMM_WORLD);
    const auto &partitioner = v.get_partitioner();
  
diff --cc tests/lac/parallel_vector_25.cc
index f4ae375c78,88f3ca0f5e..91858b67e8
--- a/tests/lac/parallel_vector_25.cc
+++ b/tests/lac/parallel_vector_25.cc
@@@ -53,15 -53,15 +53,15 @@@ test(
    vec_ref.compress(VectorOperation::insert);
  
    auto partitioner = vec_ref.get_partitioner();
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device> vec_dev(
 -  LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> vec_dev(
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default> vec_dev(
      partitioner);
    LinearAlgebra::distributed::Vector<double, MemorySpace::Host> vec_host(
      partitioner);
  
-   // Assignment from Host to Device
 -  // Assignment from Host to CUDA
++  // Assignment from Host to Default
    vec_dev.import(vec_ref, VectorOperation::insert);
  
-   // Assignment from Device to Host
 -  // Assignment from CUDA to Host
++  // Assignment from Default to Host
    vec_host.import(vec_dev, VectorOperation::insert);
  
    for (unsigned int i = 0; i < ghost_size; ++i)