#include <deal.II/base/config.h>
-#ifdef DEAL_II_WITH_CUDA
+#include <deal.II/base/memory_space.h>
+#include <deal.II/base/tensor.h>
+#include <deal.II/base/utilities.h>
-# include <deal.II/base/memory_space.h>
-# include <deal.II/base/tensor.h>
-# include <deal.II/base/utilities.h>
+#include <deal.II/lac/cuda_atomic.h>
+#include <deal.II/lac/cuda_vector.h>
-# include <deal.II/lac/cuda_atomic.h>
-# include <deal.II/lac/cuda_vector.h>
+#include <deal.II/matrix_free/cuda_hanging_nodes_internal.h>
+#include <deal.II/matrix_free/cuda_matrix_free.h>
+#include <deal.II/matrix_free/cuda_matrix_free.templates.h>
+#include <deal.II/matrix_free/cuda_tensor_product_kernels.h>
-# include <deal.II/matrix_free/cuda_hanging_nodes_internal.h>
-# include <deal.II/matrix_free/cuda_matrix_free.h>
-# include <deal.II/matrix_free/cuda_matrix_free.templates.h>
-# include <deal.II/matrix_free/cuda_tensor_product_kernels.h>
-
-# include <Kokkos_Core.hpp>
+#include <Kokkos_Core.hpp>
DEAL_II_NAMESPACE_OPEN
DEAL_II_NAMESPACE_CLOSE
#endif
-
-#endif
#include <deal.II/base/config.h>
-#ifdef DEAL_II_WITH_CUDA
+#include <deal.II/base/cuda_size.h>
-# include <deal.II/base/cuda_size.h>
+#include <deal.II/matrix_free/hanging_nodes_internal.h>
-# include <deal.II/matrix_free/hanging_nodes_internal.h>
-
-# include <Kokkos_Macros.hpp>
+#include <Kokkos_Macros.hpp>
DEAL_II_NAMESPACE_OPEN
namespace CUDAWrappers
DEAL_II_NAMESPACE_CLOSE
#endif
-
-#endif
#include <deal.II/base/config.h>
-#ifdef DEAL_II_WITH_CUDA
-
-# include <deal.II/base/cuda_size.h>
-# include <deal.II/base/memory_space.h>
-# include <deal.II/base/mpi_stub.h>
-# include <deal.II/base/partitioner.h>
-# include <deal.II/base/quadrature.h>
-# include <deal.II/base/tensor.h>
-# include <deal.II/base/utilities.h>
+#include <deal.II/base/cuda_size.h>
+#include <deal.II/base/memory_space.h>
+#include <deal.II/base/mpi_stub.h>
+#include <deal.II/base/partitioner.h>
+#include <deal.II/base/quadrature.h>
+#include <deal.II/base/tensor.h>
+#include <deal.II/base/utilities.h>
-# include <deal.II/dofs/dof_handler.h>
+#include <deal.II/dofs/dof_handler.h>
-# include <deal.II/fe/fe_update_flags.h>
-# include <deal.II/fe/mapping.h>
+#include <deal.II/fe/fe_update_flags.h>
+#include <deal.II/fe/mapping.h>
-# include <deal.II/grid/filtered_iterator.h>
+#include <deal.II/grid/filtered_iterator.h>
-# include <deal.II/lac/affine_constraints.h>
-# include <deal.II/lac/cuda_vector.h>
-# include <deal.II/lac/la_parallel_vector.h>
+#include <deal.II/lac/affine_constraints.h>
+#include <deal.II/lac/cuda_vector.h>
+#include <deal.II/lac/la_parallel_vector.h>
-# include <Kokkos_Core.hpp>
+#include <Kokkos_Core.hpp>
namespace CUDAWrappers
{
// forward declaration
-# ifndef DOXYGEN
+#ifndef DOXYGEN
namespace internal
{
template <int dim, typename Number>
class ReinitHelper;
}
-# endif
+#endif
/**
* This class collects all the data that is stored for the matrix free
, use_coloring(use_coloring)
, overlap_communication_computation(overlap_communication_computation)
{
-# ifndef DEAL_II_MPI_WITH_DEVICE_SUPPORT
+#ifndef DEAL_II_MPI_WITH_DEVICE_SUPPORT
AssertThrow(
overlap_communication_computation == false,
ExcMessage(
"Overlapping communication and computation requires CUDA-aware MPI."));
-# endif
+#endif
if (overlap_communication_computation == true)
AssertThrow(
use_coloring == false || overlap_communication_computation == false,
void
set_constrained_values(const Number val, VectorType &dst) const;
+#ifdef DEAL_II_WITH_CUDA
/**
* Initialize a serial vector. The size corresponds to the number of degrees
* of freedom in the DoFHandler object.
void
initialize_dof_vector(
LinearAlgebra::CUDAWrappers::Vector<Number> &vec) const;
+#endif
/**
* Initialize a distributed vector. The local elements correspond to the
*/
void
initialize_dof_vector(
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &vec) const;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &vec)
+ const;
/**
* Return the colored graph of locally owned active cells.
template <typename Functor>
void
distributed_cell_loop(
- const Functor & func,
- const LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &src,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &dst) const;
+ const Functor &func,
+ const LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>
+ & src,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &dst)
+ const;
+#ifdef DEAL_II_WITH_CUDA
/**
* This function should never be called. Calling it results in an internal
* error. This function exists only because cell_loop needs
const Functor & func,
const LinearAlgebra::CUDAWrappers::Vector<Number> &src,
LinearAlgebra::CUDAWrappers::Vector<Number> & dst) const;
+#endif
/**
* Unique ID associated with the object.
Kokkos::View<Number *, MemorySpace::Default::kokkos_space>
constraint_weights;
- /**
- * Grid dimensions associated to the different colors. The grid dimensions
- * are used to launch the CUDA kernels.
- */
- std::vector<dim3> grid_dim;
-
- /**
- * Block dimensions associated to the different colors. The block dimensions
- * are used to launch the CUDA kernels.
- */
- std::vector<dim3> block_dim;
-
/**
* Shared pointer to a Partitioner for distributed Vectors used in
* cell_loop. When MPI is not used the pointer is null.
data_host.row_start = data.row_start;
data_host.use_coloring = data.use_coloring;
- const unsigned int n_elements =
- data_host.n_cells * data_host.padding_length;
if (update_flags & update_quadrature_points)
{
data_host.q_points = Kokkos::create_mirror(data.q_points);
/*----------------------- Inline functions ---------------------------------*/
-# ifndef DOXYGEN
+#ifndef DOXYGEN
template <int dim, typename Number>
inline const std::vector<std::vector<
return *dof_handler;
}
-# endif
+#endif
} // namespace CUDAWrappers
DEAL_II_NAMESPACE_CLOSE
#endif
-#endif
#include <deal.II/base/config.h>
-#ifdef DEAL_II_WITH_CUDA
-
-# include <deal.II/base/cuda.h>
-# include <deal.II/base/cuda_size.h>
-# include <deal.II/base/graph_coloring.h>
-# include <deal.II/base/memory_space.h>
+#include <deal.II/base/cuda.h>
+#include <deal.II/base/cuda_size.h>
+#include <deal.II/base/graph_coloring.h>
+#include <deal.II/base/memory_space.h>
-# include <deal.II/dofs/dof_tools.h>
+#include <deal.II/dofs/dof_tools.h>
-# include <deal.II/fe/fe_dgq.h>
-# include <deal.II/fe/fe_values.h>
-# include <deal.II/fe/mapping_q1.h>
+#include <deal.II/fe/fe_dgq.h>
+#include <deal.II/fe/fe_values.h>
+#include <deal.II/fe/mapping_q1.h>
-# include <deal.II/matrix_free/cuda_hanging_nodes_internal.h>
-# include <deal.II/matrix_free/cuda_matrix_free.h>
-# include <deal.II/matrix_free/shape_info.h>
+#include <deal.II/matrix_free/cuda_hanging_nodes_internal.h>
+#include <deal.II/matrix_free/cuda_matrix_free.h>
+#include <deal.II/matrix_free/shape_info.h>
-# include <Kokkos_Core.hpp>
+#include <Kokkos_Core.hpp>
+#ifdef DEAL_II_WITH_CUDA
# include <cuda_runtime_api.h>
+#endif
-# include <cmath>
-# include <functional>
-# include <string>
+#include <cmath>
+#include <functional>
+#include <string>
DEAL_II_NAMESPACE_OPEN
// We need at least three colors when we are using CUDA-aware MPI and
// overlapping the communication
data->n_cells.resize(std::max(n_colors, 3U), 0);
- data->grid_dim.resize(n_colors);
- data->block_dim.resize(n_colors);
data->local_to_global.resize(n_colors);
data->constraint_mask.resize(n_colors);
}
};
+#ifdef DEAL_II_WITH_CUDA
template <>
struct VectorLocalSize<LinearAlgebra::CUDAWrappers::Vector<double>>
{
return vec.size();
}
};
+#endif
}
-
+#ifdef DEAL_II_WITH_CUDA
template <int dim, typename Number>
void
MatrixFree<dim, Number>::initialize_dof_vector(
{
vec.reinit(n_dofs);
}
+#endif
template <int dim, typename Number>
void
MatrixFree<dim, Number>::initialize_dof_vector(
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &vec) const
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &vec) const
{
if (partitioner)
vec.reinit(partitioner);
// First compute the size of n_cells, row_starts, kernel launch parameters,
// and constrained_dofs
std::size_t bytes = n_cells.size() * sizeof(unsigned int) * 2 +
- 2 * n_colors * sizeof(dim3) +
n_constrained_dofs * sizeof(unsigned int);
// For each color, add local_to_global, inv_jacobian, JxW, and q_points.
{
dof_handler = &dof_handler_;
- if (typeid(Number) == typeid(double))
- cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
-
UpdateFlags update_flags = additional_data.mapping_update_flags;
if (update_flags & update_gradients)
update_flags |= update_JxW_values;
template <typename Functor>
void
MatrixFree<dim, Number>::distributed_cell_loop(
- const Functor & func,
- const LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &src,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> &dst) const
+ const Functor &func,
+ const LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &src,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> &dst) const
{
MemorySpace::Default::kokkos_space::execution_space exec;
else
{
// Create the ghosted source and the ghosted destination
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>
ghosted_src(partitioner);
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>
ghosted_dst(ghosted_src);
ghosted_src = src;
ghosted_dst = dst;
+ ghosted_dst.zero_out_ghost_values();
// Execute the loop on the cells
for (unsigned int i = 0; i < n_colors; ++i)
team_policy(exec, n_cells[i], Kokkos::AUTO);
internal::ApplyKernel<dim, Number, Functor> apply_kernel(
- func, get_data(i), src.get_values(), dst.get_values());
+ func,
+ get_data(i),
+ ghosted_src.get_values(),
+ ghosted_dst.get_values());
Kokkos::parallel_for(
"dealii::MatrixFree::distributed_cell_loop_" +
}
}
-
-
+#ifdef DEAL_II_WITH_CUDA
template <int dim, typename Number>
template <typename Functor>
void
{
Assert(false, ExcInternalError());
}
+#endif
} // namespace CUDAWrappers
DEAL_II_NAMESPACE_CLOSE
#endif
-
-#endif
#include <deal.II/matrix_free/cuda_matrix_free.templates.h>
-#ifdef DEAL_II_WITH_CUDA
-
DEAL_II_NAMESPACE_OPEN
DEAL_II_NAMESPACE_CLOSE
#endif
-
-#endif
## ---------------------------------------------------------------------
set(_src
+ cuda_matrix_free.cc
dof_info.cc
evaluation_template_factory.cc
evaluation_template_factory_inst2.cc
${CMAKE_SOURCE_DIR}/include/deal.II/matrix_free/*.h
)
-if(DEAL_II_WITH_CUDA)
- set(_src
- cuda_matrix_free.cc
- ${_src}
- )
- set(_inst
- ${_inst}
- )
-endif()
-
define_object_library(object_matrix_free OBJECT ${_src} ${_header} ${_inst})
expand_instantiations(object_matrix_free "${_inst}")
#include <deal.II/matrix_free/cuda_matrix_free.templates.h>
-#ifdef DEAL_II_WITH_CUDA
-
DEAL_II_NAMESPACE_OPEN
} // namespace CUDAWrappers
DEAL_II_NAMESPACE_CLOSE
-
-#endif
ExcInternalError());
}
+#ifdef DEAL_II_WITH_CUDA
template <typename Number>
void
check(const LinearAlgebra::CUDAWrappers::Vector<Number> &vector,
{
AssertDimension(vector.size(), reference_partitioner.size());
}
+#endif
template <int dim, int fe_degree, typename VectorType>
void
unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD);
deallog.push(Utilities::int_to_string(myid));
- init_cuda(true);
+ // init_cuda(true);
MPILogInitAll mpi_inilog;
test<2,
1,
LinearAlgebra::distributed::Vector<double, MemorySpace::Default>>();
+
+#ifdef DEAL_II_WITH_CUDA
if (Utilities::MPI::n_mpi_processes(MPI_COMM_WORLD) == 1)
test<2, 1, LinearAlgebra::CUDAWrappers::Vector<double>>();
-
+#endif
return 0;
}
DEAL:0:0::locally relevant dofs :
{[0,24]}
DEAL:0:0::OK
-DEAL:0:0::locally owned dofs :
-{[0,24]}
-DEAL:0:0::locally relevant dofs :
-{[0,24]}
-DEAL:0:0::OK
--- /dev/null
+
+
+-
+-DEAL:0:0::locally owned dofs :
+-{[0,24]}
+-DEAL:0:0::locally relevant dofs :
+-{[0,24]}
+-DEAL:0:0::OK
+-DEAL:0:0::locally owned dofs :
+-{[0,24]}
+-DEAL:0:0::locally relevant dofs :
+-{[0,24]}
+-DEAL:0:0::OK
#include "../tests.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
void
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::CUDAWrappers::Vector<Number>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells());
}
#include "../tests.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
void
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::CUDAWrappers::Vector<Number>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells());
}
#include "../tests.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
void
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::CUDAWrappers::Vector<Number>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells());
}
#include "../tests.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
void
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::CUDAWrappers::Vector<Number>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells(), true, true);
}
#include "../tests.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::CUDAWrappers::Vector<Number>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells());
}
#include "../tests.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::CUDAWrappers::Vector<Number>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells());
}
#include "../tests.h"
#include "create_mesh.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
void
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::CUDAWrappers::Vector<Number>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells());
}
#include "../tests.h"
#include "create_mesh.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
void
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells());
}
#include "../tests.h"
#include "create_mesh.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
void
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells(), false);
}
#include "../tests.h"
#include "create_mesh.h"
-#include "matrix_vector_common.h"
+#include "matrix_vector_device_common.h"
template <int dim, int fe_degree, typename Number>
void
do_test<dim,
fe_degree,
Number,
- LinearAlgebra::CUDAWrappers::Vector<Number>,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>,
fe_degree + 1>(dof, constraints, tria.n_active_cells());
}
#include "../tests.h"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
const unsigned int coef_size =
tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim);
- MatrixFreeTest<dim,
- fe_degree,
- Number,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>
+ MatrixFreeTest<
+ dim,
+ fe_degree,
+ Number,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>
mf(mf_data, coef_size);
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev;
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev;
mf_data.initialize_dof_vector(in_dev);
mf_data.initialize_dof_vector(out_dev);
unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD);
deallog.push(Utilities::int_to_string(myid));
- init_cuda(true);
-
if (myid == 0)
{
initlog();
#include "../tests.h"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
template <int dim>
const unsigned int coef_size =
tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim);
- MatrixFreeTest<dim,
- fe_degree,
- Number,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>
+ MatrixFreeTest<
+ dim,
+ fe_degree,
+ Number,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>
mf(mf_data, coef_size, false);
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev;
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev;
mf_data.initialize_dof_vector(in_dev);
mf_data.initialize_dof_vector(out_dev);
LinearAlgebra::ReadWriteVector<Number> rw_in(owned_set);
unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD);
deallog.push(Utilities::int_to_string(myid));
- init_cuda(true);
-
if (myid == 0)
{
initlog();
#include "../tests.h"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
const unsigned int coef_size =
tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim);
- MatrixFreeTest<dim,
- fe_degree,
- Number,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>
+ MatrixFreeTest<
+ dim,
+ fe_degree,
+ Number,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>
mf(mf_data, coef_size);
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev(
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev(
owned_set, MPI_COMM_WORLD);
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev(
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev(
owned_set, MPI_COMM_WORLD);
LinearAlgebra::ReadWriteVector<Number> rw_in(owned_set);
unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD);
deallog.push(Utilities::int_to_string(myid));
- init_cuda(true);
-
if (myid == 0)
{
initlog();
#include "../tests.h"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
const unsigned int coef_size =
tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim);
- MatrixFreeTest<dim,
- fe_degree,
- Number,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>
+ MatrixFreeTest<
+ dim,
+ fe_degree,
+ Number,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>
mf(mf_data, coef_size);
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev;
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev;
mf_data.initialize_dof_vector(in_dev);
mf_data.initialize_dof_vector(out_dev);
unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD);
deallog.push(Utilities::int_to_string(myid));
- init_cuda(true);
-
if (myid == 0)
{
initlog();
#include "../tests.h"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
const unsigned int coef_size =
tria.n_locally_owned_active_cells() * std::pow(fe_degree + 1, dim);
- MatrixFreeTest<dim,
- fe_degree,
- Number,
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA>>
+ MatrixFreeTest<
+ dim,
+ fe_degree,
+ Number,
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default>>
mf(mf_data, coef_size);
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> in_dev;
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> out_dev;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> in_dev;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> out_dev;
mf_data.initialize_dof_vector(in_dev);
mf_data.initialize_dof_vector(out_dev);
unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD);
deallog.push(Utilities::int_to_string(myid));
- init_cuda(true);
-
if (myid == 0)
{
initlog();
#include "../tests.h"
#include "Kokkos_Core.hpp"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
template <int fe_degree, int n_q_points_1d>
void
-do_test(const DoFHandler<2> & dof,
- MatrixFreeTest<2,
- fe_degree,
- double,
- LinearAlgebra::CUDAWrappers::Vector<double>,
- n_q_points_1d> & mf,
+do_test(const DoFHandler<2> &dof,
+ MatrixFreeTest<
+ 2,
+ fe_degree,
+ double,
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Default>,
+ n_q_points_1d> & mf,
unsigned int n_dofs,
MappingQ<2> & mapping,
const AffineConstraints<double> &constraints)
{
- Vector<double> in_host(n_dofs), out_host(n_dofs);
- LinearAlgebra::ReadWriteVector<double> in(n_dofs), out(n_dofs);
- LinearAlgebra::CUDAWrappers::Vector<double> in_device(n_dofs);
- LinearAlgebra::CUDAWrappers::Vector<double> out_device(n_dofs);
+ Vector<double> in_host(n_dofs), out_host(n_dofs);
+ LinearAlgebra::ReadWriteVector<double> in(n_dofs), out(n_dofs);
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Default> in_device(
+ n_dofs);
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Default> out_device(
+ n_dofs);
for (unsigned int i = 0; i < n_dofs; ++i)
{
in_device.import_elements(in, VectorOperation::insert);
mf.vmult(out_device, in_device);
- cudaDeviceSynchronize();
+ Kokkos::fence();
out.import_elements(out_device, VectorOperation::insert);
// assemble sparse matrix with (\nabla v, \nabla u) + (v, 10 * u)
deallog << std::setprecision(3);
Kokkos::initialize();
- init_cuda();
{
Triangulation<2> tria;
const QGauss<1> quad_1(n_q_points_1d_1);
mf_data_1.reinit(mapping_1, dof_1, constraints, quad_1, additional_data_1);
const unsigned int n_dofs_1 = dof_1.n_dofs();
- MatrixFreeTest<2,
- fe_degree_1,
- double,
- LinearAlgebra::CUDAWrappers::Vector<double>,
- n_q_points_1d_1>
+ MatrixFreeTest<
+ 2,
+ fe_degree_1,
+ double,
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Default>,
+ n_q_points_1d_1>
mf_1(mf_data_1,
n_dofs_1 * std::pow(n_q_points_1d_1, 2),
constant_coefficient);
const QGauss<1> quad_2(n_q_points_1d_2);
mf_data_2.reinit(mapping_2, dof_2, constraints, quad_2, additional_data_2);
const unsigned int n_dofs_2 = dof_2.n_dofs();
- MatrixFreeTest<2,
- fe_degree_2,
- double,
- LinearAlgebra::CUDAWrappers::Vector<double>,
- n_q_points_1d_2>
+ MatrixFreeTest<
+ 2,
+ fe_degree_2,
+ double,
+ LinearAlgebra::distributed::Vector<double, MemorySpace::Default>,
+ n_q_points_1d_2>
mf_2(mf_data_2,
n_dofs_2 * std::pow(n_q_points_1d_2, 2),
constant_coefficient);
void
test() const
{
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> dst_dummy;
- LinearAlgebra::distributed::Vector<Number, MemorySpace::CUDA> src_dummy;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> dst_dummy;
+ LinearAlgebra::distributed::Vector<Number, MemorySpace::Default> src_dummy;
data.cell_loop(*this, src_dummy, dst_dummy);
- // Check that the kernel was launched correctly
- AssertCuda(cudaPeekAtLastError());
- // Check that there was no problem during the execution of the kernel
- AssertCuda(cudaDeviceSynchronize());
+ Kokkos::fence();
deallog << "OK" << std::endl;
};
initlog();
Kokkos::initialize();
- init_cuda();
test<2, 1>();
unsigned int myid = Utilities::MPI::this_mpi_process(MPI_COMM_WORLD);
deallog.push(Utilities::int_to_string(myid));
- init_cuda(true);
MPILogInitAll mpi_inilog;
test<2, 1>();
#include "../tests.h"
#include "Kokkos_Core.hpp"
-#include "matrix_vector_mf.h"
+#include "matrix_vector_device_mf.h"
// forward declare this function. will be implemented in .cc files
in_device.import_elements(in, VectorOperation::insert);
mf.vmult(out_device, in_device);
- cudaDeviceSynchronize();
+ Kokkos::fence();
out.import_elements(out_device, VectorOperation::insert);
// assemble sparse matrix with (\nabla v, \nabla u) + (v, 10 * u)
Kokkos::initialize();
- init_cuda();
-
{
deallog.push("double");
deallog.push("2d");
projects / dealii.git / commitdiff