From: Daniel Arndt <arndtd@ornl.gov>
Date: Thu, 1 Dec 2022 22:31:04 +0000 (+0000)
Subject: Rename values_dev -> values
X-Git-Tag: v9.5.0-rc1~697^2~37
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=a4b66f3904a95e4dedaa87e0d8ac15cb69dce3cf;p=dealii.git

Rename values_dev -> values
---

diff --git a/include/deal.II/base/memory_space_data.h b/include/deal.II/base/memory_space_data.h
index cc6d03a182..dc9e6d12ea 100644
--- a/include/deal.II/base/memory_space_data.h
+++ b/include/deal.II/base/memory_space_data.h
@@ -50,7 +50,7 @@ namespace MemorySpace
     }
 
     /**
-     * Copy the active data (values for Host and values_dev for Device) to @p begin.
+     * Copy the active data (values for Host and values for Device) to @p begin.
      * If the data is on the device it is moved to the host.
      */
     void
@@ -62,7 +62,7 @@ namespace MemorySpace
 
     /**
      * Copy the data in @p begin to the active data of the structure (values for
-     * Host and values_dev for Device). The pointer @p begin must be on the host.
+     * Host and values for Device). The pointer @p begin must be on the host.
      */
     void
     copy_from(T *begin, std::size_t n_elements)
@@ -71,15 +71,10 @@ namespace MemorySpace
       (void)n_elements;
     }
 
-    /**
-     * Pointer to data on the host.
-     */
-    Kokkos::View<T *, Kokkos::HostSpace> values;
-
     /**
      * Pointer to data on the device.
      */
