Functions
__global__ void	amg_influence_trivial_kernel (const unsigned int row_indices, const unsigned int column_indices, unsigned int size1, unsigned int nnz, unsigned int influences_row, unsigned int influences_id, unsigned int *influences_values)

template<typename NumericT >
void	amg_influence_trivial (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	Routine for taking all connections in the matrix as strong. More...

template<typename NumericT >
void	amg_influence_advanced (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	Routine for extracting strongly connected points considering a user-provided threshold value. More...

template<typename NumericT >
void	amg_influence (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	Dispatcher for influence processing. More...

void	enumerate_coarse_points (viennacl::linalg::detail::amg::amg_level_context &amg_context)
	Assign IDs to coarse points. More...

template<typename IndexT >
__global__ void	amg_pmis2_init_workdata (IndexT work_state, IndexT work_random, IndexT work_index, IndexT const point_types, IndexT const *random_weights, unsigned int size)
	CUDA kernel initializing the work vectors at each PMIS iteration. More...

template<typename IndexT >
__global__ void	amg_pmis2_max_neighborhood (IndexT const work_state, IndexT const work_random, IndexT const work_index, IndexT work_state2, IndexT work_random2, IndexT work_index2, IndexT const influences_row, IndexT const influences_id, unsigned int size)
	CUDA kernel propagating the state triple (status, weight, nodeID) to neighbors using a max()-operation. More...

template<typename IndexT >
__global__ void	amg_pmis2_mark_mis_nodes (IndexT const work_state, IndexT const work_index, IndexT point_types, IndexT undecided_buffer, unsigned int size)
	CUDA kernel for marking MIS and non-MIS nodes. More...

__global__ void	amg_pmis2_reset_state (unsigned int *point_types, unsigned int size)
	CUDA kernel for resetting non-MIS (i.e. coarse) points to undecided so that subsequent kernels work. More...

template<typename NumericT >
void	amg_coarse_ag_stage1_mis2 (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	AG (aggregation based) coarsening, single-threaded version of stage 1. More...

template<typename IndexT >
__global__ void	amg_agg_propagate_coarse_indices (IndexT point_types, IndexT coarse_ids, IndexT const influences_row, IndexT const influences_id, unsigned int size)

template<typename IndexT >
__global__ void	amg_agg_merge_undecided (IndexT point_types, IndexT coarse_ids, IndexT const influences_row, IndexT const influences_id, unsigned int size)

__global__ void	amg_agg_merge_undecided_2 (unsigned int *point_types, unsigned int size)

template<typename NumericT >
void	amg_coarse_ag (compressed_matrix< NumericT > const &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	AG (aggregation based) coarsening. Partially single-threaded version (VIENNACL_AMG_COARSE_AG) More...

template<typename InternalT1 >
void	amg_coarse (InternalT1 &A, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	Calls the right coarsening procedure. More...

template<typename NumericT >
__global__ void	amg_interpol_ag_kernel (unsigned int P_row_buffer, unsigned int P_col_buffer, NumericT P_elements, unsigned int coarse_ids, unsigned int size)

template<typename NumericT >
void	amg_interpol_ag (compressed_matrix< NumericT > const &A, compressed_matrix< NumericT > &P, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	AG (aggregation based) interpolation. Multi-Threaded! (VIENNACL_INTERPOL_SA) More...

template<typename NumericT >
__global__ void	amg_interpol_sa_kernel (const unsigned int A_row_indices, const unsigned int A_col_indices, const NumericT A_elements, unsigned int A_size1, unsigned int A_nnz, unsigned int Jacobi_row_indices, unsigned int Jacobi_col_indices, NumericT Jacobi_elements, NumericT omega)

template<typename NumericT >
void	amg_interpol_sa (compressed_matrix< NumericT > const &A, compressed_matrix< NumericT > &P, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	Smoothed aggregation interpolation. (VIENNACL_INTERPOL_SA) More...

template<typename MatrixT >
void	amg_interpol (MatrixT const &A, MatrixT &P, viennacl::linalg::detail::amg::amg_level_context &amg_context, viennacl::linalg::amg_tag &tag)
	Dispatcher for building the interpolation matrix. More...

template<typename NumericT >
__global__ void	compressed_matrix_assign_to_dense (const unsigned int row_indices, const unsigned int column_indices, const NumericT elements, NumericT B, unsigned int B_row_start, unsigned int B_col_start, unsigned int B_row_inc, unsigned int B_col_inc, unsigned int B_row_size, unsigned int B_col_size, unsigned int B_internal_rows, unsigned int B_internal_cols)

template<typename NumericT , unsigned int AlignmentV>
void	assign_to_dense (viennacl::compressed_matrix< NumericT, AlignmentV > const &A, viennacl::matrix_base< NumericT > &B)

template<typename NumericT >
__global__ void	compressed_matrix_smooth_jacobi_kernel (const unsigned int row_indices, const unsigned int column_indices, const NumericT elements, NumericT weight, const NumericT x_old, NumericT x_new, const NumericT rhs, unsigned int size)

template<typename NumericT >
void	smooth_jacobi (unsigned int iterations, compressed_matrix< NumericT > const &A, vector< NumericT > &x, vector< NumericT > &x_backup, vector< NumericT > const &rhs_smooth, NumericT weight)
	Damped Jacobi Smoother (CUDA version) More...

Function Documentation

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_agg_merge_undecided	(	IndexT *	point_types,
		IndexT *	coarse_ids,
		IndexT const *	influences_row,
		IndexT const *	influences_id,
		unsigned int	size
	)

Definition at line 416 of file amg_operations.hpp.

__global__ void viennacl::linalg::cuda::amg::amg_agg_merge_undecided_2	(	unsigned int *	point_types,
		unsigned int	size
	)

Definition at line 445 of file amg_operations.hpp.

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_agg_propagate_coarse_indices	(	IndexT *	point_types,
		IndexT *	coarse_ids,
		IndexT const *	influences_row,
		IndexT const *	influences_id,
		unsigned int	size
	)

Definition at line 386 of file amg_operations.hpp.

template<typename InternalT1 >

void viennacl::linalg::cuda::amg::amg_coarse	(	InternalT1 &	A,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

Calls the right coarsening procedure.

Parameters

A	Operator matrix on all levels
amg_context	AMG hierarchy datastructures
tag	AMG preconditioner tag

Definition at line 526 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_coarse_ag	(	compressed_matrix< NumericT > const &	A,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

AG (aggregation based) coarsening. Partially single-threaded version (VIENNACL_AMG_COARSE_AG)

Parameters

A	Operator matrix
amg_context	AMG hierarchy datastructures
tag	AMG preconditioner tag

Definition at line 466 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_coarse_ag_stage1_mis2	(	compressed_matrix< NumericT > const &	A,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

AG (aggregation based) coarsening, single-threaded version of stage 1.

Parameters

A	Operator matrix on all levels
amg_context	AMG hierarchy datastructures
tag	AMG preconditioner tag

Definition at line 290 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_influence	(	compressed_matrix< NumericT > const &	A,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

Dispatcher for influence processing.

Definition at line 103 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_influence_advanced	(	compressed_matrix< NumericT > const &	A,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

Routine for extracting strongly connected points considering a user-provided threshold value.

Definition at line 94 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_influence_trivial	(	compressed_matrix< NumericT > const &	A,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

Routine for taking all connections in the matrix as strong.

Definition at line 74 of file amg_operations.hpp.

__global__ void viennacl::linalg::cuda::amg::amg_influence_trivial_kernel	(	const unsigned int *	row_indices,
		const unsigned int *	column_indices,
		unsigned int	size1,
		unsigned int	nnz,
		unsigned int *	influences_row,
		unsigned int *	influences_id,
		unsigned int *	influences_values
	)

Definition at line 44 of file amg_operations.hpp.

template<typename MatrixT >

void viennacl::linalg::cuda::amg::amg_interpol	(	MatrixT const &	A,
		MatrixT &	P,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

Dispatcher for building the interpolation matrix.

Parameters

A	Operator matrix
P	Prolongation matrix
amg_context	AMG hierarchy datastructures
tag	AMG configuration tag

Definition at line 693 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_interpol_ag	(	compressed_matrix< NumericT > const &	A,
		compressed_matrix< NumericT > &	P,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

AG (aggregation based) interpolation. Multi-Threaded! (VIENNACL_INTERPOL_SA)

Parameters

A	Operator matrix
P	Prolongation matrix
amg_context	AMG hierarchy datastructures
tag	AMG configuration tag

Definition at line 572 of file amg_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::amg::amg_interpol_ag_kernel	(	unsigned int *	P_row_buffer,
		unsigned int *	P_col_buffer,
		NumericT *	P_elements,
		unsigned int *	coarse_ids,
		unsigned int	size
	)

Definition at line 543 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::amg_interpol_sa	(	compressed_matrix< NumericT > const &	A,
		compressed_matrix< NumericT > &	P,
		viennacl::linalg::detail::amg::amg_level_context &	amg_context,
		viennacl::linalg::amg_tag &	tag
	)

Smoothed aggregation interpolation. (VIENNACL_INTERPOL_SA)

Parameters

A	Operator matrix
P	Prolongation matrix
amg_context	AMG hierarchy datastructures
tag	AMG configuration tag

Definition at line 654 of file amg_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::amg::amg_interpol_sa_kernel	(	const unsigned int *	A_row_indices,
		const unsigned int *	A_col_indices,
		const NumericT *	A_elements,
		unsigned int	A_size1,
		unsigned int	A_nnz,
		unsigned int *	Jacobi_row_indices,
		unsigned int *	Jacobi_col_indices,
		NumericT *	Jacobi_elements,
		NumericT	omega
	)

Definition at line 594 of file amg_operations.hpp.

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_pmis2_init_workdata	(	IndexT *	work_state,
		IndexT *	work_random,
		IndexT *	work_index,
		IndexT const *	point_types,
		IndexT const *	random_weights,
		unsigned int	size
	)

CUDA kernel initializing the work vectors at each PMIS iteration.

Definition at line 139 of file amg_operations.hpp.

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_pmis2_mark_mis_nodes	(	IndexT const *	work_state,
		IndexT const *	work_index,
		IndexT *	point_types,
		IndexT *	undecided_buffer,
		unsigned int	size
	)

CUDA kernel for marking MIS and non-MIS nodes.

Definition at line 229 of file amg_operations.hpp.

template<typename IndexT >

__global__ void viennacl::linalg::cuda::amg::amg_pmis2_max_neighborhood	(	IndexT const *	work_state,
		IndexT const *	work_random,
		IndexT const *	work_index,
		IndexT *	work_state2,
		IndexT *	work_random2,
		IndexT *	work_index2,
		IndexT const *	influences_row,
		IndexT const *	influences_id,
		unsigned int	size
	)

CUDA kernel propagating the state triple (status, weight, nodeID) to neighbors using a max()-operation.

Definition at line 167 of file amg_operations.hpp.

__global__ void viennacl::linalg::cuda::amg::amg_pmis2_reset_state	(	unsigned int *	point_types,
		unsigned int	size
	)

CUDA kernel for resetting non-MIS (i.e. coarse) points to undecided so that subsequent kernels work.

Definition at line 271 of file amg_operations.hpp.

template<typename NumericT , unsigned int AlignmentV>

void viennacl::linalg::cuda::amg::assign_to_dense	(	viennacl::compressed_matrix< NumericT, AlignmentV > const &	A,
		viennacl::matrix_base< NumericT > &	B
	)

Definition at line 730 of file amg_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::amg::compressed_matrix_assign_to_dense	(	const unsigned int *	row_indices,
		const unsigned int *	column_indices,
		const NumericT *	elements,
		NumericT *	B,
		unsigned int	B_row_start,
		unsigned int	B_col_start,
		unsigned int	B_row_inc,
		unsigned int	B_col_inc,
		unsigned int	B_row_size,
		unsigned int	B_col_size,
		unsigned int	B_internal_rows,
		unsigned int	B_internal_cols
	)

Definition at line 708 of file amg_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::amg::compressed_matrix_smooth_jacobi_kernel	(	const unsigned int *	row_indices,
		const unsigned int *	column_indices,
		const NumericT *	elements,
		NumericT	weight,
		const NumericT *	x_old,
		NumericT *	x_new,
		const NumericT *	rhs,
		unsigned int	size
	)

Definition at line 749 of file amg_operations.hpp.

void viennacl::linalg::cuda::amg::enumerate_coarse_points ( viennacl::linalg::detail::amg::amg_level_context & amg_context )

inline

Assign IDs to coarse points.

TODO: Use exclusive_scan on GPU for this.

Definition at line 115 of file amg_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::amg::smooth_jacobi	(	unsigned int	iterations,
		compressed_matrix< NumericT > const &	A,
		vector< NumericT > &	x,
		vector< NumericT > &	x_backup,
		vector< NumericT > const &	rhs_smooth,
		NumericT	weight
	)

Damped Jacobi Smoother (CUDA version)

Parameters

iterations	Number of smoother iterations
A	Operator matrix for the smoothing
x	The vector smoothing is applied to
x_backup	(Different) Vector holding the same values as x
rhs_smooth	The right hand side of the equation for the smoother
weight	Damping factor. 0: No effect of smoother. 1: Undamped Jacobi iteration

Definition at line 791 of file amg_operations.hpp.

Functions

Function Documentation