CUDA Wrappers

The classes in this module are concerned with the description of features to be run on GPUs using CUDA. More...

Collaboration diagram for CUDA Wrappers:

Classes
struct	LinearAlgebra::CUDAWrappers::kernel::Binop_Addition< Number >
struct	LinearAlgebra::CUDAWrappers::kernel::Binop_Subtraction< Number >
struct	LinearAlgebra::CUDAWrappers::kernel::Binop_Max< Number >
struct	LinearAlgebra::CUDAWrappers::kernel::Binop_Min< Number >
struct	LinearAlgebra::CUDAWrappers::kernel::ElemSum< Number >
struct	LinearAlgebra::CUDAWrappers::kernel::L1Norm< Number >
struct	LinearAlgebra::CUDAWrappers::kernel::LInfty< Number >
struct	LinearAlgebra::CUDAWrappers::kernel::DotProduct< Number >
class	CUDAWrappers::PreconditionIC< Number >
class	CUDAWrappers::PreconditionILU< Number >
class	CUDAWrappers::SolverDirect< Number >
class	CUDAWrappers::FEEvaluation< dim, fe_degree, n_q_points_1d, n_components_, Number >
class	CUDAWrappers::MatrixFree< dim, Number >
struct	CUDAWrappers::internal::EvaluatorTensorProduct< variant, dim, fe_degree, n_q_points_1d, Number >
struct	CUDAWrappers::internal::EvaluatorTensorProduct< evaluate_general, dim, fe_degree, n_q_points_1d, Number >

Enumerations
enum	CUDAWrappers::internal::EvaluatorVariant { CUDAWrappers::internal::evaluate_general, CUDAWrappers::internal::evaluate_symmetric, CUDAWrappers::internal::evaluate_evenodd }

