sparse_matrix_operations.hpp File Reference

Implementations of operations using sparse matrices using CUDA. More...

#include "viennacl/forwards.h"
#include "viennacl/scalar.hpp"
#include "viennacl/vector.hpp"
#include "viennacl/tools/tools.hpp"
#include "viennacl/linalg/cuda/common.hpp"
#include "viennacl/linalg/cuda/vector_operations.hpp"
#include "viennacl/linalg/cuda/sparse_matrix_operations_solve.hpp"
#include "viennacl/linalg/cuda/spgemm_rmerge.hpp"

Go to the source code of this file.

Classes
struct	viennacl::linalg::cuda::detail::spmv_pure
struct	viennacl::linalg::cuda::detail::spmv_alpha_beta
struct	viennacl::linalg::cuda::mat_mult_matrix_index< LayoutT >
	Helper struct for accessing an element of a row- or column-major matrix. More...

Namespaces
	viennacl
	Main namespace in ViennaCL. Holds all the basic types such as vector, matrix, etc. and defines operations upon them.
	viennacl::linalg
	Provides all linear algebra operations which are not covered by operator overloads.
	viennacl::linalg::cuda
	Holds all CUDA compute kernels used by ViennaCL.
	viennacl::linalg::cuda::detail
	Helper functions for the CUDA linear algebra backend.

Functions
template<typename NumericT >
__global__ void	viennacl::linalg::cuda::detail::csr_row_info_extractor_kernel (const unsigned int *row_indices, const unsigned int *column_indices, const NumericT *elements, NumericT *result, unsigned int size, unsigned int option)
template<typename NumericT , unsigned int AligmentV>
void	viennacl::linalg::cuda::detail::row_info (compressed_matrix< NumericT, AligmentV > const &mat, vector_base< NumericT > &vec, viennacl::linalg::detail::row_info_types info_selector)
template<unsigned int SubWarpSizeV, typename AlphaBetaHandlerT , typename NumericT >
__global__ void	viennacl::linalg::cuda::compressed_matrix_vec_mul_kernel (const unsigned int *row_indices, const unsigned int *column_indices, const NumericT *elements, const NumericT *x, unsigned int start_x, unsigned int inc_x, NumericT alpha, NumericT *result, unsigned int start_result, unsigned int inc_result, unsigned int size_result, NumericT beta)
template<typename AlphaBetaHandlerT , typename NumericT >
__global__ void	viennacl::linalg::cuda::compressed_matrix_vec_mul_adaptive_kernel (const unsigned int *row_indices, const unsigned int *column_indices, const unsigned int *row_blocks, const NumericT *elements, unsigned int num_blocks, const NumericT *x, unsigned int start_x, unsigned int inc_x, NumericT alpha, NumericT *result, unsigned int start_result, unsigned int inc_result, unsigned int size_result, NumericT beta)
template<class NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::compressed_matrix< NumericT, AlignmentV > &mat, const viennacl::vector_base< NumericT > &vec, NumericT alpha, viennacl::vector_base< NumericT > &result, NumericT beta)
	Carries out matrix-vector multiplication with a compressed_matrix. More...
template<typename DMatIndexT , typename ResultIndexT , typename NumericT >
__global__ void	viennacl::linalg::cuda::compressed_matrix_d_mat_mul_kernel (const unsigned int *sp_mat_row_indices, const unsigned int *sp_mat_col_indices, const NumericT *sp_mat_elements, const NumericT *d_mat, unsigned int d_mat_row_start, unsigned int d_mat_col_start, unsigned int d_mat_row_inc, unsigned int d_mat_col_inc, unsigned int d_mat_row_size, unsigned int d_mat_col_size, unsigned int d_mat_internal_rows, unsigned int d_mat_internal_cols, NumericT *result, unsigned int result_row_start, unsigned int result_col_start, unsigned int result_row_inc, unsigned int result_col_inc, unsigned int result_row_size, unsigned int result_col_size, unsigned int result_internal_rows, unsigned int result_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::compressed_matrix< NumericT, AlignmentV > &sp_mat, const viennacl::matrix_base< NumericT > &d_mat, viennacl::matrix_base< NumericT > &result)
	Carries out sparse_matrix-dense_matrix multiplication first matrix being compressed. More...
