All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
viennacl::linalg::cuda::detail Namespace Reference

Helper functions for the CUDA linear algebra backend. More...

## Namespaces

| fft

## Classes

struct | spmv_alpha_beta
struct | spmv_pure
struct | type_to_type2
struct | type_to_type2< double >
struct | type_to_type2< float >

## Functions

unsigned int | make_options (vcl_size_t length, bool reciprocal, bool flip_sign)
void | cuda_last_error_check (const char *message, const char *file, const int line)
template<typename NumericT , typename OtherT >
viennacl::backend::mem_handle::cuda_handle_type & | arg_reference (viennacl::scalar< NumericT > &s, OtherT)
template<typename NumericT , typename OtherT >
viennacl::backend::mem_handle::cuda_handle_type const & | arg_reference (viennacl::scalar< NumericT > const &s, OtherT)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, char const & > ::type | arg_reference (ArgT, char const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, unsigned char const & >::type | arg_reference (ArgT, unsigned char const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, short const & > ::type | arg_reference (ArgT, short const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, unsigned short const & >::type | arg_reference (ArgT, unsigned short const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, int const & > ::type | arg_reference (ArgT, int const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, unsigned int const & >::type | arg_reference (ArgT, unsigned int const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, long const & > ::type | arg_reference (ArgT, long const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, unsigned long const & >::type | arg_reference (ArgT, unsigned long const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, float const & > ::type | arg_reference (ArgT, float const &val)
template<typename ArgT >
viennacl::enable_if < viennacl::is_cpu_scalar < ArgT >::value, double const & > ::type | arg_reference (ArgT, double const &val)
template<typename TagT >
bool | is_unit_solve (TagT const &tag)
bool | is_unit_solve (viennacl::linalg::unit_lower_tag)
bool | is_unit_solve (viennacl::linalg::unit_upper_tag)
template<typename TagT >
bool | is_upper_solve (TagT const &tag)
bool | is_upper_solve (viennacl::linalg::upper_tag)
bool | is_upper_solve (viennacl::linalg::unit_upper_tag)
template<typename Matrix1T , typename Matrix2T , typename SolverTagT >
void | inplace_solve_impl (Matrix1T const &A, Matrix2T &B, SolverTagT const &tag)
unsigned int | get_option_for_solver_tag (viennacl::linalg::upper_tag)
unsigned int | get_option_for_solver_tag (viennacl::linalg::unit_upper_tag)
unsigned int | get_option_for_solver_tag (viennacl::linalg::lower_tag)
unsigned int | get_option_for_solver_tag (viennacl::linalg::unit_lower_tag)
template<typename MatrixT , typename VectorT >
void | inplace_solve_vector_impl (MatrixT const &mat, VectorT &vec, unsigned int options)
template<typename MatrixT1 , typename MatrixT2 , typename MatrixT3 , typename ScalarT >
void | prod_slow_kernel (const MatrixT1 &A, bool transposed_A, const MatrixT2 &B, bool transposed_B, MatrixT3 &C, ScalarT alpha, ScalarT beta)
template<typename MatrixT1 , typename MatrixT2 , typename MatrixT3 , typename ScalarT >
void | prod (const MatrixT1 &A, bool transposed_A, const MatrixT2 &B, bool transposed_B, MatrixT3 &C, ScalarT alpha, ScalarT beta)
template<typename NumericT >
__global__ void | level_scheduling_substitute_kernel (const unsigned int *row_index_array, const unsigned int *row_indices, const unsigned int *column_indices, const NumericT *elements, NumericT *vec, unsigned int size)
template<typename NumericT >
void | level_scheduling_substitute (vector< NumericT > &vec, viennacl::backend::mem_handle const &row_index_array, viennacl::backend::mem_handle const &row_buffer, viennacl::backend::mem_handle const &col_buffer, viennacl::backend::mem_handle const &element_buffer, vcl_size_t num_rows)
template<typename NumericT >
__global__ void | 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 | row_info (compressed_matrix< NumericT, AligmentV > const &mat, vector_base< NumericT > &vec, viennacl::linalg::detail::row_info_types info_selector)
template<typename NumericT , unsigned int AlignmentV>
void | 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 | 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 | 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 | row_info (coordinate_matrix< NumericT, AlignmentV > const &mat, vector_base< NumericT > &vec, viennacl::linalg::detail::row_info_types info_selector)
template<typename NumericT >
void | scan_impl (vector_base< NumericT > const &input, vector_base< NumericT > &output, bool is_inclusive)
| Worker routine for scan routines. More...

## Detailed Description

Helper functions for the CUDA linear algebra backend.

## Function Documentation

template<typename NumericT , typename OtherT >

viennacl::backend::mem_handle::cuda_handle_type& viennacl::linalg::cuda::detail::arg_reference | ( | viennacl::scalar< NumericT > & | s,
|  | OtherT |
| ) |  |

Definition at line 188 of file common.hpp.

template<typename NumericT , typename OtherT >

viennacl::backend::mem_handle::cuda_handle_type const& viennacl::linalg::cuda::detail::arg_reference | ( | viennacl::scalar< NumericT > const & | s,
|  | OtherT |
| ) |  |

Definition at line 191 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, char const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | char const & | val
| ) |  |

Definition at line 197 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, unsigned char const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | unsigned char const & | val
| ) |  |

Definition at line 202 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, short const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | short const & | val
| ) |  |

Definition at line 207 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, unsigned short const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | unsigned short const & | val
| ) |  |

Definition at line 212 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, int const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | int const & | val
| ) |  |

Definition at line 217 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, unsigned int const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | unsigned int const & | val
| ) |  |

Definition at line 222 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, long const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | long const & | val
| ) |  |

Definition at line 227 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, unsigned long const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | unsigned long const & | val
| ) |  |