Functions
float	LinearAlgebra::CUDAWrappers::atomicAdd_wrapper (float *address, float val)
double	LinearAlgebra::CUDAWrappers::atomicAdd_wrapper (double *address, double val)
float	LinearAlgebra::CUDAWrappers::atomicMax_wrapper (float *address, float val)
double	LinearAlgebra::CUDAWrappers::atomicMax_wrapper (double *address, double val)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::vec_scale (Number *val, const Number a, const size_type N)
template<typename Number , template< typename > class Binop>
__global__ void	LinearAlgebra::CUDAWrappers::kernel::vector_bin_op (Number *v1, const Number *v2, const size_type N)
template<typename Number , template< typename > class Binop>
__global__ void	LinearAlgebra::CUDAWrappers::kernel::masked_vector_bin_op (const unsigned int *mask, Number *v1, const Number *v2, const size_type N)
template<typename Number , typename Operation >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::reduction (Number *result, const Number *v, const size_type N)
template<typename Number , typename Operation >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::double_vector_reduction (Number *result, const Number *v1, const Number *v2, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::vec_add (Number *val, const Number a, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::add_aV (Number *val, const Number a, const Number *V_val, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::add_aVbW (Number *val, const Number a, const Number *V_val, const Number b, const Number *W_val, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::sadd (const Number s, Number *val, const Number a, const Number *V_val, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::sadd (const Number s, Number *val, const Number a, const Number *V_val, const Number b, const Number *W_val, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::scale (Number *val, const Number *V_val, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::equ (Number *val, const Number a, const Number *V_val, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::equ (Number *val, const Number a, const Number *V_val, const Number b, const Number *W_val, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::add_and_dot (Number *res, Number *v1, const Number *v2, const Number *v3, const Number a, const size_type N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::set (Number *val, const Number s, const size_type N)
template<typename Number , typename IndexType >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::set_permutated (const IndexType *indices, Number *val, const Number *v, const IndexType N)
template<typename Number , typename IndexType >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::gather (Number *val, const IndexType *indices, const Number *v, const IndexType N)
template<typename Number >
__global__ void	LinearAlgebra::CUDAWrappers::kernel::add_permutated (const size_type *indices, Number *val, const Number *v, const size_type N)

Detailed Description

The classes in this module are concerned with the description of features to be run on GPUs using CUDA.

Enumeration Type Documentation

EvaluatorVariant

enum CUDAWrappers::internal::EvaluatorVariant

In this namespace, the evaluator routines that evaluate the tensor products are implemented.

Enumerator
evaluate_general
evaluate_symmetric
evaluate_evenodd

Definition at line 42 of file cuda_tensor_product_kernels.h.

Function Documentation

atomicAdd_wrapper() [1/2]

float LinearAlgebra::CUDAWrappers::atomicAdd_wrapper ( float *  address, float  val  )

inline

Provide atomicAdd for floats.

Definition at line 35 of file cuda_atomic.h.

atomicAdd_wrapper() [2/2]

double LinearAlgebra::CUDAWrappers::atomicAdd_wrapper ( double *  address, double  val  )

inline

Provide atomicAdd for doubles.

Definition at line 48 of file cuda_atomic.h.

atomicMax_wrapper() [1/2]

float LinearAlgebra::CUDAWrappers::atomicMax_wrapper ( float *  address, float  val  )

inline

Provide atomicMax for floats.

Definition at line 81 of file cuda_atomic.h.

atomicMax_wrapper() [2/2]

double LinearAlgebra::CUDAWrappers::atomicMax_wrapper ( double *  address, double  val  )

inline

Provide atomicMax for doubles.

Definition at line 105 of file cuda_atomic.h.

vec_scale()

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::vec_scale	(	Number *	val,
		const Number	a,
		const size_type	N
	)

Multiply each entry of val of size N by a.

vector_bin_op()

template<typename Number , template< typename > class Binop>

__global__ void LinearAlgebra::CUDAWrappers::kernel::vector_bin_op	(	Number *	v1,
		const Number *	v2,
		const size_type	N
	)

Apply the functor Binop to each element of v1 and v2.

masked_vector_bin_op()

template<typename Number , template< typename > class Binop>

__global__ void LinearAlgebra::CUDAWrappers::kernel::masked_vector_bin_op	(	const unsigned int *	mask,
		Number *	v1,
		const Number *	v2,
		const size_type	N
	)

Apply the functor Binop to the elements of v1 that have indices in mask and v2. The size of mask should be greater than the size of v1. mask and v2 should have the same size N.

reduction()

template<typename Number , typename Operation >

__global__ void LinearAlgebra::CUDAWrappers::kernel::reduction	(	Number *	result,
		const Number *	v,
		const size_type	N
	)

Perform a reduction on v using Operation.

double_vector_reduction()

template<typename Number , typename Operation >

__global__ void LinearAlgebra::CUDAWrappers::kernel::double_vector_reduction	(	Number *	result,
		const Number *	v1,
		const Number *	v2,
		const size_type	N
	)

Perform a binary operation on each element of v1 and v2 followed by reduction on the resulting array.

vec_add()

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::vec_add	(	Number *	val,
		const Number	a,
		const size_type	N
	)

Add a to each element of val.

add_aV()

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::add_aV	(	Number *	val,
		const Number	a,
		const Number *	V_val,
		const size_type	N
	)

Addition of a multiple of a vector, i.e., val += a*V_val.

add_aVbW()

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::add_aVbW	(	Number *	val,
		const Number	a,
		const Number *	V_val,
		const Number	b,
		const Number *	W_val,
		const size_type	N
	)

Addition of multiple scaled vector, i.e., val += a*V_val + b*W_val.

sadd() [1/2]

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::sadd	(	const Number	s,
		Number *	val,
		const Number	a,
		const Number *	V_val,
		const size_type	N
	)

Scaling and simple addition of a multiple of a vector, i.e. val = = s*val + a*V_val

sadd() [2/2]

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::sadd	(	const Number	s,
		Number *	val,
		const Number	a,
		const Number *	V_val,
		const Number	b,
		const Number *	W_val,
		const size_type	N
	)

Scaling and multiple additions of scaled vectors, i.e. val = = s*val + a*V_val + b*W_val

scale()

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::scale	(	Number *	val,
		const Number *	V_val,
		const size_type	N
	)

Scale each element of this vector by the corresponding element in the argument.

equ() [1/2]

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::equ	(	Number *	val,
		const Number	a,
		const Number *	V_val,
		const size_type	N
	)

Assignment val = a*V_val.

equ() [2/2]

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::equ	(	Number *	val,
		const Number	a,
		const Number *	V_val,
		const Number	b,
		const Number *	W_val,
		const size_type	N
	)

Assignment val = a*V_val + b*W_val.

add_and_dot()

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::add_and_dot	(	Number *	res,
		Number *	v1,
		const Number *	v2,
		const Number *	v3,
		const Number	a,
		const size_type	N
	)

Perform a combined operation of a vector addition and a subsequent inner product, returning the value of the inner product.

set()

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::set	(	Number *	val,
		const Number	s,
		const size_type	N
	)

Set each element of val to s.

set_permutated()

template<typename Number , typename IndexType >

__global__ void LinearAlgebra::CUDAWrappers::kernel::set_permutated	(	const IndexType *	indices,
		Number *	val,
		const Number *	v,
		const IndexType	N
	)

Set each element in val to v using indices as permutation, i.e., val[indices[i]] = v[i].

gather()

template<typename Number , typename IndexType >

__global__ void LinearAlgebra::CUDAWrappers::kernel::gather	(	Number *	val,
		const IndexType *	indices,
		const Number *	v,
		const IndexType	N
	)

Set each element in val to v using indices as permutation, i.e., val[i] = v[indices[i]].

add_permutated()

template<typename Number >

__global__ void LinearAlgebra::CUDAWrappers::kernel::add_permutated	(	const size_type *	indices,
		Number *	val,
		const Number *	v,
		const size_type	N
	)

Add each element in val to v using indices as permutation, i.e., val[indices[i]] += v[i].