}
/**
- * 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
/**
* 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)
(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
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::
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 *,
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();
}
// 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;
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))
{}
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::
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 *,
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;
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
begin(::dealii::MemorySpace::
MemorySpaceData<Number, ::dealii::MemorySpace::Host> &data)
{
- return data.values_dev.data();
+ return data.values.data();
}
static inline
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();
}
};
begin(::dealii::MemorySpace::
MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
{
- return data.values_dev.data();
+ return data.values.data();
}
static inline
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
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)];
}
// (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)];
}
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];
}
partitioner->locally_owned_size() +
partitioner->n_ghost_indices());
- return data.values_dev[local_index];
+ return data.values[local_index];
}
{
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
{
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>>(
{
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);
}
};
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;
}
}
::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);
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,
// 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);
// 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,
// 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;
}
// 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;
}
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));
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());
}
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());
}
}
// 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,
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
// 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);
# 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);
}
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());
}
}
// 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,
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);
}
// 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),
{
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);
}
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;
const size_type /*size*/,
real_type & /*sum*/,
Number * /*values*/,
- Number * /*values_dev*/)
+ Number * /*values*/)
{}
template <typename real_type>
::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);
}
::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);
}
::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);
}
::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);
}
::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);
}
::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);
}
&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);
}
::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);
}
::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);
&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);
}
::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);
}
::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);
}
::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);
}
&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);
}
{
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);
::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);
}
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;
::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);
}
::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);
}
&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);
{
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);
::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);
{
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();
}
{
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();
}
{
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();
}
{
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();
}
{
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();
}
{
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();
}
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();
}
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();
}
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();
}
{
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();
}
{
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();
}
{
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();
}
{
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();
}
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();
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
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
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);
{
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);