-    Kokkos::View<T *, typename MemorySpace::kokkos_space> values_dev;
+    Kokkos::View<T *, typename MemorySpace::kokkos_space> values;
 
     /**
      * Pointer to data on the host. The pointer points to the same data as
@@ -121,14 +116,14 @@ namespace MemorySpace
     using MemorySpace = Host;
 
     MemorySpaceData()
-    : values_dev((dealii::Impl::ensure_kokkos_initialized(),
+    : values((dealii::Impl::ensure_kokkos_initialized(),
               Kokkos::View<T *, Kokkos::HostSpace>("host data", 0)))
     {}
 
     void
     copy_to(T *begin, std::size_t n_elements)
     {
-      Assert(n_elements <= values_dev.extent(0),
+      Assert(n_elements <= values.extent(0),
              ExcMessage("n_elements greater than the size of values."));
       using ExecutionSpace = typename MemorySpace::kokkos_space::execution_space;
       Kokkos::
@@ -137,14 +132,14 @@ namespace MemorySpace
       Kokkos::deep_copy(
         ExecutionSpace{},
         begin_view,
-        Kokkos::subview(values_dev, Kokkos::make_pair(std::size_t(0), n_elements)));
+        Kokkos::subview(values, Kokkos::make_pair(std::size_t(0), n_elements)));
       ExecutionSpace{}.fence();
     }
 
     void
     copy_from(T *begin, std::size_t n_elements)
     {
-      Assert(n_elements <= values_dev.extent(0),
+      Assert(n_elements <= values.extent(0),
              ExcMessage("n_elements greater than the size of values."));
       using ExecutionSpace = typename MemorySpace::kokkos_space::execution_space;
       Kokkos::View<const T *,
@@ -153,7 +148,7 @@ namespace MemorySpace
         begin_view(begin, n_elements);
       Kokkos::deep_copy(
         ExecutionSpace{},
-        Kokkos::subview(values_dev, Kokkos::make_pair(std::size_t(0), n_elements)),
+        Kokkos::subview(values, Kokkos::make_pair(std::size_t(0), n_elements)),
         begin_view);
       ExecutionSpace{}.fence();
     }
@@ -161,7 +156,7 @@ namespace MemorySpace
     // unused
     Kokkos::View<T *, Kokkos::HostSpace> values_host_buffer;
 
-    Kokkos::View<T *, typename MemorySpace::kokkos_space> values_dev;
+    Kokkos::View<T *, typename MemorySpace::kokkos_space> values;
 
     std::shared_ptr<T> values_sm_ptr;
 
@@ -188,7 +183,7 @@ namespace MemorySpace
     MemorySpaceData()
     : values_host_buffer((dealii::Impl::ensure_kokkos_initialized(),
               Kokkos::View<T *, Kokkos::HostSpace>("host data", 0)))
-    , values_dev(Kokkos::View<T *, typename MemorySpace::kokkos_space>(
+    , values(Kokkos::View<T *, typename MemorySpace::kokkos_space>(
         "memoryspace data",
         0))
     {}
@@ -196,7 +191,7 @@ namespace MemorySpace
     void
     copy_to(T *begin, std::size_t n_elements)
     {
-      Assert(n_elements <= values_dev.extent(0),
+      Assert(n_elements <= values.extent(0),
              ExcMessage("n_elements greater than the size of values."));
       using ExecutionSpace = typename MemorySpace::kokkos_space::execution_space;
       Kokkos::
@@ -205,14 +200,14 @@ namespace MemorySpace
       Kokkos::deep_copy(
         ExecutionSpace{},
         begin_view,
-        Kokkos::subview(values_dev, Kokkos::make_pair(std::size_t(0), n_elements)));
+        Kokkos::subview(values, Kokkos::make_pair(std::size_t(0), n_elements)));
       ExecutionSpace{}.fence();
     }
 
     void
     copy_from(T *begin, std::size_t n_elements)
     {
-      Assert(n_elements <= values_dev.extent(0),
+      Assert(n_elements <= values.extent(0),
              ExcMessage("n_elements greater than the size of values."));
       using ExecutionSpace = typename MemorySpace::kokkos_space::execution_space;
       Kokkos::View<const T *,
@@ -221,14 +216,14 @@ namespace MemorySpace
         begin_view(begin, n_elements);
       Kokkos::deep_copy(
         ExecutionSpace{},
-        Kokkos::subview(values_dev, Kokkos::make_pair(std::size_t(0), n_elements)),
+        Kokkos::subview(values, Kokkos::make_pair(std::size_t(0), n_elements)),
         begin_view);
       ExecutionSpace{}.fence();
     }
 
     Kokkos::View<T *, Kokkos::HostSpace> values_host_buffer;
     
-    Kokkos::View<T *, typename MemorySpace::kokkos_space> values_dev;
+    Kokkos::View<T *, typename MemorySpace::kokkos_space> values;
     
     // unused
     std::shared_ptr<T> values_sm_ptr;
@@ -244,7 +239,7 @@ namespace MemorySpace
   swap(MemorySpaceData<T, Device> &u, MemorySpaceData<T, Device> &v)
   {
     std::swap(u.values, v.values);
-    std::swap(u.values_dev, v.values_dev);
+    std::swap(u.values_host_buffer, v.values_host_buffer);
   }
 
 #endif
diff --git a/include/deal.II/lac/la_parallel_vector.h b/include/deal.II/lac/la_parallel_vector.h
index 09bcefad0d..a4fb949a74 100644
--- a/include/deal.II/lac/la_parallel_vector.h
+++ b/include/deal.II/lac/la_parallel_vector.h
@@ -1474,7 +1474,7 @@ namespace LinearAlgebra
           begin(::dealii::MemorySpace::
                   MemorySpaceData<Number, ::dealii::MemorySpace::Host> &data)
         {
-          return data.values_dev.data();
+          return data.values.data();
         }
 
         static inline
@@ -1482,14 +1482,14 @@ namespace LinearAlgebra
           begin(const ::dealii::MemorySpace::
                   MemorySpaceData<Number, ::dealii::MemorySpace::Host> &data)
         {
-          return data.values_dev.data();
+          return data.values.data();
         }
 
         static inline Number *
         get_values(::dealii::MemorySpace::
                      MemorySpaceData<Number, ::dealii::MemorySpace::Host> &data)
         {
-          return data.values_dev.data();
+          return data.values.data();
         }
       };
 
@@ -1503,7 +1503,7 @@ namespace LinearAlgebra
           begin(::dealii::MemorySpace::
                   MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
         {
-          return data.values_dev.data();
+          return data.values.data();
         }
 
         static inline
@@ -1511,14 +1511,14 @@ namespace LinearAlgebra
           begin(const ::dealii::MemorySpace::
                   MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
         {
-          return data.values_dev.data();
+          return data.values.data();
         }
 
         static inline Number *
         get_values(::dealii::MemorySpace::
                      MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
         {
-          return data.values_dev.data();
+          return data.values.data();
         }
       };
     } // namespace internal
@@ -1652,7 +1652,7 @@ namespace LinearAlgebra
                vector_is_ghosted == true,
              ExcMessage("You tried to read a ghost element of this vector, "
                         "but it has not imported its ghost values."));
-      return data.values_dev[partitioner->global_to_local(global_index)];
+      return data.values[partitioner->global_to_local(global_index)];
     }
 
 
@@ -1679,7 +1679,7 @@ namespace LinearAlgebra
       // (then, the compiler picks this method according to the C++ rule book
       // even if a human would pick the const method when this subsequent use
       // is just a read)
-      return data.values_dev[partitioner->global_to_local(global_index)];
+      return data.values[partitioner->global_to_local(global_index)];
     }
 
 
@@ -1718,7 +1718,7 @@ namespace LinearAlgebra
              ExcMessage("You tried to read a ghost element of this vector, "
                         "but it has not imported its ghost values."));
 
-      return data.values_dev[local_index];
+      return data.values[local_index];
     }
 
 
@@ -1735,7 +1735,7 @@ namespace LinearAlgebra
                        partitioner->locally_owned_size() +
                          partitioner->n_ghost_indices());
 
-      return data.values_dev[local_index];
+      return data.values[local_index];
     }
 
 
diff --git a/include/deal.II/lac/la_parallel_vector.templates.h b/include/deal.II/lac/la_parallel_vector.templates.h
index 78298e8e5b..0a08efb5e1 100644
--- a/include/deal.II/lac/la_parallel_vector.templates.h
+++ b/include/deal.II/lac/la_parallel_vector.templates.h
@@ -135,12 +135,12 @@ namespace LinearAlgebra
         {
           if (comm_shared == MPI_COMM_SELF)
             {
-	      Kokkos::resize(data.values_dev, new_alloc_size);    	    
+	      Kokkos::resize(data.values, new_alloc_size);    	    
 
               allocated_size = new_alloc_size;
 
               data.values_sm = {
-                ArrayView<const Number>(data.values_dev.data(), new_alloc_size)};
+                ArrayView<const Number>(data.values.data(), new_alloc_size)};
             }
           else
             {
@@ -224,7 +224,7 @@ namespace LinearAlgebra
                 data.values_sm[i] =
                   ArrayView<const Number>(others[i], new_alloc_sizes[i]);
 
-              data.values_dev =
+              data.values =
                 Kokkos::View<Number *,
                              Kokkos::HostSpace,
                              Kokkos::MemoryTraits<Kokkos::Unmanaged>>(
@@ -312,7 +312,7 @@ namespace LinearAlgebra
         {
           for (size_type i = 0; i < size; ++i)
             max =
-              std::max(numbers::NumberTraits<Number>::abs(data.values_dev[i]), max);
+              std::max(numbers::NumberTraits<Number>::abs(data.values[i]), max);
         }
       };
 
@@ -339,17 +339,17 @@ namespace LinearAlgebra
 
           if (new_alloc_size > allocated_size)
             {
-              Assert(((allocated_size > 0 && data.values_dev.size() != 0) ||
-                      data.values_dev.size() == 0),
+              Assert(((allocated_size > 0 && data.values.size() != 0) ||
+                      data.values.size() == 0),
                      ExcInternalError());
 
-	      Kokkos::resize(data.values_dev, new_alloc_size);
+	      Kokkos::resize(data.values, new_alloc_size);
 
               allocated_size = new_alloc_size;
             }
           else if (new_alloc_size == 0)
             {
-              Kokkos::resize(data.values_dev, 0);
+              Kokkos::resize(data.values, 0);
               allocated_size = 0;
             }
         }
@@ -423,14 +423,14 @@ namespace LinearAlgebra
             ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_permutated<
               Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
               indices_dev,
-              data.values_dev.data(),
+              data.values.data(),
               tmp_vector.begin(),
               tmp_n_elements);
           else
             ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<
               Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
               indices_dev,
-              data.values_dev.data(),
+              data.values.data(),
               tmp_vector.begin(),
               tmp_n_elements);
 
@@ -460,7 +460,7 @@ namespace LinearAlgebra
             Number,
             ::dealii::LinearAlgebra::CUDAWrappers::kernel::LInfty<Number>>
             <<<dim3(n_blocks, 1), dim3(::dealii::CUDAWrappers::block_size)>>>(
-              result_device, data.values_dev.data(), size);
+              result_device, data.values.data(), size);
 
           // Copy the result back to the host
           error_code = cudaMemcpy(&result,
@@ -529,7 +529,7 @@ namespace LinearAlgebra
 
       // delete previous content in import data
       Kokkos::resize(import_data.values_host_buffer, 0);
-      Kokkos::resize(import_data.values_dev, 0);
+      Kokkos::resize(import_data.values, 0);
 
       // set partitioner to serial version
       partitioner = std::make_shared<Utilities::MPI::Partitioner>(size);
@@ -560,7 +560,7 @@ namespace LinearAlgebra
 
       // delete previous content in import data
       Kokkos::resize(import_data.values_host_buffer, 0);
-      Kokkos::resize(import_data.values_dev, 0);
+      Kokkos::resize(import_data.values, 0);
 
       // create partitioner
       partitioner = std::make_shared<Utilities::MPI::Partitioner>(local_size,
@@ -606,7 +606,7 @@ namespace LinearAlgebra
       // update_ghost_values, and we might have vectors where we never
       // call these methods and hence do not need to have the storage.
       Kokkos::resize(import_data.values_host_buffer, 0);
-      Kokkos::resize(import_data.values_dev, 0);
+      Kokkos::resize(import_data.values, 0);
 
       thread_loop_partitioner = v.thread_loop_partitioner;
     }
@@ -669,7 +669,7 @@ namespace LinearAlgebra
       // update_ghost_values, and we might have vectors where we never
       // call these methods and hence do not need to have the storage.
       Kokkos::resize(import_data.values_host_buffer, 0);
-      Kokkos::resize(import_data.values_dev, 0);
+      Kokkos::resize(import_data.values, 0);
 
       vector_is_ghosted = false;
     }
@@ -930,13 +930,13 @@ namespace LinearAlgebra
     void
     Vector<Number, MemorySpaceType>::zero_out_ghost_values() const
     {
-      if (data.values_dev.size() != 0)
+      if (data.values.size() != 0)
       {
 #ifdef DEAL_II_COMPILER_CUDA_AWARE
         if (std::is_same_v<MemorySpaceType, MemorySpace::CUDA>)
         {
           const cudaError_t cuda_error_code =
-            cudaMemset(data.values_dev.data() +
+            cudaMemset(data.values.data() +
                          partitioner->locally_owned_size(),
                        0,
                        partitioner->n_ghost_indices() * sizeof(Number));
@@ -945,7 +945,7 @@ namespace LinearAlgebra
 	else
 #endif
 	{
-          std::fill_n(data.values_dev.data() + partitioner->locally_owned_size(),
+          std::fill_n(data.values.data() + partitioner->locally_owned_size(),
                     partitioner->n_ghost_indices(),
                     Number());
         }
@@ -989,8 +989,8 @@ namespace LinearAlgebra
                 std::is_same<MemorySpaceType, dealii::MemorySpace::Host>::value,
                 "This code path should only be compiled for CUDA-aware-MPI for MemorySpace::Host!");
 #  endif
-              if (import_data.values_dev.size() == 0)
-                Kokkos::resize(import_data.values_dev, partitioner->n_import_indices());
+              if (import_data.values.size() == 0)
+                Kokkos::resize(import_data.values, partitioner->n_import_indices());
             }
         }
 
@@ -1002,7 +1002,7 @@ namespace LinearAlgebra
           // device. We use values to store the elements because the function
           // uses a view of the array and thus we need the data on the host to
           // outlive the scope of the function.
-          data.values_host_buffer = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, data.values_dev);
+          data.values_host_buffer = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, data.values);
           partitioner->import_from_ghosted_array_start(
             operation,
             communication_channel,
@@ -1020,10 +1020,10 @@ namespace LinearAlgebra
             operation,
             communication_channel,
             ArrayView<Number, MemorySpaceType>(
-              data.values_dev.data() + partitioner->locally_owned_size(),
+              data.values.data() + partitioner->locally_owned_size(),
               partitioner->n_ghost_indices()),
             ArrayView<Number, MemorySpaceType>(
-              import_data.values_dev.data(), partitioner->n_import_indices()),
+              import_data.values.data(), partitioner->n_import_indices()),
             compress_requests);
         }
 #else
@@ -1070,7 +1070,7 @@ namespace LinearAlgebra
 	  // The communication is done on the host, so we need to
           // move the data back to the device.
 	  cudaError_t cuda_error_code =
-            cudaMemcpy(data.values_dev.data(),
+            cudaMemcpy(data.values.data(),
                        data.values_host_buffer.data(),
                        allocated_size * sizeof(Number),
                        cudaMemcpyHostToDevice);
@@ -1082,17 +1082,17 @@ namespace LinearAlgebra
 #  endif
         {
           Assert(partitioner->n_import_indices() == 0 ||
-                   import_data.values_dev.size() != 0,
+                   import_data.values.size() != 0,
                  ExcNotInitialized());
           partitioner
             ->import_from_ghosted_array_finish<Number, MemorySpaceType>(
               operation,
               ArrayView<const Number, MemorySpaceType>(
-                import_data.values_dev.data(), partitioner->n_import_indices()),
+                import_data.values.data(), partitioner->n_import_indices()),
               ArrayView<Number, MemorySpaceType>(
-                data.values_dev.data(), partitioner->locally_owned_size()),
+                data.values.data(), partitioner->locally_owned_size()),
               ArrayView<Number, MemorySpaceType>(
-                data.values_dev.data() + partitioner->locally_owned_size(),
+                data.values.data() + partitioner->locally_owned_size(),
                 partitioner->n_ghost_indices()),
               compress_requests);
         }
@@ -1136,8 +1136,8 @@ namespace LinearAlgebra
             std::is_same<MemorySpaceType, dealii::MemorySpace::Host>::value,
             "This code path should only be compiled for CUDA-aware-MPI for MemorySpace::Host!");
 #    endif
-          if (import_data.values_dev.size() == 0)
-            Kokkos::resize(import_data.values_dev, partitioner->n_import_indices());
+          if (import_data.values.size() == 0)
+            Kokkos::resize(import_data.values, partitioner->n_import_indices());
 		     }
         }
 
@@ -1148,7 +1148,7 @@ namespace LinearAlgebra
        // device. We use values to store the elements because the function
        // uses a view of the array and thus we need the data on the host to
        // outlive the scope of the function.
-       data.values_host_buffer = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, data.values_dev);
+       data.values_host_buffer = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, data.values);
        
        partitioner->export_to_ghosted_array_start<Number, MemorySpace::Host>(
         communication_channel,
@@ -1166,11 +1166,11 @@ namespace LinearAlgebra
       partitioner->export_to_ghosted_array_start<Number, MemorySpaceType>(
         communication_channel,
         ArrayView<const Number, MemorySpaceType>(
-          data.values_dev.data(), partitioner->locally_owned_size()),
-        ArrayView<Number, MemorySpaceType>(import_data.values_dev.data(),
+          data.values.data(), partitioner->locally_owned_size()),
+        ArrayView<Number, MemorySpaceType>(import_data.values.data(),
                                              partitioner->n_import_indices()),
         ArrayView<Number, MemorySpaceType>(
-          data.values_dev.data() + partitioner->locally_owned_size(),
+          data.values.data() + partitioner->locally_owned_size(),
           partitioner->n_ghost_indices()),
         update_ghost_values_requests);
  }
@@ -1210,7 +1210,7 @@ namespace LinearAlgebra
 // The communication is done on the host, so we need to
       // move the data back to the device.
  cudaError_t cuda_error_code =
-            cudaMemcpy(data.values_dev.data() +
+            cudaMemcpy(data.values.data() +
                          partitioner->locally_owned_size(),
                        data.values_host_buffer.data() + partitioner->locally_owned_size(),
                        partitioner->n_ghost_indices() * sizeof(Number),
@@ -1223,7 +1223,7 @@ namespace LinearAlgebra
 		{
           partitioner->export_to_ghosted_array_finish(
             ArrayView<Number, MemorySpaceType>(
-              data.values_dev.data() + partitioner->locally_owned_size(),
+              data.values.data() + partitioner->locally_owned_size(),
               partitioner->n_ghost_indices()),
             update_ghost_values_requests);
 		}
@@ -2033,7 +2033,7 @@ namespace LinearAlgebra
       if (partitioner.use_count() > 0)
         memory +=
           partitioner->memory_consumption() / partitioner.use_count() + 1;
-      if (import_data.values_host_buffer.size() != 0 || import_data.values_dev.size() != 0)
+      if (import_data.values_host_buffer.size() != 0 || import_data.values.size() != 0)
         memory += (static_cast<std::size_t>(partitioner->n_import_indices()) *
                    sizeof(Number));
       return memory;
diff --git a/include/deal.II/lac/vector_operations_internal.h b/include/deal.II/lac/vector_operations_internal.h
index 80754c2412..cde2c0b491 100644
--- a/include/deal.II/lac/vector_operations_internal.h
+++ b/include/deal.II/lac/vector_operations_internal.h
@@ -1677,7 +1677,7 @@ namespace internal
         const size_type /*size*/,
         real_type & /*sum*/,
         Number * /*values*/,