template<typename DMatIndexT , typename ResultIndexT , typename NumericT >
__global__ void	viennacl::linalg::cuda::compressed_matrix_d_tr_mat_mul_kernel (const unsigned int *sp_mat_row_indices, const unsigned int *sp_mat_col_indices, const NumericT *sp_mat_elements, const NumericT *d_mat, unsigned int d_mat_row_start, unsigned int d_mat_col_start, unsigned int d_mat_row_inc, unsigned int d_mat_col_inc, unsigned int d_mat_row_size, unsigned int d_mat_col_size, unsigned int d_mat_internal_rows, unsigned int d_mat_internal_cols, NumericT *result, unsigned int result_row_start, unsigned int result_col_start, unsigned int result_row_inc, unsigned int result_col_inc, unsigned int result_row_size, unsigned int result_col_size, unsigned int result_internal_rows, unsigned int result_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::compressed_matrix< NumericT, AlignmentV > &sp_mat, const viennacl::matrix_expression< const viennacl::matrix_base< NumericT >, const viennacl::matrix_base< NumericT >, viennacl::op_trans > &d_mat, viennacl::matrix_base< NumericT > &result)
	Carries out matrix-trans(matrix) multiplication first matrix being compressed and the second transposed. More...
template<typename NumericT >
__global__ void	viennacl::linalg::cuda::compressed_matrix_diagonal_kernel (const unsigned int *row_indices, const unsigned int *column_indices, const NumericT *elements, NumericT *result, unsigned int size)
template<typename SparseMatrixT , typename NumericT >
viennacl::enable_if < viennacl::is_any_sparse_matrix < SparseMatrixT >::value > ::type	viennacl::linalg::cuda::inplace_solve (const SparseMatrixT &mat, viennacl::vector_base< NumericT > &vec, viennacl::linalg::unit_lower_tag)
	Carries out triangular inplace solves. More...
template<typename SparseMatrixT , typename NumericT >
viennacl::enable_if < viennacl::is_any_sparse_matrix < SparseMatrixT >::value > ::type	viennacl::linalg::cuda::inplace_solve (const SparseMatrixT &mat, viennacl::vector_base< NumericT > &vec, viennacl::linalg::lower_tag)
	Carries out triangular inplace solves. More...
template<typename SparseMatrixT , typename NumericT >
viennacl::enable_if < viennacl::is_any_sparse_matrix < SparseMatrixT >::value > ::type	viennacl::linalg::cuda::inplace_solve (const SparseMatrixT &mat, viennacl::vector_base< NumericT > &vec, viennacl::linalg::unit_upper_tag)
	Carries out triangular inplace solves. More...
template<typename SparseMatrixT , typename NumericT >
viennacl::enable_if < viennacl::is_any_sparse_matrix < SparseMatrixT >::value > ::type	viennacl::linalg::cuda::inplace_solve (const SparseMatrixT &mat, viennacl::vector_base< NumericT > &vec, viennacl::linalg::upper_tag)
	Carries out triangular inplace solves. More...
template<typename SparseMatrixT , typename NumericT >
viennacl::enable_if < viennacl::is_any_sparse_matrix < SparseMatrixT >::value > ::type	viennacl::linalg::cuda::inplace_solve (const matrix_expression< const SparseMatrixT, const SparseMatrixT, op_trans > &mat, viennacl::vector_base< NumericT > &vec, viennacl::linalg::unit_lower_tag)
	Carries out triangular inplace solves. More...
template<typename SparseMatrixT , typename NumericT >
viennacl::enable_if < viennacl::is_any_sparse_matrix < SparseMatrixT >::value > ::type	viennacl::linalg::cuda::inplace_solve (const matrix_expression< const SparseMatrixT, const SparseMatrixT, op_trans > &mat, viennacl::vector_base< NumericT > &vec, viennacl::linalg::lower_tag)
	Carries out triangular inplace solves. More...
template<typename SparseMatrixT , typename NumericT >
viennacl::enable_if < viennacl::is_any_sparse_matrix < SparseMatrixT >::value > ::type	viennacl::linalg::cuda::inplace_solve (const matrix_expression< const SparseMatrixT, const SparseMatrixT, op_trans > &mat, viennacl::vector_base< NumericT > &vec, viennacl::linalg::unit_upper_tag)
	Carries out triangular inplace solves. More...