Definition at line 232 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, float const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | float const & | val
| ) |  |

Definition at line 237 of file common.hpp.

template<typename ArgT >

viennacl::enable_if< viennacl::is_cpu_scalar<ArgT>::value, double const &>::type viennacl::linalg::cuda::detail::arg_reference | ( | ArgT | ,
|  | double const & | val
| ) |  |

Definition at line 242 of file common.hpp.

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 |
| ) |  |

Definition at line 970 of file sparse_matrix_operations.hpp.

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 |
| ) |  |

Definition at line 989 of file sparse_matrix_operations.hpp.

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
| ) |  |

Definition at line 1095 of file sparse_matrix_operations.hpp.

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
| ) |  |

Definition at line 54 of file sparse_matrix_operations.hpp.

void viennacl::linalg::cuda::detail::cuda_last_error_check | ( | const char * | message,
|  | const char * | file,
|  | const int | line
| ) |  |

inline

Definition at line 165 of file common.hpp.

unsigned int viennacl::linalg::cuda::detail::get_option_for_solver_tag | ( | viennacl::linalg::upper_tag |  | ) |

inline

Definition at line 348 of file direct_solve.hpp.

unsigned int viennacl::linalg::cuda::detail::get_option_for_solver_tag | ( | viennacl::linalg::unit_upper_tag |  | ) |

inline

Definition at line 349 of file direct_solve.hpp.

unsigned int viennacl::linalg::cuda::detail::get_option_for_solver_tag | ( | viennacl::linalg::lower_tag |  | ) |

inline

Definition at line 350 of file direct_solve.hpp.

unsigned int viennacl::linalg::cuda::detail::get_option_for_solver_tag | ( | viennacl::linalg::unit_lower_tag |  | ) |

inline

Definition at line 351 of file direct_solve.hpp.

template<typename Matrix1T , typename Matrix2T , typename SolverTagT >

void viennacl::linalg::cuda::detail::inplace_solve_impl | ( | Matrix1T const & | A,
|  | Matrix2T & | B,
|  | SolverTagT const & | tag
| ) |  |

Definition at line 189 of file direct_solve.hpp.

template<typename MatrixT , typename VectorT >