-        Number * /*values_dev*/)
+        Number * /*values*/)
       {}
 
       template <typename real_type>
@@ -1734,8 +1734,8 @@ namespace internal
                                                   ::dealii::MemorySpace::Host>
              &data)
       {
-        Vector_copy<Number, Number2> copier(v_data.values_dev.data(),
-                                            data.values_dev.data());
+        Vector_copy<Number, Number2> copier(v_data.values.data(),
+                                            data.values.data());
         parallel_for(copier, 0, size, thread_loop_partitioner);
       }
 
@@ -1748,7 +1748,7 @@ namespace internal
                                                  ::dealii::MemorySpace::Host>
             &data)
       {
-        Vector_set<Number> setter(s, data.values_dev.data());
+        Vector_set<Number> setter(s, data.values.data());
         parallel_for(setter, 0, size, thread_loop_partitioner);
       }
 
@@ -1763,8 +1763,8 @@ namespace internal
                                                ::dealii::MemorySpace::Host>
           &data)
       {
-        Vectorization_add_v<Number> vector_add(data.values_dev.data(),
-                                               v_data.values_dev.data());
+        Vectorization_add_v<Number> vector_add(data.values.data(),
+                                               v_data.values.data());
         parallel_for(vector_add, 0, size, thread_loop_partitioner);
       }
 