template<typename SparseMatrixT , typename NumericT >
viennacl::enable_if < viennacl::is_any_sparse_matrix < SparseMatrixT >::value > ::type	viennacl::linalg::cuda::inplace_solve (const matrix_expression< const SparseMatrixT, const SparseMatrixT, op_trans > &mat, viennacl::vector_base< NumericT > &vec, viennacl::linalg::upper_tag)
	Carries out triangular inplace solves. More...
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::detail::block_inplace_solve (const matrix_expression< const compressed_matrix< NumericT, AlignmentV >, const compressed_matrix< NumericT, AlignmentV >, op_trans > &L, viennacl::backend::mem_handle const &block_indices, vcl_size_t num_blocks, vector_base< NumericT > const &, vector_base< NumericT > &vec, viennacl::linalg::unit_lower_tag)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::detail::block_inplace_solve (const matrix_expression< const compressed_matrix< NumericT, AlignmentV >, const compressed_matrix< NumericT, AlignmentV >, op_trans > &U, viennacl::backend::mem_handle const &block_indices, vcl_size_t num_blocks, vector_base< NumericT > const &U_diagonal, vector_base< NumericT > &vec, viennacl::linalg::upper_tag)
template<typename NumericT >
__global__ void	viennacl::linalg::cuda::compressed_compressed_matrix_vec_mul_kernel (const unsigned int *row_jumper, const unsigned int *row_indices, const unsigned int *column_indices, const NumericT *elements, unsigned int nonzero_rows, const NumericT *x, unsigned int start_x, unsigned int inc_x, NumericT alpha, NumericT *result, unsigned int start_result, unsigned int inc_result, unsigned int size_result, NumericT beta)
template<typename NumericT >
void	viennacl::linalg::cuda::prod_impl (const viennacl::compressed_compressed_matrix< NumericT > &mat, const viennacl::vector_base< NumericT > &vec, NumericT alpha, viennacl::vector_base< NumericT > &result, NumericT beta)
	Carries out matrix-vector multiplication with a compressed_compressed_matrix. More...
template<typename NumericT >
__global__ void	viennacl::linalg::cuda::detail::coo_row_info_extractor (const unsigned int *coords, const NumericT *elements, const unsigned int *group_boundaries, NumericT *result, unsigned int option)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::detail::row_info (coordinate_matrix< NumericT, AlignmentV > const &mat, vector_base< NumericT > &vec, viennacl::linalg::detail::row_info_types info_selector)
template<typename NumericT >
__global__ void	viennacl::linalg::cuda::coordinate_matrix_vec_mul_kernel (const unsigned int *coords, const NumericT *elements, const unsigned int *group_boundaries, const NumericT *x, unsigned int start_x, unsigned int inc_x, NumericT alpha, NumericT *result, unsigned int start_result, unsigned int inc_result, NumericT beta)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::coordinate_matrix< NumericT, AlignmentV > &mat, const viennacl::vector_base< NumericT > &vec, NumericT alpha, viennacl::vector_base< NumericT > &result, NumericT beta)
	Carries out matrix-vector multiplication with a coordinate_matrix. More...
template<typename DMatIndexT , typename ResultIndexT , typename NumericT >
__global__ void	viennacl::linalg::cuda::coordinate_matrix_d_mat_mul_kernel (const unsigned int *coords, const NumericT *elements, const unsigned int *group_boundaries, const NumericT *d_mat, unsigned int d_mat_row_start, unsigned int d_mat_col_start, unsigned int d_mat_row_inc, unsigned int d_mat_col_inc, unsigned int d_mat_row_size, unsigned int d_mat_col_size, unsigned int d_mat_internal_rows, unsigned int d_mat_internal_cols, NumericT *result, unsigned int result_row_start, unsigned int result_col_start, unsigned int result_row_inc, unsigned int result_col_inc, unsigned int result_row_size, unsigned int result_col_size, unsigned int result_internal_rows, unsigned int result_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::coordinate_matrix< NumericT, AlignmentV > &sp_mat, const viennacl::matrix_base< NumericT > &d_mat, viennacl::matrix_base< NumericT > &result)
	Carries out Compressed Matrix(COO)-Dense Matrix multiplication. More...
