From a4b66f3904a95e4dedaa87e0d8ac15cb69dce3cf Mon Sep 17 00:00:00 2001 From: Daniel Arndt Date: Thu, 1 Dec 2022 22:31:04 +0000 Subject: [PATCH] Rename values_dev -> values --- include/deal.II/base/memory_space_data.h | 37 +++-- include/deal.II/lac/la_parallel_vector.h | 20 +-- .../lac/la_parallel_vector.templates.h | 74 +++++----- .../deal.II/lac/vector_operations_internal.h | 136 +++++++++--------- 4 files changed, 131 insertions(+), 136 deletions(-) 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 values; - /** * Pointer to data on the device. */ - Kokkos::View values_dev; + Kokkos::View 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("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 values_host_buffer; - Kokkos::View values_dev; + Kokkos::View values; std::shared_ptr values_sm_ptr; @@ -188,7 +183,7 @@ namespace MemorySpace MemorySpaceData() : values_host_buffer((dealii::Impl::ensure_kokkos_initialized(), Kokkos::View("host data", 0))) - , values_dev(Kokkos::View( + , values(Kokkos::View( "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 values_host_buffer; - Kokkos::View values_dev; + Kokkos::View values; // unused std::shared_ptr values_sm_ptr; @@ -244,7 +239,7 @@ namespace MemorySpace swap(MemorySpaceData &u, MemorySpaceData &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 &data) { - return data.values_dev.data(); + return data.values.data(); } static inline @@ -1482,14 +1482,14 @@ namespace LinearAlgebra begin(const ::dealii::MemorySpace:: MemorySpaceData &data) { - return data.values_dev.data(); + return data.values.data(); } static inline Number * get_values(::dealii::MemorySpace:: MemorySpaceData &data) { - return data.values_dev.data(); + return data.values.data(); } }; @@ -1503,7 +1503,7 @@ namespace LinearAlgebra begin(::dealii::MemorySpace:: MemorySpaceData &data) { - return data.values_dev.data(); + return data.values.data(); } static inline @@ -1511,14 +1511,14 @@ namespace LinearAlgebra begin(const ::dealii::MemorySpace:: MemorySpaceData &data) { - return data.values_dev.data(); + return data.values.data(); } static inline Number * get_values(::dealii::MemorySpace:: MemorySpaceData &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(data.values_dev.data(), new_alloc_size)}; + ArrayView(data.values.data(), new_alloc_size)}; } else { @@ -224,7 +224,7 @@ namespace LinearAlgebra data.values_sm[i] = ArrayView(others[i], new_alloc_sizes[i]); - data.values_dev = + data.values = Kokkos::View>( @@ -312,7 +312,7 @@ namespace LinearAlgebra { for (size_type i = 0; i < size; ++i) max = - std::max(numbers::NumberTraits::abs(data.values_dev[i]), max); + std::max(numbers::NumberTraits::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><<>>( indices_dev, - data.values_dev.data(), + data.values.data(), tmp_vector.begin(), tmp_n_elements); else ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated< Number><<>>( 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> <<>>( - 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(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(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::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) { 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::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( - data.values_dev.data() + partitioner->locally_owned_size(), + data.values.data() + partitioner->locally_owned_size(), partitioner->n_ghost_indices()), ArrayView( - 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( operation, ArrayView( - import_data.values_dev.data(), partitioner->n_import_indices()), + import_data.values.data(), partitioner->n_import_indices()), ArrayView( - data.values_dev.data(), partitioner->locally_owned_size()), + data.values.data(), partitioner->locally_owned_size()), ArrayView( - 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::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( communication_channel, @@ -1166,11 +1166,11 @@ namespace LinearAlgebra partitioner->export_to_ghosted_array_start( communication_channel, ArrayView( - data.values_dev.data(), partitioner->locally_owned_size()), - ArrayView(import_data.values_dev.data(), + data.values.data(), partitioner->locally_owned_size()), + ArrayView(import_data.values.data(), partitioner->n_import_indices()), ArrayView( - 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( - 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(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 @@ -1734,8 +1734,8 @@ namespace internal ::dealii::MemorySpace::Host> &data) { - Vector_copy copier(v_data.values_dev.data(), - data.values_dev.data()); + Vector_copy 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 setter(s, data.values_dev.data()); + Vector_set 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 vector_add(data.values_dev.data(), - v_data.values_dev.data()); + Vectorization_add_v 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 vector_subtract(data.values_dev.data(), - v_data.values_dev.data()); + Vectorization_subtract_v 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 vector_add(data.values_dev.data(), a); + Vectorization_add_factor 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 vector_add(data.values_dev.data(), - v_data.values_dev.data(), + Vectorization_add_av 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 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 vector_sadd(data.values_dev.data(), - v_data.values_dev.data(), + Vectorization_sadd_xv 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 vector_sadd(data.values_dev.data(), - v_data.values_dev.data(), + Vectorization_sadd_xav 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 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 vector_multiply(data.values_dev.data(), + Vectorization_multiply_factor 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 vector_scale(data.values_dev.data(), - v_data.values_dev.data()); + Vectorization_scale 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 vector_equ(data.values_dev.data(), - v_data.values_dev.data(), + Vectorization_equ_au 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 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 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 norm2(data.values_dev.data()); + Norm2 norm2(data.values.data()); parallel_reduce(norm2, 0, size, sum, thread_loop_partitioner); } @@ -2004,7 +2004,7 @@ namespace internal MemorySpaceData &data) { Number sum; - MeanValue mean(data.values_dev.data()); + MeanValue 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 norm1(data.values_dev.data()); + Norm1 norm1(data.values.data()); parallel_reduce(norm1, 0, size, sum, thread_loop_partitioner); } @@ -2035,7 +2035,7 @@ namespace internal ::dealii::MemorySpace::Host> &data) { - NormP normp(data.values_dev.data(), p); + NormP normp(data.values.data(), p); parallel_reduce(normp, 0, size, sum, thread_loop_partitioner); } @@ -2054,9 +2054,9 @@ namespace internal &data) { Number sum; - AddAndDot adder(data.values_dev.data(), - v_data.values_dev.data(), - w_data.values_dev.data(), + AddAndDot 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 - <<>>(data.values_dev.data(), s, 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 - <<>>(data.values_dev.data(), + <<>>(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 - <<>>(data.values_dev.data(), + <<>>(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 - <<>>(data.values_dev.data(), a, 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 - <<>>(data.values_dev.data(), + <<>>(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 - <<>>(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(); } @@ -2280,7 +2280,7 @@ namespace internal const int n_blocks = 1 + size / (chunk_size * block_size); ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd <<>>( - 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 <<>>( - 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 <<>>(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 - <<>>(data.values_dev.data(), factor, 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 - <<>>(data.values_dev.data(), - v_data.values_dev.data(), + <<>>(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 - <<>>(data.values_dev.data(), + <<>>(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 - <<>>(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(); } @@ -2428,8 +2428,8 @@ namespace internal Number, ::dealii::LinearAlgebra::CUDAWrappers::kernel::DotProduct> <<>>(result_device, - data.values_dev.data(), - v_data.values_dev.data(), + data.values.data(), + v_data.values.data(), static_cast( size)); AssertCudaKernel(); @@ -2480,7 +2480,7 @@ namespace internal Number, ::dealii::LinearAlgebra::CUDAWrappers::kernel::ElemSum> <<>>(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> <<>>(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 <<>>(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); -- 2.39.5