ViennaCL - The Vienna Computing Library  1.7.0
Free open-source GPU-accelerated linear algebra and solver library.
Namespaces | Classes | Functions
viennacl::linalg::cuda::detail Namespace Reference

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

Namespaces

 fft
 

Classes

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 863 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 882 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 981 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 53 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 106 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 1108 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.