Reference documentation for deal.II version Git 9585ce37e8 2019-11-21 18:02:11 -0500
\(\newcommand{\dealcoloneq}{\mathrel{\vcenter{:}}=}\)
Classes | Public Types | Public Member Functions | Private Member Functions | Private Attributes | List of all members
CUDAWrappers::MatrixFree< dim, Number > Class Template Reference

#include <deal.II/matrix_free/cuda_matrix_free.h>

Inheritance diagram for CUDAWrappers::MatrixFree< dim, Number >:
[legend]

Classes

struct  AdditionalData
 
struct  Data
 

Public Types

enum  ParallelizationScheme
 

Public Member Functions

 MatrixFree ()
 
unsigned int get_padding_length () const
 
void reinit (const Mapping< dim > &mapping, const DoFHandler< dim > &dof_handler, const AffineConstraints< Number > &constraints, const Quadrature< 1 > &quad, const AdditionalData &additional_data=AdditionalData())
 
void reinit (const DoFHandler< dim > &dof_handler, const AffineConstraints< Number > &constraints, const Quadrature< 1 > &quad, const AdditionalData &AdditionalData=AdditionalData())
 
Data get_data (unsigned int color) const
 
template<typename Functor , typename VectorType >
void cell_loop (const Functor &func, const VectorType &src, VectorType &dst) const
 
template<typename Functor >
void evaluate_coefficients (Functor func) const
 
template<typename VectorType >
void copy_constrained_values (const VectorType &src, VectorType &dst) const
 
template<typename VectorType >
void set_constrained_values (const Number val, VectorType &dst) const
 
void initialize_dof_vector (LinearAlgebra::CUDAWrappers::Vector< Number > &vec) const
 
void initialize_dof_vector (LinearAlgebra::distributed::Vector< Number, MemorySpace::CUDA > &vec) const
 
void free ()
 
std::size_t memory_consumption () const
 
- Public Member Functions inherited from Subscriptor
 Subscriptor ()
 
 Subscriptor (const Subscriptor &)
 
 Subscriptor (Subscriptor &&) noexcept
 
virtual ~Subscriptor ()
 
Subscriptoroperator= (const Subscriptor &)
 
Subscriptoroperator= (Subscriptor &&) noexcept
 
void subscribe (std::atomic< bool > *const validity, const std::string &identifier="") const
 
void unsubscribe (std::atomic< bool > *const validity, const std::string &identifier="") const
 
unsigned int n_subscriptions () const
 
template<typename StreamType >
void list_subscribers (StreamType &stream) const
 
void list_subscribers () const
 
template<class Archive >
void serialize (Archive &ar, const unsigned int version)
 

Private Member Functions

void internal_reinit (const Mapping< dim > &mapping, const DoFHandler< dim > &dof_handler, const AffineConstraints< Number > &constraints, const Quadrature< 1 > &quad, std::shared_ptr< const MPI_Comm > comm, const AdditionalData additional_data)
 
template<typename Functor , typename VectorType >
void serial_cell_loop (const Functor &func, const VectorType &src, VectorType &dst) const
 
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
 
template<typename Functor >
void distributed_cell_loop (const Functor &func, const LinearAlgebra::CUDAWrappers::Vector< Number > &src, LinearAlgebra::CUDAWrappers::Vector< Number > &dst) const
 
template<typename VectorType >
void serial_copy_constrained_values (const VectorType &src, VectorType &dst) const
 
void distributed_copy_constrained_values (const LinearAlgebra::distributed::Vector< Number, MemorySpace::CUDA > &src, LinearAlgebra::distributed::Vector< Number, MemorySpace::CUDA > &dst) const
 
void distributed_copy_constrained_values (const LinearAlgebra::CUDAWrappers::Vector< Number > &src, LinearAlgebra::CUDAWrappers::Vector< Number > &dst) const
 
template<typename VectorType >
void serial_set_constrained_values (const Number val, VectorType &dst) const
 
void distributed_set_constrained_values (const Number val, LinearAlgebra::distributed::Vector< Number, MemorySpace::CUDA > &dst) const
 
void distributed_set_constrained_values (const Number val, LinearAlgebra::CUDAWrappers::Vector< Number > &dst) const
 

Private Attributes

ParallelizationScheme parallelization_scheme
 
