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/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::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 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 *result, unsigned int start_result, unsigned int inc_result, unsigned int size_result) |
template<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 *result, unsigned int start_result, unsigned int inc_result, unsigned int size_result) |
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, viennacl::vector_base< NumericT > &result) |
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 *result, unsigned int start_result, unsigned int inc_result, unsigned int size_result) |
template<typename NumericT > | |
void | viennacl::linalg::cuda::prod_impl (const viennacl::compressed_compressed_matrix< NumericT > &mat, const viennacl::vector_base< NumericT > &vec, viennacl::vector_base< NumericT > &result) |
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 *result, unsigned int start_result, unsigned int inc_result) |
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, viennacl::vector_base< NumericT > &result) |
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 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 *result, unsigned int start_result, unsigned int inc_result, 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, viennacl::vector_base< NumericT > &result) |
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 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 *result, unsigned int start_result, unsigned int inc_result, unsigned int size_result, 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, viennacl::vector_base< NumericT > &result) |
Carries out matrix-vector multiplication with a sliced_ell_matrix. More... | |
template<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 *result, unsigned int start_result, unsigned int inc_result, 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, viennacl::vector_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_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... | |
Implementations of operations using sparse matrices using CUDA.
Definition in file sparse_matrix_operations.hpp.