@@ -1779,8 +1779,8 @@ namespace internal
                                                ::dealii::MemorySpace::Host>
           &data)
       {
-        Vectorization_subtract_v<Number> vector_subtract(data.values_dev.data(),
-                                                         v_data.values_dev.data());
+        Vectorization_subtract_v<Number> vector_subtract(data.values.data(),
+                                                         v_data.values.data());
         parallel_for(vector_subtract, 0, size, thread_loop_partitioner);
       }
 
@@ -1794,7 +1794,7 @@ namespace internal
                                                ::dealii::MemorySpace::Host>
           &data)
       {
-        Vectorization_add_factor<Number> vector_add(data.values_dev.data(), a);
+        Vectorization_add_factor<Number> vector_add(data.values.data(), a);
         parallel_for(vector_add, 0, size, thread_loop_partitioner);
       }
 
@@ -1809,8 +1809,8 @@ namespace internal
                                                     ::dealii::MemorySpace::Host>
                &data)
       {
-        Vectorization_add_av<Number> vector_add(data.values_dev.data(),
-                                                v_data.values_dev.data(),
+        Vectorization_add_av<Number> vector_add(data.values.data(),
+                                                v_data.values.data(),
                                                 a);
         parallel_for(vector_add, 0, size, thread_loop_partitioner);
       }
@@ -1831,7 +1831,7 @@ namespace internal
           &data)
       {
         Vectorization_add_avpbw<Number> vector_add(
-          data.values_dev.data(), v_data.values_dev.data(), w_data.values_dev.data(), a, b);
+          data.values.data(), v_data.values.data(), w_data.values.data(), a, b);
         parallel_for(vector_add, 0, size, thread_loop_partitioner);
       }
 
