begin(::dealii::MemorySpace::
MemorySpaceData<Number, ::dealii::MemorySpace::Host> &data)
{
- return data.values.get();
+ return data.values.data();
}
static inline
begin(const ::dealii::MemorySpace::
MemorySpaceData<Number, ::dealii::MemorySpace::Host> &data)
{
- return data.values.get();
+ return data.values.data();
}
static inline Number *
get_values(::dealii::MemorySpace::
MemorySpaceData<Number, ::dealii::MemorySpace::Host> &data)
{
- return data.values.get();
+ return data.values.data();
}
};
begin(::dealii::MemorySpace::
MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
{
- return data.values_dev.get();
+ return data.values_dev.data();
}
static inline
begin(const ::dealii::MemorySpace::
MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
{
- return data.values_dev.get();
+ return data.values_dev.data();
}
static inline Number *
get_values(::dealii::MemorySpace::
MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
{
- return data.values_dev.get();
+ return data.values_dev.data();
}
};
} // namespace internal
{
if (comm_shared == MPI_COMM_SELF)
{
- Number *new_val;
- Utilities::System::posix_memalign(
- reinterpret_cast<void **>(&new_val),
- 64,
- sizeof(Number) * new_alloc_size);
- data.values = {new_val, [](Number *data) { std::free(data); }};
+ Kokkos::resize(data.values, new_alloc_size);
allocated_size = new_alloc_size;
data.values_sm = {
- ArrayView<const Number>(data.values.get(), 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 = {ptr_aligned, [mpi_window](Number *) mutable {
- // note: we are creating here a copy of the
- // window other approaches led to segmentation
- // faults
- const auto ierr = MPI_Win_free(&mpi_window);
- AssertThrowMPI(ierr);
- }};
+ data.values =
+ Kokkos::View<Number *,
+ Kokkos::HostSpace,
+ Kokkos::MemoryTraits<Kokkos::Unmanaged>>(
+ ptr_aligned, new_alloc_size);
+
+ // Kokkos will not free the memory because the memory is
+ // unmanaged. Instead we use a shared pointer to take care of
+ // that.
+ data.values_sm_ptr = {ptr_aligned,
+ [mpi_window](Number *) mutable {
+ // note: we are creating here a copy of
+ // the window other approaches led to
+ // segmentation faults
+ const auto ierr =
+ MPI_Win_free(&mpi_window);
+ AssertThrowMPI(ierr);
+ }};
+
#else
Assert(false, ExcInternalError());
#endif
if (new_alloc_size > allocated_size)
{
- Assert(((allocated_size > 0 && data.values_dev != nullptr) ||
- data.values_dev == nullptr),
+ Assert(((allocated_size > 0 && data.values_dev.size() != 0) ||
+ data.values_dev.size() == 0),
ExcInternalError());
- Number *new_val_dev;
- Utilities::CUDA::malloc(new_val_dev, new_alloc_size);
- data.values_dev.reset(new_val_dev);
+ Kokkos::resize(data.values_dev, new_alloc_size);
allocated_size = new_alloc_size;
}
else if (new_alloc_size == 0)
{
- data.values_dev.reset();
+ Kokkos::resize(data.values_dev, 0);
allocated_size = 0;
}
}
::dealii::LinearAlgebra::CUDAWrappers::kernel::add_permutated<
Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
indices_dev,
- data.values_dev.get(),
+ data.values_dev.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.get(),
+ data.values_dev.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.get(), size);
+ result_device, data.values_dev.data(), size);
// Copy the result back to the host
error_code = cudaMemcpy(&result,
resize_val(size, comm_sm);
// delete previous content in import data
- import_data.values.reset();
- import_data.values_dev.reset();
+ Kokkos::resize(import_data.values, 0);
+ Kokkos::resize(import_data.values_dev, 0);
// set partitioner to serial version
partitioner = std::make_shared<Utilities::MPI::Partitioner>(size);
resize_val(local_size + ghost_size, comm_sm);
// delete previous content in import data
- import_data.values.reset();
- import_data.values_dev.reset();
+ Kokkos::resize(import_data.values, 0);
+ Kokkos::resize(import_data.values_dev, 0);
// create partitioner
partitioner = std::make_shared<Utilities::MPI::Partitioner>(local_size,
// is only used as temporary storage for compress() and
// update_ghost_values, and we might have vectors where we never
// call these methods and hence do not need to have the storage.
- import_data.values.reset();
- import_data.values_dev.reset();
+ Kokkos::resize(import_data.values, 0);
+ Kokkos::resize(import_data.values_dev, 0);
thread_loop_partitioner = v.thread_loop_partitioner;
}
// is only used as temporary storage for compress() and
// update_ghost_values, and we might have vectors where we never
// call these methods and hence do not need to have the storage.
- import_data.values.reset();
- import_data.values_dev.reset();
+ Kokkos::resize(import_data.values, 0);
+ Kokkos::resize(import_data.values_dev, 0);
vector_is_ghosted = false;
}
void
Vector<Number, MemorySpaceType>::zero_out_ghost_values() const
{
- if (data.values != nullptr)
- std::fill_n(data.values.get() + partitioner->locally_owned_size(),
+ if (data.values.size() != 0)
+ std::fill_n(data.values.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices(),
Number());
#ifdef DEAL_II_COMPILER_CUDA_AWARE
- if (data.values_dev != nullptr)
+ if (data.values_dev.size() != 0)
{
const cudaError_t cuda_error_code =
- cudaMemset(data.values_dev.get() +
+ cudaMemset(data.values_dev.data() +
partitioner->locally_owned_size(),
0,
partitioner->n_ghost_indices() * sizeof(Number));
defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
if (std::is_same<MemorySpaceType, dealii::MemorySpace::CUDA>::value)
{
- if (import_data.values_dev == nullptr)
- import_data.values_dev.reset(
- Utilities::CUDA::allocate_device_data<Number>(
- partitioner->n_import_indices()));
+ if (import_data.values_dev.size() == 0)
+ Kokkos::resize(import_data.values_dev, partitioner->n_import_indices());
}
else
# endif
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 == nullptr)
- {
- Number *new_val;
- Utilities::System::posix_memalign(
- reinterpret_cast<void **>(&new_val),
- 64,
- sizeof(Number) * partitioner->n_import_indices());
- import_data.values.reset(new_val);
- }
+ 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.
- Number *new_val;
- Utilities::System::posix_memalign(reinterpret_cast<void **>(&new_val),
- 64,
- sizeof(Number) * allocated_size);
-
- data.values = {new_val, [](Number *data) { std::free(data); }};
-
- cudaError_t cuda_error_code =
- cudaMemcpy(data.values.get(),
- data.values_dev.get(),
- allocated_size * sizeof(Number),
- cudaMemcpyDeviceToHost);
- AssertCuda(cuda_error_code);
+ data.values = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, data.values_dev);
}
# endif
operation,
communication_channel,
ArrayView<Number, MemorySpace::CUDA>(
- data.values_dev.get() + partitioner->locally_owned_size(),
+ data.values_dev.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices()),
ArrayView<Number, MemorySpace::CUDA>(
- import_data.values_dev.get(), partitioner->n_import_indices()),
+ import_data.values_dev.data(), partitioner->n_import_indices()),
compress_requests);
}
else
operation,
communication_channel,
ArrayView<Number, MemorySpace::Host>(
- data.values.get() + partitioner->locally_owned_size(),
+ data.values.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices()),
ArrayView<Number, MemorySpace::Host>(
- import_data.values.get(), partitioner->n_import_indices()),
+ import_data.values.data(), partitioner->n_import_indices()),
compress_requests);
}
#else
if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
{
Assert(partitioner->n_import_indices() == 0 ||
- import_data.values_dev != nullptr,
+ import_data.values_dev.size() != 0,
ExcNotInitialized());
partitioner
->import_from_ghosted_array_finish<Number, MemorySpace::CUDA>(
operation,
ArrayView<const Number, MemorySpace::CUDA>(
- import_data.values_dev.get(), partitioner->n_import_indices()),
+ import_data.values_dev.data(), partitioner->n_import_indices()),
ArrayView<Number, MemorySpace::CUDA>(
- data.values_dev.get(), partitioner->locally_owned_size()),
+ data.values_dev.data(), partitioner->locally_owned_size()),
ArrayView<Number, MemorySpace::CUDA>(
- data.values_dev.get() + partitioner->locally_owned_size(),
+ data.values_dev.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices()),
compress_requests);
}
# endif
{
Assert(partitioner->n_import_indices() == 0 ||
- import_data.values != nullptr,
+ import_data.values.size() != 0,
ExcNotInitialized());
partitioner
->import_from_ghosted_array_finish<Number, MemorySpace::Host>(
operation,
ArrayView<const Number, MemorySpace::Host>(
- import_data.values.get(), partitioner->n_import_indices()),
+ import_data.values.data(), partitioner->n_import_indices()),
ArrayView<Number, MemorySpace::Host>(
- data.values.get(), partitioner->locally_owned_size()),
+ data.values.data(), partitioner->locally_owned_size()),
ArrayView<Number, MemorySpace::Host>(
- data.values.get() + partitioner->locally_owned_size(),
+ data.values.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices()),
compress_requests);
}
if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
{
cudaError_t cuda_error_code =
- cudaMemcpy(data.values_dev.get(),
- data.values.get(),
+ cudaMemcpy(data.values_dev.data(),
+ data.values.data(),
allocated_size * sizeof(Number),
cudaMemcpyHostToDevice);
AssertCuda(cuda_error_code);
- data.values.reset();
+ Kokkos::resize(data.values, 0);
}
# endif
#else
(std::is_same<MemorySpaceType, dealii::MemorySpace::CUDA>::value),
ExcMessage(
"Using MemorySpace::CUDA only allowed if the code is compiled with a CUDA compiler!"));
- if (import_data.values_dev == nullptr)
- import_data.values_dev.reset(
- Utilities::CUDA::allocate_device_data<Number>(
- partitioner->n_import_indices()));
+ if (import_data.values_dev.size() == 0)
+ Kokkos::resize(import_data.values_dev, partitioner->n_import_indices());
# else
# ifdef DEAL_II_MPI_WITH_CUDA_SUPPORT
static_assert(
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 == nullptr)
- {
- Number *new_val;
- Utilities::System::posix_memalign(
- reinterpret_cast<void **>(&new_val),
- 64,
- sizeof(Number) * partitioner->n_import_indices());
- import_data.values.reset(new_val);
- }
+ if (import_data.values.size() == 0)
+ Kokkos::resize(import_data.values, partitioner->n_import_indices());
# endif
}
// 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.
- Number *new_val;
- Utilities::System::posix_memalign(reinterpret_cast<void **>(&new_val),
- 64,
- sizeof(Number) * allocated_size);
-
- data.values = {new_val, [](Number *data) { std::free(data); }};
-
- cudaError_t cuda_error_code = cudaMemcpy(data.values.get(),
- data.values_dev.get(),
- allocated_size * sizeof(Number),
- cudaMemcpyDeviceToHost);
- AssertCuda(cuda_error_code);
+ data.values = Kokkos::create_mirror_view_and_copy(Kokkos::HostSpace{}, data.values_dev);
# endif
# if !(defined(DEAL_II_COMPILER_CUDA_AWARE) && \
partitioner->export_to_ghosted_array_start<Number, MemorySpace::Host>(
communication_channel,
ArrayView<const Number, MemorySpace::Host>(
- data.values.get(), partitioner->locally_owned_size()),
- ArrayView<Number, MemorySpace::Host>(import_data.values.get(),
+ data.values.data(), partitioner->locally_owned_size()),
+ ArrayView<Number, MemorySpace::Host>(import_data.values.data(),
partitioner->n_import_indices()),
ArrayView<Number, MemorySpace::Host>(
- data.values.get() + partitioner->locally_owned_size(),
+ data.values.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices()),
update_ghost_values_requests);
# else
partitioner->export_to_ghosted_array_start<Number, MemorySpace::CUDA>(
communication_channel,
ArrayView<const Number, MemorySpace::CUDA>(
- data.values_dev.get(), partitioner->locally_owned_size()),
- ArrayView<Number, MemorySpace::CUDA>(import_data.values_dev.get(),
+ data.values_dev.data(), partitioner->locally_owned_size()),
+ ArrayView<Number, MemorySpace::CUDA>(import_data.values_dev.data(),
partitioner->n_import_indices()),
ArrayView<Number, MemorySpace::CUDA>(
- data.values_dev.get() + partitioner->locally_owned_size(),
+ data.values_dev.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices()),
update_ghost_values_requests);
# endif
defined(DEAL_II_MPI_WITH_CUDA_SUPPORT))
partitioner->export_to_ghosted_array_finish(
ArrayView<Number, MemorySpace::Host>(
- data.values.get() + partitioner->locally_owned_size(),
+ data.values.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices()),
update_ghost_values_requests);
# else
partitioner->export_to_ghosted_array_finish(
ArrayView<Number, MemorySpace::CUDA>(
- data.values_dev.get() + partitioner->locally_owned_size(),
+ data.values_dev.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices()),
update_ghost_values_requests);
# endif
if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
{
cudaError_t cuda_error_code =
- cudaMemcpy(data.values_dev.get() +
+ cudaMemcpy(data.values_dev.data() +
partitioner->locally_owned_size(),
- data.values.get() + partitioner->locally_owned_size(),
+ data.values.data() + partitioner->locally_owned_size(),
partitioner->n_ghost_indices() * sizeof(Number),
cudaMemcpyHostToDevice);
AssertCuda(cuda_error_code);
- data.values.reset();
+ Kokkos::resize(data.values, 0);
}
# endif
if (partitioner.use_count() > 0)
memory +=
partitioner->memory_consumption() / partitioner.use_count() + 1;
- if (import_data.values != nullptr || import_data.values_dev != nullptr)
+ if (import_data.values.size() != 0 || import_data.values_dev.size() != 0)
memory += (static_cast<std::size_t>(partitioner->n_import_indices()) *
sizeof(Number));
return memory;
::dealii::MemorySpace::Host>
&data)
{
- Vector_copy<Number, Number2> copier(v_data.values.get(),
- data.values.get());
+ 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.get());
+ 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.get(),
- v_data.values.get());
+ 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.get(),
- v_data.values.get());
+ 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.get(), 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.get(),
- v_data.values.get(),
+ 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.get(), v_data.values.get(), w_data.values.get(), 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.get(),
- v_data.values.get(),
+ 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.get(),
- v_data.values.get(),
+ 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.get(), v_data.values.get(), w_data.values.get(), 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.get(),
+ 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.get(),
- v_data.values.get());
+ 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.get(),
- v_data.values.get(),
+ 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.get(), v_data.values.get(), w_data.values.get(), 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.get(), v_data.values.get());
+ 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.get());
+ 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.get());
+ 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.get());
+ 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.get(), 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.get(),
- v_data.values.get(),
- w_data.values.get(),
+ 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.get(),
- v_data.values_dev.get(),
+ cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
+ v_data.values_dev.data(),
size * sizeof(Number),
cudaMemcpyDeviceToHost);
AssertCuda(cuda_error_code);
::dealii::MemorySpace::CUDA>
&data)
{
- cudaError_t cuda_error_code = cudaMemcpy(data.values_dev.get(),
- v_data.values_dev.get(),
+ cudaError_t cuda_error_code = cudaMemcpy(data.values_dev.data(),
+ v_data.values_dev.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.get(), s, size);
+ <<<n_blocks, block_size>>>(data.values_dev.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.get(),
+ <<<n_blocks, block_size>>>(data.values_dev.data(),
1.,
- v_data.values_dev.get(),
+ v_data.values_dev.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.get(),
+ <<<n_blocks, block_size>>>(data.values_dev.data(),
-1.,
- v_data.values_dev.get(),
+ v_data.values_dev.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.get(), a, size);
+ <<<n_blocks, block_size>>>(data.values_dev.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.get(),
+ <<<n_blocks, block_size>>>(data.values_dev.data(),
a,
- v_data.values_dev.get(),
+ v_data.values_dev.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.get(),
+ <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values_dev.data(),
a,
- v_data.values_dev.get(),
+ v_data.values_dev.data(),
b,
- w_data.values_dev.get(),
+ w_data.values_dev.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.get(), 1., v_data.values_dev.get(), size);
+ x, data.values_dev.data(), 1., v_data.values_dev.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.get(), a, v_data.values_dev.get(), size);
+ x, data.values_dev.data(), a, v_data.values_dev.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.get(),
+ data.values_dev.data(),
a,
- v_data.values_dev.get(),
+ v_data.values_dev.data(),
b,
- w_data.values_dev.get(),
+ w_data.values_dev.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.get(), factor, size);
+ <<<n_blocks, block_size>>>(data.values_dev.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.get(),
- v_data.values_dev.get(),
+ <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values_dev.data(),
+ v_data.values_dev.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.get(),
+ <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values_dev.data(),
a,
- v_data.values_dev.get(),
+ v_data.values_dev.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.get(),
+ <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values_dev.data(),
a,
- v_data.values_dev.get(),
+ v_data.values_dev.data(),
b,
- w_data.values_dev.get(),
+ w_data.values_dev.data(),
size);
AssertCudaKernel();
}
Number,
::dealii::LinearAlgebra::CUDAWrappers::kernel::DotProduct<Number>>
<<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
- data.values_dev.get(),
- v_data.values_dev.get(),
+ data.values_dev.data(),
+ v_data.values_dev.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.get(),
+ data.values_dev.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.get(),
+ data.values_dev.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.get(),
- v_data.values_dev.get(),
- w_data.values_dev.get(),
+ data.values_dev.data(),
+ v_data.values_dev.data(),
+ w_data.values_dev.data(),
a,
size);
{
if (operation == VectorOperation::insert)
{
- cudaError_t cuda_error_code = cudaMemcpy(data.values_dev.get(),
- v_data.values.get(),
+ cudaError_t cuda_error_code = cudaMemcpy(data.values_dev.data(),
+ v_data.values.data(),
size * sizeof(Number),
cudaMemcpyHostToDevice);
AssertCuda(cuda_error_code);