template <typename Number, typename MemorySpaceType>
struct la_parallel_vector_templates_functions
{
- static_assert(std::is_same<MemorySpaceType, MemorySpace::Host>::value ||
- std::is_same<MemorySpaceType, MemorySpace::Device>::value,
- "MemorySpace should be Host or CUDA");
+ static_assert(
+ std::is_same<MemorySpaceType, MemorySpace::Host>::value ||
+ std::is_same<MemorySpaceType, MemorySpace::Device>::value,
+ "MemorySpace should be Host or Device");
static void
resize_val(
};
template <typename Number>
- struct la_parallel_vector_templates_functions<Number,
- ::dealii::MemorySpace::Device>
+ struct la_parallel_vector_templates_functions<
+ Number,
+ ::dealii::MemorySpace::Device>
{
using size_type = types::global_dof_index;
static void
- resize_val(const types::global_dof_index new_alloc_size,
- types::global_dof_index & allocated_size,
- ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data,
- const MPI_Comm &comm_sm)
+ resize_val(
+ const types::global_dof_index new_alloc_size,
+ types::global_dof_index & allocated_size,
+ ::dealii::MemorySpace::MemorySpaceData<Number,
+ ::dealii::MemorySpace::Device>
+ & data,
+ const MPI_Comm &comm_sm)
{
(void)comm_sm;
static_assert(
std::is_same<Number, float>::value ||
std::is_same<Number, double>::value,
- "Number should be float or double for CUDA memory space");
+ "Number should be float or double for Device memory space");
if (new_alloc_size > allocated_size)
{
// including ghost entries. this is not really efficient right now
// because indices are translated twice, once by nth_index_in_set(i)
// and once for operator() of tmp_vector
- const IndexSet & v_stored = V.get_stored_elements();
- const size_type n_elements = v_stored.n_elements();
- Kokkos::DefaultHostExecutionSpace host_exec;
- Kokkos::View<size_type*, Kokkos::HostSpace> indices("indices", n_elements);
- Kokkos::parallel_for(Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec, 0, n_elements),
- KOKKOS_LAMBDA(size_type i) {
- indices[i] = communication_pattern->global_to_local(
- v_stored.nth_index_in_set(i));
- });
- host_exec.fence();
-
+ const IndexSet & v_stored = V.get_stored_elements();
+ const size_type n_elements = v_stored.n_elements();
+ Kokkos::DefaultHostExecutionSpace host_exec;
+ Kokkos::View<size_type *, Kokkos::HostSpace> indices("indices",
+ n_elements);
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec,
+ 0,
+ n_elements),
+ KOKKOS_LAMBDA(size_type i) {
+ indices[i] = communication_pattern->global_to_local(
+ v_stored.nth_index_in_set(i));
+ });
+ host_exec.fence();
+
// Move the indices to the device
::dealii::MemorySpace::Device::kokkos_space::execution_space exec;
- Kokkos::View<size_type *, ::dealii::MemorySpace::Device::kokkos_space> indices_dev("indices_dev", n_elements);
- Kokkos::deep_copy(exec, indices_dev, indices);
- exec.fence();
-
+ Kokkos::View<size_type *, ::dealii::MemorySpace::Device::kokkos_space>
+ indices_dev("indices_dev", n_elements);
+ Kokkos::deep_copy(exec, indices_dev, indices);
+ exec.fence();
+
// Move the data to the device
- Kokkos::View<Number *, ::dealii::MemorySpace::Device::kokkos_space> V_dev("V_dev", n_elements);
- Kokkos::deep_copy(exec, V_dev, Kokkos::View<Number*, Kokkos::HostSpace>(V.begin(), n_elements));
+ Kokkos::View<Number *, ::dealii::MemorySpace::Device::kokkos_space>
+ V_dev("V_dev", n_elements);
+ Kokkos::deep_copy(
+ exec,
+ V_dev,
+ Kokkos::View<Number *, Kokkos::HostSpace>(V.begin(), n_elements));
exec.fence();
// Set the values in tmp_vector
// Copy the local elements of tmp_vector to the right place in val
IndexSet tmp_index_set = tmp_vector.locally_owned_elements();
const size_type tmp_n_elements = tmp_index_set.n_elements();
- Kokkos::realloc(indices, tmp_n_elements);
- Kokkos::parallel_for(Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec, 0, n_elements),
- KOKKOS_LAMBDA(size_type i) {
- indices[i] = locally_owned_elem.index_within_set(
- tmp_index_set.nth_index_in_set(i));
+ Kokkos::realloc(indices, tmp_n_elements);
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<Kokkos::DefaultHostExecutionSpace>(host_exec,
+ 0,
+ n_elements),
+ KOKKOS_LAMBDA(size_type i) {
+ indices[i] = locally_owned_elem.index_within_set(
+ tmp_index_set.nth_index_in_set(i));
});
host_exec.fence();
- Kokkos::realloc(indices_dev, tmp_n_elements);
- Kokkos::deep_copy(indices_dev, Kokkos::View<size_type*>(indices.data(), tmp_n_elements));
+ Kokkos::realloc(indices_dev, tmp_n_elements);
+ Kokkos::deep_copy(indices_dev,
+ Kokkos::View<size_type *>(indices.data(),
+ tmp_n_elements));
- if (operation == VectorOperation::add)
+ if (operation == VectorOperation::add)
Kokkos::parallel_for(
Kokkos::RangePolicy<
::dealii::MemorySpace::Device::kokkos_space::execution_space>(
linfty_norm_local(const ::dealii::MemorySpace::MemorySpaceData<
Number,
::dealii::MemorySpace::Device> &data,
- const unsigned int size,
- RealType & result)
+ const unsigned int size,
+ RealType & result)
{
static_assert(std::is_same<Number, RealType>::value,
"RealType should be the same type as Number");
- typename ::dealii::MemorySpace::Device::kokkos_space::execution_space
+ typename ::dealii::MemorySpace::Device::kokkos_space::execution_space
exec;
Kokkos::parallel_reduce(
Kokkos::RangePolicy<
if (data.values_host_buffer.size() > 0)
Kokkos::deep_copy(Kokkos::subview(data.values_host_buffer, range), 0);
if (data.values.size() > 0)
- Kokkos::deep_copy(Kokkos::subview(data.values, range), 0);
+ Kokkos::deep_copy(Kokkos::subview(data.values, range), 0);
vector_is_ghosted = false;
}
// The communication is done on the host, so we need to
// move the data back to the device.
- auto range = Kokkos::make_pair(
- partitioner->locally_owned_size(),
- partitioner->locally_owned_size() + partitioner->n_ghost_indices());
+ auto range = Kokkos::make_pair(partitioner->locally_owned_size(),
+ partitioner->locally_owned_size() +
+ partitioner->n_ghost_indices());
Kokkos::deep_copy(Kokkos::subview(data.values, range),
- Kokkos::subview(data.values_host_buffer, range));
+ Kokkos::subview(data.values_host_buffer,
+ range));
Kokkos::resize(data.values_host_buffer, 0);
}
{
real_type sum;
-
dealii::internal::VectorOperations::
functions<Number, Number, MemorySpaceType>::norm_2(
thread_loop_partitioner,
::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace> & /*data*/)
{
static_assert(
- std::is_same<MemorySpace, ::dealii::MemorySpace::CUDA>::value &&
+ std::is_same<MemorySpace, ::dealii::MemorySpace::Device>::value &&
std::is_same<Number, Number2>::value,
- "For the CUDA MemorySpace Number and Number2 should be the same type");
+ "For the Device MemorySpace Number and Number2 should be the same type");
}
static void
}
}
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
template <typename MemorySpace2>
static void
import_elements(
::dealii::MemorySpace::Host>
&data,
std::enable_if_t<
- std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::value,
+ std::is_same<MemorySpace2, ::dealii::MemorySpace::Device>::value,
int> = 0)
{
if (operation == VectorOperation::insert)
{
- cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
- v_data.values.data(),
- size * sizeof(Number),
- cudaMemcpyDeviceToHost);
- AssertCuda(cuda_error_code);
+ Kokkos::deep_copy(
+ Kokkos::subview(data.values,
+ Kokkos::pair<size_type, size_type>(0, size)),
+ Kokkos::subview(v_data.values,
+ Kokkos::pair<size_type, size_type>(0, size)));
}
else
{
AssertThrow(false, ExcNotImplemented());
}
}
-#endif
};
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
template <typename Number>
- struct functions<Number, Number, ::dealii::MemorySpace::CUDA>
+ struct functions<Number, Number, ::dealii::MemorySpace::Device>
{
- static const int block_size =
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::block_size;
- static const int chunk_size =
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::chunk_size;
-
static void
copy(
const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
const size_type size,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
- v_data.values.data(),
- size * sizeof(Number),
- cudaMemcpyDeviceToDevice);
- AssertCuda(cuda_error_code);
+ Kokkos::deep_copy(
+ Kokkos::subview(data.values,
+ Kokkos::pair<size_type, size_type>(0, size)),
+ Kokkos::subview(v_data.values,
+ Kokkos::pair<size_type, size_type>(0, size)));
}
static void
const size_type size,
const Number s,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::set<Number>
- <<<n_blocks, block_size>>>(data.values.data(), s, size);
- AssertCudaKernel();
+ Kokkos::deep_copy(
+ Kokkos::subview(data.values,
+ Kokkos::pair<size_type, size_type>(0, size)),
+ s);
}
static void
const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
const size_type size,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
- <<<n_blocks, block_size>>>(data.values.data(),
- 1.,
- v_data.values.data(),
- size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(int i) { data.values(i) += v_data.values(i); });
+ exec.fence();
}
static void
const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
const size_type size,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
- <<<n_blocks, block_size>>>(data.values.data(),
- -1.,
- v_data.values.data(),
- size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) { data.values(i) -= v_data.values(i); });
+ exec.fence();
}
static void
const size_type size,
Number a,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_add<Number>
- <<<n_blocks, block_size>>>(data.values.data(), a, size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) { data.values(i) += a; });
+ exec.fence();
}
static void
const size_type size,
const Number a,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aV<Number>
- <<<n_blocks, block_size>>>(data.values.data(),
- a,
- v_data.values.data(),
- size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(i) += a * v_data.values(i);
+ });
+ exec.fence();
}
static void
const Number a,
const Number b,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_aVbW<Number>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
- a,
- v_data.values.data(),
- b,
- w_data.values.data(),
- size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(i) += a * v_data.values(i) + b * w_data.values(i);
+ });
+ exec.fence();
}
static void
const size_type size,
const Number x,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(
- x, data.values.data(), 1., v_data.values.data(), size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(i) = x * data.values(i) + v_data.values(i);
+ });
+ exec.fence();
}
static void
const Number x,
const Number a,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(
- x, data.values.data(), a, v_data.values.data(), size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(i) = x * data.values(i) + a * v_data.values(i);
+ });
+ exec.fence();
}
static void
const Number a,
const Number b,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd<Number>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(x,
- data.values.data(),
- a,
- v_data.values.data(),
- b,
- w_data.values.data(),
- size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(i) =
+ x * data.values(i) + a * v_data.values(i) + b * w_data.values(i);
+ });
+ exec.fence();
}
static void
const size_type size,
const Number factor,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_scale<Number>
- <<<n_blocks, block_size>>>(data.values.data(), factor, size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) { data.values(i) *= factor; });
+ exec.fence();
}
static void
const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
const size_type size,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::scale<Number>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
- v_data.values.data(),
- size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) { data.values(i) *= v_data.values(i); });
+ exec.fence();
}
static void
const size_type size,
const Number a,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::equ<Number>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
- a,
- v_data.values.data(),
- size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(i) = a * v_data.values(i);
+ });
+ exec.fence();
}
static void
const Number a,
const Number b,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::equ<Number>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(data.values.data(),
- a,
- v_data.values.data(),
- b,
- w_data.values.data(),
- size);
- AssertCudaKernel();
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(i) = a * v_data.values(i) + b * w_data.values(i);
+ });
+ exec.fence();
}
static Number
dot(const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
const size_type size,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- Number * result_device;
- cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
- AssertCuda(error_code);
- error_code = cudaMemset(result_device, 0, sizeof(Number));
- AssertCuda(error_code);
-
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::double_vector_reduction<
- Number,
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::DotProduct<Number>>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
- data.values.data(),
- v_data.values.data(),
- static_cast<unsigned int>(
- size));
- AssertCudaKernel();
-
- // Copy the result back to the host
Number result;
- error_code = cudaMemcpy(&result,
- result_device,
- sizeof(Number),
- cudaMemcpyDeviceToHost);
- AssertCuda(error_code);
- // Free the memory on the device
- error_code = cudaFree(result_device);
- AssertCuda(error_code);
- AssertIsFinite(result);
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_reduce(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i, Number & update) {
+ update += data.values(i) * v_data.values(i);
+ },
+ result);
+ AssertIsFinite(result);
return result;
}
& thread_loop_partitioner,
const size_type size,
real_type & sum,
- ::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
- &data)
+ ::dealii::MemorySpace::
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data)
{
sum = dot(thread_loop_partitioner, size, data, data);
}
const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
const size_type size,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data)
- {
- Number * result_device;
- cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
- AssertCuda(error_code);
- error_code = cudaMemset(result_device, 0, sizeof(Number));
-
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
- Number,
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::ElemSum<Number>>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
- data.values.data(),
- size);
-
- // Copy the result back to the host
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data)
+ {
Number result;
- error_code = cudaMemcpy(&result,
- result_device,
- sizeof(Number),
- cudaMemcpyDeviceToHost);
- AssertCuda(error_code);
- // Free the memory on the device
- error_code = cudaFree(result_device);
- AssertCuda(error_code);
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_reduce(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i, Number & update) {
+ update += data.values(i);
+ },
+ result);
+
+ AssertIsFinite(result);
return result;
}
const size_type size,
real_type & sum,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- Number * result_device;
- cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number));
- AssertCuda(error_code);
- error_code = cudaMemset(result_device, 0, sizeof(Number));
-
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
- Number,
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::L1Norm<Number>>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(result_device,
- data.values.data(),
- size);
-
- // Copy the result back to the host
- error_code = cudaMemcpy(&sum,
- result_device,
- sizeof(Number),
- cudaMemcpyDeviceToHost);
- AssertCuda(error_code);
- // Free the memory on the device
- error_code = cudaFree(result_device);
- AssertCuda(error_code);
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_reduce(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i, Number & update) {
+ update += Kokkos::abs(data.values(i));
+ },
+ sum);
}
template <typename real_type>
static void
norm_p(
const std::shared_ptr<::dealii::parallel::internal::TBBPartitioner> &,
- const size_type,
- real_type &,
- real_type,
+ const size_type size,
+ real_type & sum,
+ real_type exp,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA> &)
+ ::dealii::MemorySpace::Device>
+ &data)
{
- Assert(false, ExcNotImplemented());
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_reduce(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i, Number & update) {
+ update += Kokkos::pow(Kokkos::abs(data.values(i)), exp);
+ },
+ sum);
}
static Number
const size_type size,
const Number a,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &v_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &v_data,
const ::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &w_data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &w_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
- Number * res_d;
- cudaError_t error_code = cudaMalloc(&res_d, sizeof(Number));
- AssertCuda(error_code);
- error_code = cudaMemset(res_d, 0, sizeof(Number));
- AssertCuda(error_code);
-
- const int n_blocks = 1 + size / (chunk_size * block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_and_dot<Number>
- <<<dim3(n_blocks, 1), dim3(block_size)>>>(res_d,
- data.values.data(),
- v_data.values.data(),
- w_data.values.data(),
- a,
- size);
-
Number res;
- error_code =
- cudaMemcpy(&res, res_d, sizeof(Number), cudaMemcpyDeviceToHost);
- AssertCuda(error_code);
- error_code = cudaFree(res_d);
+
+ auto exec = typename ::dealii::MemorySpace::Device::kokkos_space::
+ execution_space{};
+ Kokkos::parallel_reduce(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(exec,
+ 0,
+ size),
+ KOKKOS_LAMBDA(size_type i, Number & update) {
+ data.values(i) += a * v_data.values(i);
+ update +=
+ data.values(i) * Number(numbers::NumberTraits<Number>::conjugate(
+ w_data.values(i)));
+ },
+ res);
return res;
}
const ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace2>
&v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data,
std::enable_if_t<
- std::is_same<MemorySpace2, ::dealii::MemorySpace::CUDA>::value,
+ std::is_same<MemorySpace2, ::dealii::MemorySpace::Device>::value,
int> = 0)
{
if (operation == VectorOperation::insert)
const ::dealii::MemorySpace::MemorySpaceData<Number, MemorySpace2>
&v_data,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data,
std::enable_if_t<
std::is_same<MemorySpace2, ::dealii::MemorySpace::Host>::value,
{
if (operation == VectorOperation::insert)
{
- cudaError_t cuda_error_code = cudaMemcpy(data.values.data(),
- v_data.values.data(),
- size * sizeof(Number),
- cudaMemcpyHostToDevice);
- AssertCuda(cuda_error_code);
+ Kokkos::deep_copy(
+ Kokkos::subview(data.values,
+ Kokkos::pair<size_type, size_type>(0, size)),
+ Kokkos::subview(v_data.values,
+ Kokkos::pair<size_type, size_type>(0, size)));
}
else
{
}
}
};
-#endif
} // namespace VectorOperations
} // namespace internal