@@ -1847,8 +1847,8 @@ namespace internal
                                                ::dealii::MemorySpace::Host>
           &data)
       {
-        Vectorization_sadd_xv<Number> vector_sadd(data.values_dev.data(),
-                                                  v_data.values_dev.data(),
+        Vectorization_sadd_xv<Number> vector_sadd(data.values.data(),
+                                                  v_data.values.data(),
                                                   x);
         parallel_for(vector_sadd, 0, size, thread_loop_partitioner);
       }
@@ -1866,8 +1866,8 @@ namespace internal
                                                ::dealii::MemorySpace::Host>
           &data)
       {
-        Vectorization_sadd_xav<Number> vector_sadd(data.values_dev.data(),
-                                                   v_data.values_dev.data(),
+        Vectorization_sadd_xav<Number> vector_sadd(data.values.data(),
+                                                   v_data.values.data(),
                                                    a,
                                                    x);
         parallel_for(vector_sadd, 0, size, thread_loop_partitioner);
@@ -1890,7 +1890,7 @@ namespace internal
           &data)
       {
         Vectorization_sadd_xavbw<Number> vector_sadd(
-          data.values_dev.data(), v_data.values_dev.data(), w_data.values_dev.data(), x, a, b);
+          data.values.data(), v_data.values.data(), w_data.values.data(), x, a, b);
         parallel_for(vector_sadd, 0, size, thread_loop_partitioner);
       }
 
