IndexSet local_owned(numproc * 2);
local_owned.add_range(myid * 2, myid * 2 + 2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_owned, MPI_COMM_WORLD);
// set local values
local_relevant = local_owned;
local_relevant.add_range(1, 2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
// set local values and check them
local_relevant = local_owned;
local_relevant.add_range(1, 2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
// set local values and check them
local_relevant.add_range(4, 7);
}
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
v = 0.;
local_relevant = local_owned;
local_relevant.add_range(1, 2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
// set local values and check them
local_relevant = local_owned;
local_relevant.add_range(1, 2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_owned, MPI_COMM_WORLD);
// set local values
{
const double norm_sqr = v.l2_norm() * v.l2_norm();
AssertThrow(std::fabs(v * v - norm_sqr) < 1e-15, ExcInternalError());
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
v2 = v;
AssertThrow(std::fabs(v2 * v - norm_sqr) < 1e-15, ExcInternalError());
bool allzero = v.all_zero();
if (myid == 0)
deallog << " v==0 ? " << allzero << std::endl;
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
v2.reinit(v);
allzero = v2.all_zero();
if (myid == 0)
// v has ghosts, w has none. set some entries
// on w, copy into v and check if they are
// there
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(
local_owned, local_owned, MPI_COMM_WORLD);
// set a few of the local elements
local_relevant = local_owned;
local_relevant.add_range(1, 2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v);
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v);
// set local values and check them
LinearAlgebra::ReadWriteVector<double> rw_vector(local_owned);
local_relevant = local_owned;
local_relevant.add_index(2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
AssertDimension(static_cast<unsigned int>(actual_local_size), v.local_size());
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v), x(v),
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v), x(v),
y(v);
// set local elements
if (numproc > 2)
local_relevant0.add_index(8);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v0(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v0(
local_owned0, local_relevant0, MPI_COMM_WORLD);
// vector1: local size 4
local_relevant1.add_index(10);
}
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v1(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v1(
local_owned1, local_relevant1, MPI_COMM_WORLD);
v0 = 1;
deallog << "Ghost values after re-set OK" << std::endl;
// swap with an empty vector
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
v2.swap(v0);
AssertDimension(v0.size(), 0);
AssertDimension(v2.size(), global_size1);
local_relevant = local_owned;
local_relevant.add_range(1, 2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
// set local values and check them
// and once where they have not
for (unsigned int run = 0; run < 2; ++run)
{
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
// set local values
v.import(rw_vector, VectorOperation::insert);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w(v), u(v);
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w(v), u(v);
u = 0;
v *= 2.0;
}
// copy vector content to non-ghosted vectors, manually created.
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v_dist(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v_dist(
local_owned, MPI_COMM_WORLD),
w_dist(v_dist), u_dist(v_dist);
if (numproc > 1)
local_relevant.add_range(3, 4);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
// set local values
if (myid == 0)
deallog << "v has ghost elements: " << v.has_ghost_elements() << std::endl;
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> w, x;
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> w, x;
w = v;
if (myid == 0)
deallog << "w has ghost elements: " << w.has_ghost_elements() << std::endl;
#include "../tests.h"
-__global__ void
-set_value(double *values_dev, unsigned int index, double val)
-{
- values_dev[index] = val;
-}
-
void
test()
{
local_relevant.add_range(min_index + 38, min_index + 40);
local_relevant.add_range(min_index + 41, min_index + 43);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
local_owned, local_relevant, MPI_COMM_WORLD);
+ Kokkos::View<double*, MemorySpace::Device::kokkos_space> v_view(v.get_values(), local_relevant.n_elements());
- deallog << "Local range of proc 0: " << v.local_range().first << " "
- << v.local_range().second << std::endl;
+ deallog << "Local range of proc 0: " << v.get_partitioner()->local_range().first << " "
+ << v.get_partitioner()->local_range().second << std::endl;
// set local values
for (types::global_dof_index i = 0; i < local_size; ++i)
{
- double *values_dev = v.get_values();
- set_value<<<1, 1>>>(values_dev, i, min_index + myid * local_size + i);
+ Kokkos::deep_copy(Kokkos::subview(v_view, i), min_index + myid * local_size + i);
}
deallog << "vector norm: " << v.l2_norm() << std::endl;
v.print(deallog.get_file_stream(), 12, false, false);
v.zero_out_ghost_values();
- double * values_dev = v.get_values();
const auto &partitioner = v.get_partitioner();
- set_value<<<1, 1>>>(values_dev,
- partitioner->global_to_local(min_index + 38),
- min_index);
- set_value<<<1, 1>>>(values_dev,
- partitioner->global_to_local(min_index + 39),
- min_index * 2);
- set_value<<<1, 1>>>(values_dev,
- partitioner->global_to_local(min_index + 41),
- min_index + 7);
- set_value<<<1, 1>>>(values_dev,
- partitioner->global_to_local(min_index + 42),
- -static_cast<double>(min_index));
+ Kokkos::deep_copy(Kokkos::subview(v_view, partitioner->global_to_local(min_index + 38)), min_index);
+ Kokkos::deep_copy(Kokkos::subview(v_view, partitioner->global_to_local(min_index + 39)), min_index*2);
+ Kokkos::deep_copy(Kokkos::subview(v_view, partitioner->global_to_local(min_index + 41)), min_index+7);
+ Kokkos::deep_copy(Kokkos::subview(v_view, partitioner->global_to_local(min_index + 42)), -static_cast<double>(min_index));
v.compress(VectorOperation::add);
v.update_ghost_values();
v.print(deallog.get_file_stream(), 12, false, false);
ghost_set.add_index(0);
ghost_set.add_index(2);
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v(
locally_owned, ghost_set, MPI_COMM_WORLD);
// create vector without actually setting the entries since they will be
// overwritten soon anyway
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> v2;
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> v2;
v2.reinit(v, true);
// set locally owned range of v2 manually
vec_ref.compress(VectorOperation::insert);
auto partitioner = vec_ref.get_partitioner();
- LinearAlgebra::distributed::Vector<double, MemorySpace::CUDA> vec_dev(
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Device> vec_dev(
partitioner);
LinearAlgebra::distributed::Vector<double, MemorySpace::Host> vec_host(
partitioner);
- // Assignment from Host to CUDA
+ // Assignment from Host to Device
vec_dev.import(vec_ref, VectorOperation::insert);
- // Assignment from CUDA to Host
+ // Assignment from Device to Host
vec_host.import(vec_dev, VectorOperation::insert);
for (unsigned int i = 0; i < ghost_size; ++i)
// Check for a bug where compress cannot be called for a LA::d::Vector<Host> in
-// a CUDA file
+// a Device file
#include <deal.II/base/index_set.h>
#include <deal.II/base/utilities.h>