void viennacl::linalg::cuda::detail::inplace_solve_vector_impl | ( | MatrixT const & | mat,
|  | VectorT & | vec,
|  | unsigned int | options
| ) |  |

Definition at line 354 of file direct_solve.hpp.

template<typename TagT >

bool viennacl::linalg::cuda::detail::is_unit_solve | ( | TagT const & | tag | ) |

Definition at line 177 of file direct_solve.hpp.

bool viennacl::linalg::cuda::detail::is_unit_solve | ( | viennacl::linalg::unit_lower_tag |  | ) |

inline

Definition at line 179 of file direct_solve.hpp.

bool viennacl::linalg::cuda::detail::is_unit_solve | ( | viennacl::linalg::unit_upper_tag |  | ) |

inline

Definition at line 180 of file direct_solve.hpp.

template<typename TagT >

bool viennacl::linalg::cuda::detail::is_upper_solve | ( | TagT const & | tag | ) |

Definition at line 183 of file direct_solve.hpp.

bool viennacl::linalg::cuda::detail::is_upper_solve | ( | viennacl::linalg::upper_tag |  | ) |

inline

Definition at line 185 of file direct_solve.hpp.

bool viennacl::linalg::cuda::detail::is_upper_solve | ( | viennacl::linalg::unit_upper_tag |  | ) |

inline

Definition at line 186 of file direct_solve.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::detail::level_scheduling_substitute | ( | vector< NumericT > & | vec,
|  | viennacl::backend::mem_handle const & | row_index_array,
|  | viennacl::backend::mem_handle const & | row_buffer,
|  | viennacl::backend::mem_handle const & | col_buffer,
|  | viennacl::backend::mem_handle const & | element_buffer,
|  | vcl_size_t | num_rows
| ) |  |

Definition at line 68 of file misc_operations.hpp.

template<typename NumericT >

__global__ void viennacl::linalg::cuda::detail::level_scheduling_substitute_kernel | ( | const unsigned int * | row_index_array,
|  | const unsigned int * | row_indices,
|  | const unsigned int * | column_indices,
|  | const NumericT * | elements,
|  | NumericT * | vec,
|  | unsigned int | size
| ) |  |

Definition at line 42 of file misc_operations.hpp.

unsigned int viennacl::linalg::cuda::detail::make_options | ( | vcl_size_t | length,
|  | bool | reciprocal,
|  | bool | flip_sign
| ) |  |

inline

Definition at line 160 of file common.hpp.

template<typename MatrixT1 , typename MatrixT2 , typename MatrixT3 , typename ScalarT >

void viennacl::linalg::cuda::detail::prod | ( | const MatrixT1 & | A,
|  | bool | transposed_A,
|  | const MatrixT2 & | B,
|  | bool | transposed_B,
|  | MatrixT3 & | C,
|  | ScalarT | alpha,
|  | ScalarT | beta
| ) |  |

Definition at line 2346 of file matrix_operations.hpp.

template<typename MatrixT1 , typename MatrixT2 , typename MatrixT3 , typename ScalarT >

void viennacl::linalg::cuda::detail::prod_slow_kernel | ( | const MatrixT1 & | A,
|  | bool | transposed_A,
|  | const MatrixT2 & | B,
|  | bool | transposed_B,
|  | MatrixT3 & | C,
|  | ScalarT | alpha,
|  | ScalarT | beta
| ) |  |

Definition at line 1569 of file matrix_operations.hpp.

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
| ) |  |

Definition at line 107 of file sparse_matrix_operations.hpp.

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
| ) |  |

Definition at line 1222 of file sparse_matrix_operations.hpp.

template<typename NumericT >

void viennacl::linalg::cuda::detail::scan_impl | ( | vector_base< NumericT > const & | input,
|  | vector_base< NumericT > & | output,
|  | bool | is_inclusive
| ) |  |

Worker routine for scan routines.

Note on performance: For non-in-place scans one could optimize away the temporary 'cuda_carries'-array. This, however, only provides small savings in the latency-dominated regime, yet would effectively double the amount of code to maintain.

Definition at line 3181 of file vector_operations.hpp.