@@ -1904,7 +1904,7 @@ namespace internal
                                                ::dealii::MemorySpace::Host>
           &data)
       {
-        Vectorization_multiply_factor<Number> vector_multiply(data.values_dev.data(),
+        Vectorization_multiply_factor<Number> vector_multiply(data.values.data(),
                                                               factor);
         parallel_for(vector_multiply, 0, size, thread_loop_partitioner);
       }
@@ -1919,8 +1919,8 @@ namespace internal
                                                    ::dealii::MemorySpace::Host>
               &data)
       {
-        Vectorization_scale<Number> vector_scale(data.values_dev.data(),
-                                                 v_data.values_dev.data());
+        Vectorization_scale<Number> vector_scale(data.values.data(),
+                                                 v_data.values.data());
         parallel_for(vector_scale, 0, size, thread_loop_partitioner);
       }
 
@@ -1935,8 +1935,8 @@ namespace internal
                                                     ::dealii::MemorySpace::Host>
                &data)
       {
-        Vectorization_equ_au<Number> vector_equ(data.values_dev.data(),
-                                                v_data.values_dev.data(),
+        Vectorization_equ_au<Number> vector_equ(data.values.data(),
+                                                v_data.values.data(),
                                                 a);
         parallel_for(vector_equ, 0, size, thread_loop_partitioner);
       }
@@ -1957,7 +1957,7 @@ namespace internal
           &data)
       {
         Vectorization_equ_aubv<Number> vector_equ(
-          data.values_dev.data(), v_data.values_dev.data(), w_data.values_dev.data(), a, b);
+          data.values.data(), v_data.values.data(), w_data.values.data(), a, b);
         parallel_for(vector_equ, 0, size, thread_loop_partitioner);
       }
 