template<typename DMatIndexT , typename ResultIndexT , typename NumericT >
__global__ void	viennacl::linalg::cuda::coordinate_matrix_d_tr_mat_mul_kernel (const unsigned int *coords, const NumericT *elements, const unsigned int *group_boundaries, const NumericT *d_mat, unsigned int d_mat_row_start, unsigned int d_mat_col_start, unsigned int d_mat_row_inc, unsigned int d_mat_col_inc, unsigned int d_mat_row_size, unsigned int d_mat_col_size, unsigned int d_mat_internal_rows, unsigned int d_mat_internal_cols, NumericT *result, unsigned int result_row_start, unsigned int result_col_start, unsigned int result_row_inc, unsigned int result_col_inc, unsigned int result_row_size, unsigned int result_col_size, unsigned int result_internal_rows, unsigned int result_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::coordinate_matrix< NumericT, AlignmentV > &sp_mat, const viennacl::matrix_expression< const viennacl::matrix_base< NumericT >, const viennacl::matrix_base< NumericT >, viennacl::op_trans > &d_mat, viennacl::matrix_base< NumericT > &result)
	Carries out Compressed Matrix(COO)-Dense Transposed Matrix multiplication. More...
template<typename AlphaBetaHandlerT , typename NumericT >
__global__ void	viennacl::linalg::cuda::ell_matrix_vec_mul_kernel (const unsigned int *coords, const NumericT *elements, const NumericT *x, unsigned int start_x, unsigned int inc_x, NumericT alpha, NumericT *result, unsigned int start_result, unsigned int inc_result, NumericT beta, unsigned int row_num, unsigned int col_num, unsigned int internal_row_num, unsigned int items_per_row, unsigned int aligned_items_per_row)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::ell_matrix< NumericT, AlignmentV > &mat, const viennacl::vector_base< NumericT > &vec, NumericT alpha, viennacl::vector_base< NumericT > &result, NumericT beta)
	Carries out matrix-vector multiplication with a ell_matrix. More...
template<typename DMatIndexT , typename ResultIndexT , typename NumericT >
__global__ void	viennacl::linalg::cuda::ell_matrix_d_mat_mul_kernel (const unsigned int *sp_mat_coords, const NumericT *sp_mat_elements, unsigned int sp_mat_row_num, unsigned int sp_mat_col_num, unsigned int sp_mat_internal_row_num, unsigned int sp_mat_items_per_row, unsigned int sp_mat_aligned_items_per_row, const NumericT *d_mat, unsigned int d_mat_row_start, unsigned int d_mat_col_start, unsigned int d_mat_row_inc, unsigned int d_mat_col_inc, unsigned int d_mat_row_size, unsigned int d_mat_col_size, unsigned int d_mat_internal_rows, unsigned int d_mat_internal_cols, NumericT *result, unsigned int result_row_start, unsigned int result_col_start, unsigned int result_row_inc, unsigned int result_col_inc, unsigned int result_row_size, unsigned int result_col_size, unsigned int result_internal_rows, unsigned int result_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::ell_matrix< NumericT, AlignmentV > &sp_mat, const viennacl::matrix_base< NumericT > &d_mat, viennacl::matrix_base< NumericT > &result)
	Carries out Sparse Matrix(ELL)-Dense Matrix multiplication. More...
template<typename DMatIndexT , typename ResultIndexT , typename NumericT >
__global__ void	viennacl::linalg::cuda::ell_matrix_d_tr_mat_mul_kernel (const unsigned int *sp_mat_coords, const NumericT *sp_mat_elements, unsigned int sp_mat_row_num, unsigned int sp_mat_col_num, unsigned int sp_mat_internal_row_num, unsigned int sp_mat_items_per_row, unsigned int sp_mat_aligned_items_per_row, const NumericT *d_mat, unsigned int d_mat_row_start, unsigned int d_mat_col_start, unsigned int d_mat_row_inc, unsigned int d_mat_col_inc, unsigned int d_mat_row_size, unsigned int d_mat_col_size, unsigned int d_mat_internal_rows, unsigned int d_mat_internal_cols, NumericT *result, unsigned int result_row_start, unsigned int result_col_start, unsigned int result_row_inc, unsigned int result_col_inc, unsigned int result_row_size, unsigned int result_col_size, unsigned int result_internal_rows, unsigned int result_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::ell_matrix< NumericT, AlignmentV > &sp_mat, const viennacl::matrix_expression< const viennacl::matrix_base< NumericT >, const viennacl::matrix_base< NumericT >, viennacl::op_trans > &d_mat, viennacl::matrix_base< NumericT > &result)
	Carries out Sparse Matrix(ELL)-Dense Transposed Matrix multiplication. More...