bool use_coloring
 
types::global_dof_index n_dofs
 
unsigned int fe_degree
 
unsigned int dofs_per_cell
 
unsigned int n_constrained_dofs
 
unsigned int q_points_per_cell
 
unsigned int n_colors
 
std::vector< unsigned int > n_cells
 
std::vector< point_type * > q_points
 
std::vector< types::global_dof_index * > local_to_global
 
std::vector< Number * > inv_jacobian
 
std::vector< Number * > JxW
 
types::global_dof_indexconstrained_dofs
 
std::vector< unsigned int * > constraint_mask
 
std::vector< dim3 > grid_dim
 
std::vector< dim3 > block_dim
 
std::shared_ptr< const Utilities::MPI::Partitionerpartitioner
 
unsigned int cells_per_block
 
dim3 constraint_grid_dim
 
dim3 constraint_block_dim
 
unsigned int padding_length
 
std::vector< unsigned int > row_start
 

Additional Inherited Members

- Static Public Member Functions inherited from Subscriptor
static ::ExceptionBaseExcInUse (int arg1, std::string arg2, std::string arg3)
 
static ::ExceptionBaseExcNoSubscriber (std::string arg1, std::string arg2)
 

Detailed Description

template<int dim, typename Number = double>
class CUDAWrappers::MatrixFree< dim, Number >

This class collects all the data that is stored for the matrix free implementation. The storage scheme is tailored towards several loops performed with the same data, i.e., typically doing many matrix-vector products or residual computations on the same mesh.

This class does not implement any operations involving finite element basis functions, i.e., regarding the operation performed on the cells. For these operations, the class FEEvaluation is designed to use the data collected in this class.

This class implements a loop over all cells (cell_loop()). This loop is scheduled in such a way that cells that share degrees of freedom are not worked on simultaneously, which implies that it is possible to write to vectors in parallel without having to explicitly synchronize access to these vectors and matrices. This class does not implement any shape values, all it does is to cache the respective data. To implement finite element operations, use the class CUDAWrappers::FEEvalutation.

This class traverse the cells in a different order than the usual Triangulation class in deal.II.

Note
Only float and double are supported.

Definition at line 80 of file cuda_matrix_free.h.

Member Enumeration Documentation

◆ ParallelizationScheme

template<int dim, typename Number = double>
enum CUDAWrappers::MatrixFree::ParallelizationScheme

Parallelization scheme used: parallel_in_elem (parallelism at the level of degrees of freedom) or parallel_over_elem (parallelism at the level of cells)

Definition at line 91 of file cuda_matrix_free.h.

Constructor & Destructor Documentation

◆ MatrixFree()

template<int dim, typename Number = double>
CUDAWrappers::MatrixFree< dim, Number >::MatrixFree ( )

Default constructor.

Member Function Documentation

◆ get_padding_length()

template<int dim, typename Number = double>
unsigned int CUDAWrappers::MatrixFree< dim, Number >::get_padding_length ( ) const

Return the length of the padding.

◆ reinit() [1/2]

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::reinit ( const Mapping< dim > &  mapping,
const DoFHandler< dim > &  dof_handler,
const AffineConstraints< Number > &  constraints,
const Quadrature< 1 > &  quad,
const AdditionalData additional_data = AdditionalData() 
)

Extracts the information needed to perform loops over cells. The DoFHandler and AffineConstraints objects describe the layout of degrees of freedom, the DoFHandler and the mapping describe the transformation from unit to real cell, and the finite element underlying the DoFHandler together with the quadrature formula describe the local operations.

◆ reinit() [2/2]

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::reinit ( const DoFHandler< dim > &  dof_handler,
const AffineConstraints< Number > &  constraints,
const Quadrature< 1 > &  quad,
const AdditionalData AdditionalData = AdditionalData() 
)

Initializes the data structures. Same as above but using a Q1 mapping.

◆ get_data()

template<int dim, typename Number = double>
Data CUDAWrappers::MatrixFree< dim, Number >::get_data ( unsigned int  color) const

Return the Data structure associated with color.

◆ cell_loop()

template<int dim, typename Number = double>
template<typename Functor , typename VectorType >
void CUDAWrappers::MatrixFree< dim, Number >::cell_loop ( const Functor &  func,
const VectorType &  src,
VectorType &  dst 
) const