@@ -1973,7 +1973,7 @@ namespace internal
       {
         Number                                                   sum;
         dealii::internal::VectorOperations::Dot<Number, Number2> dot(
-          data.values_dev.data(), v_data.values_dev.data());
+          data.values.data(), v_data.values.data());
         dealii::internal::VectorOperations::parallel_reduce(
           dot, 0, size, sum, thread_loop_partitioner);
         AssertIsFinite(sum);
@@ -1991,7 +1991,7 @@ namespace internal
                                                     ::dealii::MemorySpace::Host>
                &data)
       {
-        Norm2<Number, real_type> norm2(data.values_dev.data());
+        Norm2<Number, real_type> norm2(data.values.data());
         parallel_reduce(norm2, 0, size, sum, thread_loop_partitioner);
       }
 
@@ -2004,7 +2004,7 @@ namespace internal
           MemorySpaceData<Number, ::dealii::MemorySpace::Host> &data)
       {
         Number            sum;
-        MeanValue<Number> mean(data.values_dev.data());
+        MeanValue<Number> mean(data.values.data());
         parallel_reduce(mean, 0, size, sum, thread_loop_partitioner);
 
         return sum;
@@ -2020,7 +2020,7 @@ namespace internal
                                                     ::dealii::MemorySpace::Host>
                &data)
       {
-        Norm1<Number, real_type> norm1(data.values_dev.data());
+        Norm1<Number, real_type> norm1(data.values.data());
         parallel_reduce(norm1, 0, size, sum, thread_loop_partitioner);
       }
 
@@ -2035,7 +2035,7 @@ namespace internal
                                                     ::dealii::MemorySpace::Host>
                &data)
       {
-        NormP<Number, real_type> normp(data.values_dev.data(), p);
+        NormP<Number, real_type> normp(data.values.data(), p);
         parallel_reduce(normp, 0, size, sum, thread_loop_partitioner);
       }
 
@@ -2054,9 +2054,9 @@ namespace internal
           &data)
       {
         Number            sum;
-        AddAndDot<Number> adder(data.values_dev.data(),
-                                v_data.values_dev.data(),
-                                w_data.values_dev.data(),
+        AddAndDot<Number> adder(data.values.data(),
+                                v_data.values.data(),
+                                w_data.values.data(),
                                 a);
         parallel_reduce(adder, 0, size, sum, thread_loop_partitioner);
 
@@ -2112,8 +2112,8 @@ namespace internal
       {
         if (operation == VectorOperation::insert)
           {
-            cudaError_t cuda_error_code = cudaMemcpy(data.values_dev.data(),
-                                                     v_data.values_dev.data(),
+            cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
+                                                     v_data.values.data(),
                                                      size * sizeof(Number),
                                                      cudaMemcpyDeviceToHost);
             AssertCuda(cuda_error_code);
@@ -2147,8 +2147,8 @@ namespace internal
                                                ::dealii::MemorySpace::CUDA>
           &data)
       {
-        cudaError_t cuda_error_code = cudaMemcpy(data.values_dev.data(),
-                                                 v_data.values_dev.data(),
+        cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
+                                                 v_data.values.data(),
                                                  size * sizeof(Number),
                                                  cudaMemcpyDeviceToDevice);
         AssertCuda(cuda_error_code);
@@ -2164,7 +2164,7 @@ namespace internal
       {
         const int n_blocks = 1 + size / (chunk_size * block_size);
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::set<Number>
-          <<<n_blocks, block_size>>>(data.values_dev.data(), s, size);
+          <<<n_blocks, block_size>>>(data.values.data(), s, size);
         AssertCudaKernel();
       }
 
@@ -2180,9 +2180,9 @@ namespace internal
       {
         const int n_blocks = 1 + size / (chunk_size * block_size);
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
-          <<<n_blocks, block_size>>>(data.values_dev.data(),
+          <<<n_blocks, block_size>>>(data.values.data(),
                                      1.,
-                                     v_data.values_dev.data(),
+                                     v_data.values.data(),
                                      size);
         AssertCudaKernel();
       }
@@ -2199,9 +2199,9 @@ namespace internal
       {
         const int n_blocks = 1 + size / (chunk_size * block_size);
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
-          <<<n_blocks, block_size>>>(data.values_dev.data(),
+          <<<n_blocks, block_size>>>(data.values.data(),
                                      -1.,
-                                     v_data.values_dev.data(),
+                                     v_data.values.data(),
                                      size);
         AssertCudaKernel();
       }
@@ -2217,7 +2217,7 @@ namespace internal
       {
         const int n_blocks = 1 + size / (chunk_size * block_size);
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_add<Number>
-          <<<n_blocks, block_size>>>(data.values_dev.data(), a, size);
+          <<<n_blocks, block_size>>>(data.values.data(), a, size);
         AssertCudaKernel();
       }
 
@@ -2234,9 +2234,9 @@ namespace internal
       {
         const int n_blocks = 1 + size / (chunk_size * block_size);
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
-          <<<n_blocks, block_size>>>(data.values_dev.data(),
+          <<<n_blocks, block_size>>>(data.values.data(),
                                      a,
-                                     v_data.values_dev.data(),
+                                     v_data.values.data(),
                                      size);
         AssertCudaKernel();
       }
@@ -2257,11 +2257,11 @@ namespace internal
       {
         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_dev.data(),
+          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
                                                     a,
-                                                    v_data.values_dev.data(),
+                                                    v_data.values.data(),
                                                     b,
-                                                    w_data.values_dev.data(),
+                                                    w_data.values.data(),
                                                     size);
         AssertCudaKernel();
       }
@@ -2280,7 +2280,7 @@ namespace internal
         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_dev.data(), 1., v_data.values_dev.data(), size);
+            x, data.values.data(), 1., v_data.values.data(), size);
         AssertCudaKernel();
       }
 
@@ -2299,7 +2299,7 @@ namespace internal
         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_dev.data(), a, v_data.values_dev.data(), size);
+            x, data.values.data(), a, v_data.values.data(), size);
         AssertCudaKernel();
       }
 