template<typename AlphaBetaHandlerT , typename NumericT >
__global__ void	viennacl::linalg::cuda::sliced_ell_matrix_vec_mul_kernel (const unsigned int *columns_per_block, const unsigned int *column_indices, const unsigned int *block_start, const NumericT *elements, const NumericT *x, unsigned int start_x, unsigned int inc_x, unsigned int size_x, NumericT alpha, NumericT *result, unsigned int start_result, unsigned int inc_result, unsigned int size_result, NumericT beta, unsigned int block_size)
template<typename NumericT , typename IndexT >
void	viennacl::linalg::cuda::prod_impl (const viennacl::sliced_ell_matrix< NumericT, IndexT > &mat, const viennacl::vector_base< NumericT > &vec, NumericT alpha, viennacl::vector_base< NumericT > &result, NumericT beta)
	Carries out matrix-vector multiplication with a sliced_ell_matrix. More...
template<typename AlphaBetaHandlerT , typename NumericT >
__global__ void	viennacl::linalg::cuda::hyb_matrix_vec_mul_kernel (const unsigned int *ell_coords, const NumericT *ell_elements, const unsigned int *csr_rows, const unsigned int *csr_cols, const NumericT *csr_elements, const NumericT *x, unsigned int start_x, unsigned int inc_x, NumericT alpha, NumericT *result, unsigned int start_result, unsigned int inc_result, NumericT beta, unsigned int row_num, unsigned int internal_row_num, unsigned int items_per_row, unsigned int aligned_items_per_row)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::hyb_matrix< NumericT, AlignmentV > &mat, const viennacl::vector_base< NumericT > &vec, NumericT alpha, viennacl::vector_base< NumericT > &result, NumericT beta)
	Carries out matrix-vector multiplication with a hyb_matrix. More...
template<typename DMatIndexT , typename ResultIndexT , typename NumericT >
__global__ void	viennacl::linalg::cuda::hyb_matrix_d_mat_mul_kernel (const unsigned int *ell_coords, const NumericT *ell_elements, const unsigned int *csr_rows, const unsigned int *csr_cols, const NumericT *csr_elements, unsigned int row_num, unsigned int internal_row_num, unsigned int items_per_row, unsigned int aligned_items_per_row, const NumericT *d_mat, unsigned int d_mat_row_start, unsigned int d_mat_col_start, unsigned int d_mat_row_inc, unsigned int d_mat_col_inc, unsigned int d_mat_row_size, unsigned int d_mat_col_size, unsigned int d_mat_internal_rows, unsigned int d_mat_internal_cols, NumericT *result, unsigned int result_row_start, unsigned int result_col_start, unsigned int result_row_inc, unsigned int result_col_inc, unsigned int result_row_size, unsigned int result_col_size, unsigned int result_internal_rows, unsigned int result_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::hyb_matrix< NumericT, AlignmentV > &mat, const viennacl::matrix_base< NumericT > &d_mat, viennacl::matrix_base< NumericT > &result)
	Carries out matrix-vector multiplication with a hyb_matrix. More...
template<typename DMatIndexT , typename ResultIndexT , typename NumericT >
__global__ void	viennacl::linalg::cuda::hyb_matrix_d_tr_mat_mul_kernel (const unsigned int *ell_coords, const NumericT *ell_elements, const unsigned int *csr_rows, const unsigned int *csr_cols, const NumericT *csr_elements, unsigned int row_num, unsigned int internal_row_num, unsigned int items_per_row, unsigned int aligned_items_per_row, const NumericT *d_mat, unsigned int d_mat_row_start, unsigned int d_mat_col_start, unsigned int d_mat_row_inc, unsigned int d_mat_col_inc, unsigned int d_mat_row_size, unsigned int d_mat_col_size, unsigned int d_mat_internal_rows, unsigned int d_mat_internal_cols, NumericT *result, unsigned int result_row_start, unsigned int result_col_start, unsigned int result_row_inc, unsigned int result_col_inc, unsigned int result_row_size, unsigned int result_col_size, unsigned int result_internal_rows, unsigned int result_internal_cols)
template<typename NumericT , unsigned int AlignmentV>
void	viennacl::linalg::cuda::prod_impl (const viennacl::hyb_matrix< NumericT, AlignmentV > &mat, const viennacl::matrix_expression< const viennacl::matrix_base< NumericT >, const viennacl::matrix_base< NumericT >, viennacl::op_trans > &d_mat, viennacl::matrix_base< NumericT > &result)
	Carries out matrix-vector multiplication with a hyb_matrix. More...

Detailed Description

Implementations of operations using sparse matrices using CUDA.

Definition in file sparse_matrix_operations.hpp.