This method runs the loop over all cells and apply the local operation on each element in parallel. func is a functor which is applied on each color.

func needs to define

__device__ void operator()(
const unsigned int cell,
const Number * src,
Number * dst) const;
static const unsigned int n_dofs_1d;
static const unsigned int n_local_dofs;
static const unsigned int n_q_points;

◆ evaluate_coefficients()

template<int dim, typename Number = double>
template<typename Functor >
void CUDAWrappers::MatrixFree< dim, Number >::evaluate_coefficients ( Functor  func) const

This method runs the loop over all cells and apply the local operation on each element in parallel. This function is very similar to cell_loop() but it uses a simpler functor.

func needs to define

__device__ void operator()(
const unsigned int cell,
static const unsigned int n_dofs_1d;
static const unsigned int n_local_dofs;
static const unsigned int n_q_points;

◆ copy_constrained_values()

template<int dim, typename Number = double>
template<typename VectorType >
void CUDAWrappers::MatrixFree< dim, Number >::copy_constrained_values ( const VectorType &  src,
VectorType &  dst 
) const

Copy the values of the constrained entries from src to dst. This is used to impose zero Dirichlet boundary condition.

◆ set_constrained_values()

template<int dim, typename Number = double>
template<typename VectorType >
void CUDAWrappers::MatrixFree< dim, Number >::set_constrained_values ( const Number  val,
VectorType &  dst 
) const

Set the entries in dst corresponding to constrained values to val. The main purpose of this function is to set the constrained entries of the source vector used in cell_loop() to zero.

◆ initialize_dof_vector() [1/2]

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::initialize_dof_vector ( LinearAlgebra::CUDAWrappers::Vector< Number > &  vec) const

Initialize a serial vector. The size corresponds to the number of degrees of freedom in the DoFHandler object.

◆ initialize_dof_vector() [2/2]

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::initialize_dof_vector ( LinearAlgebra::distributed::Vector< Number, MemorySpace::CUDA > &  vec) const

Initialize a distributed vector. The local elements correspond to the locally owned degrees of freedom and the ghost elements correspond to the (additional) locally relevant dofs.

◆ free()

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::free ( )

Free all the memory allocated.

◆ memory_consumption()

template<int dim, typename Number = double>
std::size_t CUDAWrappers::MatrixFree< dim, Number >::memory_consumption ( ) const

Return an approximation of the memory consumption of this class in bytes.

◆ internal_reinit()

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::internal_reinit ( const Mapping< dim > &  mapping,
const DoFHandler< dim > &  dof_handler,
const AffineConstraints< Number > &  constraints,
const Quadrature< 1 > &  quad,
std::shared_ptr< const MPI_Comm >  comm,
const AdditionalData  additional_data 
)
private

Initializes the data structures.

◆ serial_cell_loop()

template<int dim, typename Number = double>
template<typename Functor , typename VectorType >
void CUDAWrappers::MatrixFree< dim, Number >::serial_cell_loop ( const Functor &  func,
const VectorType &  src,
VectorType &  dst 
) const
private

Helper function. Loop over all the cells and apply the functor on each element in parallel. This function is used when MPI is not used.

◆ distributed_cell_loop() [1/2]

template<int dim, typename Number = double>
template<typename Functor >
void CUDAWrappers::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
private

Helper function. Loop over all the cells and apply the functor on each element in parallel. This function is used when MPI is used.

◆ distributed_cell_loop() [2/2]

template<int dim, typename Number = double>
template<typename Functor >
void CUDAWrappers::MatrixFree< dim, Number >::distributed_cell_loop ( const Functor &  func,
const LinearAlgebra::CUDAWrappers::Vector< Number > &  src,
LinearAlgebra::CUDAWrappers::Vector< Number > &  dst 
) const
private

This function should never be called. Calling it results in an internal error. This function exists only because cell_loop needs distributed_cell_loop() to exist for LinearAlgebra::CUDAWrappers::Vector.

◆ serial_copy_constrained_values()

template<int dim, typename Number = double>
template<typename VectorType >
void CUDAWrappers::MatrixFree< dim, Number >::serial_copy_constrained_values ( const VectorType &  src,
VectorType &  dst 
) const
private

Helper function. Copy the values of the constrained entries of src to dst. This function is used when MPI is not used.

◆ distributed_copy_constrained_values() [1/2]

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::distributed_copy_constrained_values ( const LinearAlgebra::distributed::Vector< Number, MemorySpace::CUDA > &  src,
LinearAlgebra::distributed::Vector< Number, MemorySpace::CUDA > &  dst 
) const
private

Helper function. Copy the values of the constrained entries of src to dst. This function is used when MPI is used.

◆ distributed_copy_constrained_values() [2/2]

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::distributed_copy_constrained_values ( const LinearAlgebra::CUDAWrappers::Vector< Number > &  src,
LinearAlgebra::CUDAWrappers::Vector< Number > &  dst 
) const
private

This function should never be called. Calling it results in an internal error. This function exists only because copy_constrained_values needs distributed_copy_constrained_values() to exist for LinearAlgebra::CUDAWrappers::Vector.

◆ serial_set_constrained_values()

template<int dim, typename Number = double>
template<typename VectorType >
void CUDAWrappers::MatrixFree< dim, Number >::serial_set_constrained_values ( const Number  val,
VectorType &  dst 
) const
private

Helper function. Set the constrained entries of dst to val. This function is used when MPI is not used.

◆ distributed_set_constrained_values() [1/2]

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::distributed_set_constrained_values ( const Number  val,
LinearAlgebra::distributed::Vector< Number, MemorySpace::CUDA > &  dst 
) const
private

Helper function. Set the constrained entries of dst to val. This function is used when MPI is used.

◆ distributed_set_constrained_values() [2/2]

template<int dim, typename Number = double>
void CUDAWrappers::MatrixFree< dim, Number >::distributed_set_constrained_values ( const Number  val,
LinearAlgebra::CUDAWrappers::Vector< Number > &  dst 
) const
private

This function should never be called. Calling it results in an internal error. This function exists only because set_constrained_values needs distributed_set_constrained_values() to exist for LinearAlgebra::CUDAWrappers::Vector.

Member Data Documentation

◆ parallelization_scheme

template<int dim, typename Number = double>
ParallelizationScheme CUDAWrappers::MatrixFree< dim, Number >::parallelization_scheme
private

Parallelization scheme used, parallelization over degrees of freedom or over cells.

Definition at line 436 of file cuda_matrix_free.h.

◆ use_coloring

template<int dim, typename Number = double>
bool CUDAWrappers::MatrixFree< dim, Number >::use_coloring
private

If true, use graph coloring. Otherwise, use atomic operations. Graph coloring ensures bitwise reproducibility but is slower on Pascal and newer architectures.

Definition at line 443 of file cuda_matrix_free.h.

◆ n_dofs

template<int dim, typename Number = double>
types::global_dof_index CUDAWrappers::MatrixFree< dim, Number >::n_dofs
private

Total number of degrees of freedom.

Definition at line 448 of file cuda_matrix_free.h.

◆ fe_degree

template<int dim, typename Number = double>
unsigned int CUDAWrappers::MatrixFree< dim, Number >::fe_degree
private

Degree of the finite element used.

Definition at line 453 of file cuda_matrix_free.h.

◆ dofs_per_cell

template<int dim, typename Number = double>
unsigned int CUDAWrappers::MatrixFree< dim, Number >::dofs_per_cell
private

Number of degrees of freedom per cell.

Definition at line 458 of file cuda_matrix_free.h.

◆ n_constrained_dofs

template<int dim, typename Number = double>
unsigned int CUDAWrappers::MatrixFree< dim, Number >::n_constrained_dofs
private

Number of constrained degrees of freedom.

Definition at line 463 of file cuda_matrix_free.h.

◆ q_points_per_cell

template<int dim, typename Number = double>
unsigned int CUDAWrappers::MatrixFree< dim, Number >::q_points_per_cell
private

Number of quadrature points per cells.

Definition at line 468 of file cuda_matrix_free.h.

◆ n_colors

template<int dim, typename Number = double>
unsigned int CUDAWrappers::MatrixFree< dim, Number >::n_colors
private

Number of colors produced by the graph coloring algorithm.

Definition at line 473 of file cuda_matrix_free.h.

◆ n_cells

template<int dim, typename Number = double>
std::vector<unsigned int> CUDAWrappers::MatrixFree< dim, Number >::n_cells
private

Number of cells in each color.

Definition at line 478 of file cuda_matrix_free.h.

◆ q_points

template<int dim, typename Number = double>
std::vector<point_type *> CUDAWrappers::MatrixFree< dim, Number >::q_points
private

Vector of pointers to the quadrature points associated to the cells of each color.

Definition at line 484 of file cuda_matrix_free.h.

◆ local_to_global

template<int dim, typename Number = double>
std::vector<types::global_dof_index *> CUDAWrappers::MatrixFree< dim, Number >::local_to_global
private

Map the position in the local vector to the position in the global vector.

Definition at line 490 of file cuda_matrix_free.h.

◆ inv_jacobian

template<int dim, typename Number = double>
std::vector<Number *> CUDAWrappers::MatrixFree< dim, Number >::inv_jacobian
private

Vector of pointer to the inverse Jacobian associated to the cells of each color.

Definition at line 496 of file cuda_matrix_free.h.

◆ JxW

template<int dim, typename Number = double>
std::vector<Number *> CUDAWrappers::MatrixFree< dim, Number >::JxW
private

Vector of pointer to the Jacobian times the weights associated to the cells of each color.

Definition at line 502 of file cuda_matrix_free.h.

◆ constrained_dofs

template<int dim, typename Number = double>
types::global_dof_index* CUDAWrappers::MatrixFree< dim, Number >::constrained_dofs
private

Pointer to the constrained degrees of freedom.

Definition at line 507 of file cuda_matrix_free.h.

◆ constraint_mask

template<int dim, typename Number = double>
std::vector<unsigned int *> CUDAWrappers::MatrixFree< dim, Number >::constraint_mask
private

Mask deciding where constraints are set on a given cell.

Definition at line 512 of file cuda_matrix_free.h.

◆ grid_dim

template<int dim, typename Number = double>
std::vector<dim3> CUDAWrappers::MatrixFree< dim, Number >::grid_dim
private

Grid dimensions associated to the different colors. The grid dimensions are used to launch the CUDA kernels.

Definition at line 518 of file cuda_matrix_free.h.

◆ block_dim

template<int dim, typename Number = double>
std::vector<dim3> CUDAWrappers::MatrixFree< dim, Number >::block_dim
private

Block dimensions associated to the different colors. The block dimensions are used to launch the CUDA kernels.

Definition at line 524 of file cuda_matrix_free.h.

◆ partitioner

template<int dim, typename Number = double>
std::shared_ptr<const Utilities::MPI::Partitioner> CUDAWrappers::MatrixFree< dim, Number >::partitioner
private

Shared pointer to a Partitioner for distributed Vectors used in cell_loop. When MPI is not used the pointer is null.

Definition at line 530 of file cuda_matrix_free.h.

◆ cells_per_block

template<int dim, typename Number = double>
unsigned int CUDAWrappers::MatrixFree< dim, Number >::cells_per_block
private

Cells per block (determined by the function cells_per_block_shmem() ).

Definition at line 535 of file cuda_matrix_free.h.

◆ constraint_grid_dim

template<int dim, typename Number = double>
dim3 CUDAWrappers::MatrixFree< dim, Number >::constraint_grid_dim
private

Grid dimensions used to launch the CUDA kernels in *_constrained_values-operations.

Definition at line 541 of file cuda_matrix_free.h.

◆ constraint_block_dim

template<int dim, typename Number = double>
dim3 CUDAWrappers::MatrixFree< dim, Number >::constraint_block_dim
private

Block dimensions used to launch the CUDA kernels in *_constrained_values-operations.

Definition at line 547 of file cuda_matrix_free.h.

◆ padding_length

template<int dim, typename Number = double>
unsigned int CUDAWrappers::MatrixFree< dim, Number >::padding_length
private

Length of the padding (closest power of two larger than or equal to the number of thread).

Definition at line 553 of file cuda_matrix_free.h.

◆ row_start

template<int dim, typename Number = double>
std::vector<unsigned int> CUDAWrappers::MatrixFree< dim, Number >::row_start
private

Row start of each color.

Definition at line 558 of file cuda_matrix_free.h.


The documentation for this class was generated from the following file: