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

author Daniel Arndt <arndtd@ornl.gov>

Tue, 20 Dec 2022 00:31:34 +0000 (19:31 -0500)

committer Daniel Arndt <arndtd@ornl.gov>

Tue, 20 Dec 2022 00:31:34 +0000 (19:31 -0500)
author Daniel Arndt <arndtd@ornl.gov>
Tue, 20 Dec 2022 00:31:34 +0000 (19:31 -0500)
committer Daniel Arndt <arndtd@ornl.gov>
Tue, 20 Dec 2022 00:31:34 +0000 (19:31 -0500)
diff --cc include/deal.II/base/memory_space_data.h

index 564d7afccbe57b1032a8dda9f62e536a09908da7,19c615a628710c2d21fe6922e0df86812cb8ce03..f30bbb32a1a9067679b68ff9c471035dfef5d4fb
--- 1/include/deal.II/base/memory_space_data.h
--- 2/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 e310f207ec4810fae56b6c6d0b8f55934b27b084,21b101d18b6dbdb8390532ef5d5bb5e153d69a80..f0060403e5547e4f35ddc74b58458a686b4ed0cf
--- 1/include/deal.II/lac/la_parallel_vector.templates.h
--- 2/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 08135c2d1033a47e14328c7a60b3e90085e9ca66,d1b9d6c070666daf6b5d31af1fe114ee7b066489..c9a8f3347075f6b0871d5309ef28585aec54327a
--- 1/include/deal.II/lac/read_write_vector.templates.h
--- 2/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 4f2a2b475accbf1350e8daf7fac748432d0a02f7,21d08f4413eaa8790c16a07c52840a31121ea33d..92300d124e6a69438f580111b196f36f073319cf
--- 1/include/deal.II/lac/vector_operations_internal.h
--- 2/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 76dc6c01ebb9dc9e8e2e9a88881b0f0e2f0d4e98,e4ebd6c6f51318305df4254ae38d6a9d1034adad..91fc69a62564ff1f224675623c254fe7b8e53e4a
--- 1/source/lac/la_parallel_vector_cuda.cc
--- 2/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);
-     Vector<float, ::dealii::MemorySpace::Device>::reinit<float>(
-       const Vector<float, ::dealii::MemorySpace::Device> &,
+ +
+ +    template void
-     Vector<double, ::dealii::MemorySpace::Device>::reinit<double>(
-       const Vector<double, ::dealii::MemorySpace::Device> &,
++    Vector<float, ::dealii::MemorySpace::Default>::reinit<float>(
++      const Vector<float, ::dealii::MemorySpace::Default> &,
+ +      const bool);
+ +    template void
++    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 a6dc6febd2d640a7d22797542ae0fc66167cea56,f1ca4fee7ce39a71acb04e44e265a0c18c9ecf93..23ced9c5bf7cfd126c0d147f822e08adb65505a7
--- 1/source/lac/read_write_vector.cc
--- 2/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> &);
-     const distributed::Vector<float, ::dealii::MemorySpace::Device> &,
+ +  template void
+ +  ReadWriteVector<float>::import(
++    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> &);
-     const distributed::Vector<double, ::dealii::MemorySpace::Device> &,
+ +  template void
+ +  ReadWriteVector<double>::import(
++    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 db9ea9dfa2393e56a68d4381575e54d65674061b,461c1cf50c28440efbc24a35a712f2d5bede00bd..5ab7a4c8d683f1628067929248c7d7d745d90df6
--- 1/source/lac/vector_memory.cc
--- 2/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"
-   LinearAlgebra::distributed::Vector<float, MemorySpace::Device>>;
+ +template class VectorMemory<
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<float, MemorySpace::Default>>;
+ +template class VectorMemory<
-   LinearAlgebra::distributed::Vector<float, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default>>;
+ +template class GrowingVectorMemory<
-   LinearAlgebra::distributed::Vector<double, MemorySpace::Device>>;
++  LinearAlgebra::distributed::Vector<float, MemorySpace::Default>>;
+ +template class GrowingVectorMemory<
++  LinearAlgebra::distributed::Vector<double, MemorySpace::Default>>;
   
   namespace internal
   {
@@@ -52,12 -44,6 +52,12 @@@
       release_all_unused_memory()
       {
   #include "vector_memory_release.inst"
-         MemorySpace::Device>>::release_unused_memory();
+ +      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::Default>>::release_unused_memory();
   #ifdef DEAL_II_WITH_CUDA
         release_all_unused_cuda_memory();
   #endif
diff --cc tests/lac/parallel_vector_01.cc

index 60dfbc9f2ba23d16a6296710c5b05533790eba07,5021d1aa16d082e8255c60e1cd5a1903632c403f..cad97db078eefb049544df65efc50f5d1903b8cc
--- 1/tests/lac/parallel_vector_01.cc
--- 2/tests/cuda/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 891f0ee834464afcb866396ef7bd685d252a6d6b,e7664f4e1c158273ffc55d4dc84686e57a854fcf..79e9922ceed1b29cd3658c67bf1013c4a785f0e0
--- 1/tests/lac/parallel_vector_02.cc
--- 2/tests/cuda/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 5d686847c82b2a8850f4bc164927d2e5f53cc569,0e34c7ed54694455ca6dd1bc49f5300b8ffeb928..33145afb4ebffa541adeeffb017a01f51882e36f
--- 1/tests/lac/parallel_vector_03.cc
--- 2/tests/cuda/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 a7ff1fc723e86edceebf3651ced811b1db33bb7d,47f510ab7e9ff903e856c6c6b14d772190a5aaae..190dc37fc1f9854a9d34c2e6d25f233280a171e5
--- 1/tests/lac/parallel_vector_03a.cc
--- 2/tests/cuda/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 a3cb0938547cebf80dc2de56fc4bc92b0875c3e2,e928ed32b176e04ce59ba699ff288d43ea77bc7e..08f2e81776f27ed4bb4cdcf200c413a4782fd013
--- 1/tests/lac/parallel_vector_04.cc
--- 2/tests/cuda/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 9cf4631e2f1944025de6a85f6600bbd2d89fc09e,1519bd97e6934fe48c0cd5a7031c88858d1854ac..005e4e1c17b0b0e04946b5c4d90bfd630140182e
--- 1/tests/lac/parallel_vector_05.cc
--- 2/tests/cuda/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 01b1432a569719aad33ff65a022ed62fd2ca364d,bfe06fffaeff66fe648fadde5dd1cd007411da0d..804248706cc63dacda3543c75febd6572576822f
--- 1/tests/lac/parallel_vector_06.cc
--- 2/tests/cuda/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 c33c5dadf0625b01d2d1b08484a8b81f5caa8c4c,f0c4716232a51f064edf2acc1f9437e95ed7e7b6..b9c4b6cd97d5ef537807daab05d54627da99e002
--- 1/tests/lac/parallel_vector_08.cc
--- 2/tests/cuda/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 0e13b1b21f8045946bfbe3cfd2d366f23da59298,6b250c3292583de315dd4c7b0a7e9b377a23f60f..715a646fbd19d3c5c42f4cdfe59f3309f3b103aa
--- 1/tests/lac/parallel_vector_10.cc
--- 2/tests/cuda/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 2185db298096f9a910df24555c0fb50558bcfc4b,a46a73426e26b5ce9090e66b957e0d108f841382..df0285c862ca1c60dca7e3babd23daaee65d9235
--- 1/tests/lac/parallel_vector_11.cc
--- 2/tests/cuda/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 0cdea76c30a9b1f61101eb549a02f3f7ce2d5071,9103d9ef77dcd1566b902328ef2990ef0a768f2d..97f83361f80fb5a5c73bdc7054e95b0e794b417c
--- 1/tests/lac/parallel_vector_12.cc
--- 2/tests/cuda/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 eafc5552b5dbcf567d9ace4d17dcea3d297ab7f5,c157244982da4a3c4c47cc23a8ebd893243a2fa4..5249faa317161bec491f429cabf6808c35d7d2e6
--- 1/tests/lac/parallel_vector_13.cc
--- 2/tests/cuda/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 4ff6f1bbeca5763974121611a2ff50f74b049bf9,c3fab28bade4d73c1f4aaa827e022c14c18434ea..2d568e8cc15852df7cae636908d6df92d293d5e9
--- 1/tests/lac/parallel_vector_14.cc
--- 2/tests/cuda/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 2246d321a23a03a57f8c7de0b06b914fabeb2dd2,050c320fedf489e4390c06341505db575470b5c8..4233a9d88fee5a6d879a7ccd15c6229cd110aa23
--- 1/tests/lac/parallel_vector_15.cc
--- 2/tests/cuda/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 21e0f5cc2358566024abb528da9d51ab88095a90,68e71da7662fcca71d70be2062a81094c9d4dcd7..23670af24175933187ad5ba87ce7463756c66bc1
--- 1/tests/lac/parallel_vector_16.cc
--- 2/tests/cuda/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 ac37704758e27c6755b01e6ff5d8893592820c4b,99f867327d2b3cb838488d4760c3ebb0a4a99361..ddedd449b9038eaafe16bb383aff993ed5970e32
--- 1/tests/lac/parallel_vector_21.cc
--- 2/tests/cuda/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 1697623d1118672a0d3967ebebba4a4413748f71,6456af6e9b6fbf96175d13701bd3a2c16d0ec4e4..cfde38daf6b4621555fe6de7bafa9b36f0d1b392
--- 1/tests/lac/parallel_vector_22.cc
--- 2/tests/cuda/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 16bcfa90e6a4895c816976478309c401be9deca4,a9bb9c5d9ce18bff86554a2dc23f900bd1233da2..064e50e71256d520196408fca644995eb5ccbc5d
--- 1/tests/lac/parallel_vector_23.cc
--- 2/tests/cuda/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 f4ae375c78960cb98543e306fa70edd84455f677,88f3ca0f5eaa4225550be4d58dbe5e10049e48a2..91858b67e81b0d07f4a77486a1127de88628fd2e
--- 1/tests/lac/parallel_vector_25.cc
--- 2/tests/cuda/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)
author	Daniel Arndt <arndtd@ornl.gov>
	Tue, 20 Dec 2022 00:31:34 +0000 (19:31 -0500)
committer	Daniel Arndt <arndtd@ornl.gov>
	Tue, 20 Dec 2022 00:31:34 +0000 (19:31 -0500)
		1	2
include/deal.II/base/memory_space_data.h	patch \|	diff1 \|	diff2 \|	blob \| history
include/deal.II/lac/la_parallel_vector.templates.h	patch \|	diff1 \|	diff2 \|	blob \| history
include/deal.II/lac/read_write_vector.templates.h	patch \|	diff1 \|	diff2 \|	blob \| history
include/deal.II/lac/vector_operations_internal.h	patch \|	diff1 \|	diff2 \|	blob \| history
source/lac/la_parallel_vector_cuda.cc	patch \|	diff1 \|	diff2 \|	blob \| history
source/lac/read_write_vector.cc	patch \|	diff1 \|	diff2 \|	blob \| history
source/lac/vector_memory.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_01.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_02.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_03.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_03a.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_04.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_05.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_06.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_08.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_10.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_11.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_12.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_13.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_14.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_15.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_16.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_21.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_22.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_23.cc	patch \|	diff1 \|	diff2 \|	blob \| history
tests/lac/parallel_vector_25.cc	patch \|	diff1 \|	diff2 \|	blob \| history