struct la_parallel_vector_templates_functions
{
static_assert(std::is_same<MemorySpaceType, MemorySpace::Host>::value ||
- std::is_same<MemorySpaceType, MemorySpace::CUDA>::value,
+ std::is_same<MemorySpaceType, MemorySpace::Device>::value,
"MemorySpace should be Host or CUDA");
static void
}
};
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
template <typename Number>
struct la_parallel_vector_templates_functions<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
{
using size_type = types::global_dof_index;
resize_val(const types::global_dof_index new_alloc_size,
types::global_dof_index & allocated_size,
::dealii::MemorySpace::
- MemorySpaceData<Number, ::dealii::MemorySpace::CUDA> &data,
+ MemorySpaceData<Number, ::dealii::MemorySpace::Device> &data,
const MPI_Comm &comm_sm)
{
(void)comm_sm;
communication_pattern,
const IndexSet &locally_owned_elem,
::dealii::MemorySpace::MemorySpaceData<Number,
- ::dealii::MemorySpace::CUDA>
+ ::dealii::MemorySpace::Device>
&data)
{
Assert(
"Only VectorOperation::add and VectorOperation::insert are allowed"));
::dealii::LinearAlgebra::distributed::
- Vector<Number, ::dealii::MemorySpace::CUDA>
+ Vector<Number, ::dealii::MemorySpace::Device>
tmp_vector(communication_pattern);
// fill entries from ReadWriteVector into the distributed vector,
// and once for operator() of tmp_vector
const IndexSet & v_stored = V.get_stored_elements();
const size_type n_elements = v_stored.n_elements();
- std::vector<size_type> indices(n_elements);
- for (size_type i = 0; i < n_elements; ++i)
+ 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
- size_type *indices_dev;
- ::dealii::Utilities::CUDA::malloc(indices_dev, n_elements);
- ::dealii::Utilities::CUDA::copy_to_dev(indices, indices_dev);
+ ::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();
+
// Move the data to the device
- Number *V_dev;
- ::dealii::Utilities::CUDA::malloc(V_dev, n_elements);
- cudaError_t cuda_error_code = cudaMemcpy(V_dev,
- V.begin(),
- n_elements * sizeof(Number),
- cudaMemcpyHostToDevice);
- AssertCuda(cuda_error_code);
+ 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
- const int n_blocks =
- 1 + n_elements / (::dealii::CUDAWrappers::chunk_size *
- ::dealii::CUDAWrappers::block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<Number>
- <<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
- indices_dev, tmp_vector.begin(), V_dev, n_elements);
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
+ exec, 0, n_elements),
+ KOKKOS_LAMBDA(size_type i) {
+ tmp_vector(indices_dev(i)) = V_dev(i);
+ });
+ exec.fence();
tmp_vector.compress(operation);
// 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();
- indices.resize(tmp_n_elements);
- for (size_type i = 0; i < tmp_n_elements; ++i)
- indices[i] = locally_owned_elem.index_within_set(
+ 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));
- ::dealii::Utilities::CUDA::free(indices_dev);
- ::dealii::Utilities::CUDA::malloc(indices_dev, tmp_n_elements);
- ::dealii::Utilities::CUDA::copy_to_dev(indices, indices_dev);
-
- if (operation == VectorOperation::add)
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_permutated<
- Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
- indices_dev,
- data.values.data(),
- tmp_vector.begin(),
- tmp_n_elements);
+ });
+ host_exec.fence();
+ 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)
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
+ exec, 0, n_elements),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(indices_dev(i)) += tmp_vector(i);
+ });
else
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<
- Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
- indices_dev,
- data.values.data(),
- tmp_vector.begin(),
- tmp_n_elements);
-
- ::dealii::Utilities::CUDA::free(indices_dev);
- ::dealii::Utilities::CUDA::free(V_dev);
+ Kokkos::parallel_for(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
+ exec, 0, n_elements),
+ KOKKOS_LAMBDA(size_type i) {
+ data.values(indices_dev(i)) = tmp_vector(i);
+ });
+ exec.fence();
}
template <typename RealType>
static void
linfty_norm_local(const ::dealii::MemorySpace::MemorySpaceData<
Number,
- ::dealii::MemorySpace::CUDA> &data,
+ ::dealii::MemorySpace::Device> &data,
const unsigned int size,
RealType & result)
{
static_assert(std::is_same<Number, RealType>::value,
"RealType should be the same type as Number");
- 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 / (::dealii::CUDAWrappers::chunk_size *
- ::dealii::CUDAWrappers::block_size);
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::reduction<
- Number,
- ::dealii::LinearAlgebra::CUDAWrappers::kernel::LInfty<Number>>
- <<<dim3(n_blocks, 1), dim3(::dealii::CUDAWrappers::block_size)>>>(
- result_device, data.values.data(), size);
-
- // Copy the result back to the host
- 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);
+ typename ::dealii::MemorySpace::Device::kokkos_space::execution_space
+ exec;
+ Kokkos::parallel_reduce(
+ Kokkos::RangePolicy<
+ ::dealii::MemorySpace::Device::kokkos_space::execution_space>(
+ exec, 0, size),
+ KOKKOS_LAMBDA(size_type i, RealType & update) {
+ update += Kokkos::abs(data.values(i));
+ },
+ result);
}
};
-#endif
} // namespace internal
void
Vector<Number, MemorySpaceType>::zero_out_ghost_values() const
{
- if (data.values.size() != 0)
- {
-#ifdef DEAL_II_COMPILER_CUDA_AWARE
- if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
- {
- const cudaError_t cuda_error_code =
- cudaMemset(data.values.data() +
- partitioner->locally_owned_size(),
- 0,
- partitioner->n_ghost_indices() * sizeof(Number));
- AssertCuda(cuda_error_code);
- }
- else
-#endif
- {
- std::fill_n(data.values.data() +
- partitioner->locally_owned_size(),
- partitioner->n_ghost_indices(),
- Number());
- }
- }
+ Kokkos::pair<size_type, size_type> range(
+ partitioner->locally_owned_size(),
+ partitioner->locally_owned_size() + partitioner->n_ghost_indices());
+ 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);
vector_is_ghosted = false;
}
// allocate import_data in case it is not set up yet
if (partitioner->n_import_indices() > 0)
{
-# if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
- !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
- if (std::is_same<MemorySpaceType, dealii::MemorySpace::CUDA>::value)
+# if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+ if (std::is_same<MemorySpaceType, dealii::MemorySpace::Device>::value)
{
if (import_data.values_host_buffer.size() == 0)
Kokkos::resize(import_data.values_host_buffer,
else
# endif
{
-# if !defined(DEAL_II_COMPILER_CUDA_AWARE) && \
- defined(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.size() == 0)
Kokkos::resize(import_data.values,
partitioner->n_import_indices());
}
}
-# if defined DEAL_II_COMPILER_CUDA_AWARE && \
- !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
- if (std::is_same<MemorySpaceType, dealii::MemorySpace::CUDA>::value)
+# if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+ if (std::is_same<MemorySpaceType, dealii::MemorySpace::Device>::value)
{
// Move the data to the host and then move it back to the
// device. We use values to store the elements because the function
// make this function thread safe
std::lock_guard<std::mutex> lock(mutex);
-# if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
- !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
- if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
+# if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+ if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
{
Assert(partitioner->n_import_indices() == 0 ||
import_data.values_host_buffer.size() != 0,
// 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.data(),
- data.values_host_buffer.data(),
- allocated_size * sizeof(Number),
- cudaMemcpyHostToDevice);
- AssertCuda(cuda_error_code);
+ Kokkos::deep_copy(data.values, data.values_host_buffer);
Kokkos::resize(data.values_host_buffer, 0);
}
// allocate import_data in case it is not set up yet
if (partitioner->n_import_indices() > 0)
{
-# if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
- !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
- if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
+# if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+ if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
{
if (import_data.values_host_buffer.size() == 0)
Kokkos::resize(import_data.values_host_buffer,
}
}
-# if defined DEAL_II_COMPILER_CUDA_AWARE && \
- !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
- if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
+# if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+ if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
{
// Move the data to the host and then move it back to the
// device. We use values to store the elements because the function
// make this function thread safe
std::lock_guard<std::mutex> lock(mutex);
-# if defined(DEAL_II_COMPILER_CUDA_AWARE) && \
- !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
- if (std::is_same<MemorySpaceType, MemorySpace::CUDA>::value)
+# if !defined(DEAL_II_MPI_WITH_CUDA_SUPPORT)
+ if (std::is_same<MemorySpaceType, MemorySpace::Device>::value)
{
partitioner->export_to_ghosted_array_finish(
ArrayView<Number, MemorySpace::Host>(
// 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.data() +
- partitioner->locally_owned_size(),
- data.values_host_buffer.data() +
- partitioner->locally_owned_size(),
- partitioner->n_ghost_indices() * sizeof(Number),
- cudaMemcpyHostToDevice);
- AssertCuda(cuda_error_code);
+ 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::resize(data.values_host_buffer, 0);
}
projects / dealii.git / commitdiff