@@ -2321,11 +2321,11 @@ namespace internal
         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_dev.data(),
+                                                    data.values.data(),
                                                     a,
-                                                    v_data.values_dev.data(),
+                                                    v_data.values.data(),
                                                     b,
-                                                    w_data.values_dev.data(),
+                                                    w_data.values.data(),
                                                     size);
         AssertCudaKernel();
       }
@@ -2341,7 +2341,7 @@ namespace internal
       {
         const int n_blocks = 1 + size / (chunk_size * block_size);
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_scale<Number>
-          <<<n_blocks, block_size>>>(data.values_dev.data(), factor, size);
+          <<<n_blocks, block_size>>>(data.values.data(), factor, size);
         AssertCudaKernel();
       }
 
@@ -2357,8 +2357,8 @@ namespace internal
       {
         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_dev.data(),
-                                                    v_data.values_dev.data(),
+          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
+                                                    v_data.values.data(),
                                                     size);
         AssertCudaKernel();
       }
@@ -2376,9 +2376,9 @@ namespace internal
       {
         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_dev.data(),
+          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
                                                     a,
-                                                    v_data.values_dev.data(),
+                                                    v_data.values.data(),
                                                     size);
         AssertCudaKernel();
       }
@@ -2399,11 +2399,11 @@ namespace internal
       {
         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_dev.data(),
+          <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
                                                     a,
-                                                    v_data.values_dev.data(),
+                                                    v_data.values.data(),
                                                     b,
-                                                    w_data.values_dev.data(),
+                                                    w_data.values.data(),
                                                     size);
         AssertCudaKernel();
       }
@@ -2428,8 +2428,8 @@ namespace internal
           Number,
           ::dealii::LinearAlgebra::CUDAWrappers::kernel::DotProduct<Number>>
           <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
-                                                    data.values_dev.data(),
-                                                    v_data.values_dev.data(),
+                                                    data.values.data(),
+                                                    v_data.values.data(),
                                                     static_cast<unsigned int>(
                                                       size));
         AssertCudaKernel();
@@ -2480,7 +2480,7 @@ namespace internal
           Number,
           ::dealii::LinearAlgebra::CUDAWrappers::kernel::ElemSum<Number>>
           <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
-                                                    data.values_dev.data(),
+                                                    data.values.data(),
                                                     size);
 
         // Copy the result back to the host
@@ -2517,7 +2517,7 @@ namespace internal
           Number,
           ::dealii::LinearAlgebra::CUDAWrappers::kernel::L1Norm<Number>>
           <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
-                                                    data.values_dev.data(),
+                                                    data.values.data(),
                                                     size);
 
         // Copy the result back to the host
@@ -2566,9 +2566,9 @@ namespace internal
         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_dev.data(),
-                                                    v_data.values_dev.data(),
-                                                    w_data.values_dev.data(),
+                                                    data.values.data(),
+                                                    v_data.values.data(),
+                                                    w_data.values.data(),
                                                     a,
                                                     size);
 
@@ -2629,8 +2629,8 @@ namespace internal
       {
         if (operation == VectorOperation::insert)
           {
-            cudaError_t cuda_error_code = cudaMemcpy(data.values_dev.data(),
-                                                     v_data.values_dev.data(),
+            cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
+                                                     v_data.values.data(),
                                                      size * sizeof(Number),
                                                      cudaMemcpyHostToDevice);
             AssertCuda(cuda_error_code);