All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
utils.hpp
Go to the documentation of this file.
#ifndef VIENNACL_DEVICE_SPECIFIC_UTILS_HPP
#define VIENNACL_DEVICE_SPECIFIC_UTILS_HPP
/* =========================================================================
Copyright (c) 2010-2016, Institute for Microelectronics,
Institute for Analysis and Scientific Computing,
TU Wien.
Portions of this software are copyright by UChicago Argonne, LLC.
-----------------
ViennaCL - The Vienna Computing Library
-----------------
Project Head: Karl Rupp rupp@iue.tuwien.ac.at
(A list of authors and contributors can be found in the manual)
License: MIT (X11), see file LICENSE in the base directory
============================================================================= */
#include <sstream>
#include "viennacl/detail/matrix_def.hpp"
#include "viennacl/detail/vector_def.hpp"
#include "viennacl/device_specific/forwards.h"
#include "viennacl/ocl/forwards.h"
#include "viennacl/scheduler/forwards.h"
#include "viennacl/traits/size.hpp"
#include "viennacl/traits/handle.hpp"
#include "viennacl/traits/row_major.hpp"
#include "viennacl/tools/tools.hpp"
namespace viennacl
{
namespace device_specific
{
namespace utils
{
//CUDA Conversion
inline std::string opencl_source_to_cuda_source(std::string const & opencl_src)
{
std::string res = opencl_src;
viennacl::tools::find_and_replace(res,"__attribute__","//__attribute__");
//Pointer
viennacl::tools::find_and_replace(res, "__global float*", "float*");
viennacl::tools::find_and_replace(res, "__local float*", "float*");
viennacl::tools::find_and_replace(res, "__global double*", "double*");
viennacl::tools::find_and_replace(res, "__local double*", "double*");
//Qualifiers
viennacl::tools::find_and_replace(res,"__global","__device__");
viennacl::tools::find_and_replace(res,"__kernel","__global__");
viennacl::tools::find_and_replace(res,"__constant","__constant__");
viennacl::tools::find_and_replace(res,"__local","__shared__");
//Indexing
viennacl::tools::find_and_replace(res,"get_num_groups(0)","gridDim.x");
viennacl::tools::find_and_replace(res,"get_num_groups(1)","gridDim.y");
viennacl::tools::find_and_replace(res,"get_local_size(0)","blockDim.x");
viennacl::tools::find_and_replace(res,"get_local_size(1)","blockDim.y");
viennacl::tools::find_and_replace(res,"get_group_id(0)","blockIdx.x");
viennacl::tools::find_and_replace(res,"get_group_id(1)","blockIdx.y");
viennacl::tools::find_and_replace(res,"get_local_id(0)","threadIdx.x");
viennacl::tools::find_and_replace(res,"get_local_id(1)","threadIdx.y");
viennacl::tools::find_and_replace(res,"get_global_id(0)","(blockIdx.x*blockDim.x + threadIdx.x)");
viennacl::tools::find_and_replace(res,"get_global_id(1)","(blockIdx.y*blockDim.y + threadIdx.y)");
//Synchronization
viennacl::tools::find_and_replace(res,"barrier(CLK_LOCAL_MEM_FENCE)","__syncthreads()");
viennacl::tools::find_and_replace(res,"barrier(CLK_GLOBAL_MEM_FENCE)","__syncthreads()");
return res;
}
static std::string numeric_type_to_string(scheduler::statement_node_numeric_type const & type){
switch (type)
{
//case scheduler::CHAR_TYPE: return "char";
//case scheduler::UCHAR_TYPE: return "unsigned char";
//case scheduler::SHORT_TYPE: return "short";
//case scheduler::USHORT_TYPE: return "unsigned short";
case scheduler::INT_TYPE: return "int";
case scheduler::UINT_TYPE: return "unsigned int";
case scheduler::LONG_TYPE: return "long";
case scheduler::ULONG_TYPE: return "unsigned long";
case scheduler::FLOAT_TYPE : return "float";
case scheduler::DOUBLE_TYPE : return "double";
default : throw generator_not_supported_exception("Unsupported Scalartype");
}
}
template<class Fun>
static typename Fun::result_type call_on_host_scalar(scheduler::lhs_rhs_element element, Fun const & fun){
assert(element.type_family == scheduler::SCALAR_TYPE_FAMILY && bool("Must be called on a host scalar"));
switch (element.numeric_type)
{
//case scheduler::CHAR_TYPE: return fun(element.host_char);
//case scheduler::UCHAR_TYPE: return fun(element.host_uchar);
//case scheduler::SHORT_TYPE: return fun(element.host_short);
//case scheduler::USHORT_TYPE: return fun(element.host_ushort);
case scheduler::INT_TYPE: return fun(element.host_int);
case scheduler::UINT_TYPE: return fun(element.host_uint);
case scheduler::LONG_TYPE: return fun(element.host_long);
case scheduler::ULONG_TYPE: return fun(element.host_ulong);
case scheduler::FLOAT_TYPE : return fun(element.host_float);
case scheduler::DOUBLE_TYPE : return fun(element.host_double);
default : throw generator_not_supported_exception("Unsupported Scalartype");
}
}
template<class Fun>
static typename Fun::result_type call_on_scalar(scheduler::lhs_rhs_element element, Fun const & fun){
assert(element.type_family == scheduler::SCALAR_TYPE_FAMILY && bool("Must be called on a scalar"));
switch (element.numeric_type)
{
//case scheduler::CHAR_TYPE: return fun(*element.scalar_char);
//case scheduler::UCHAR_TYPE: return fun(*element.scalar_uchar);
//case scheduler::SHORT_TYPE: return fun(*element.scalar_short);
//case scheduler::USHORT_TYPE: return fun(*element.scalar_ushort);
case scheduler::INT_TYPE: return fun(*element.scalar_int);
case scheduler::UINT_TYPE: return fun(*element.scalar_uint);
case scheduler::LONG_TYPE: return fun(*element.scalar_long);
case scheduler::ULONG_TYPE: return fun(*element.scalar_ulong);
case scheduler::FLOAT_TYPE : return fun(*element.scalar_float);
case scheduler::DOUBLE_TYPE : return fun(*element.scalar_double);
default : throw generator_not_supported_exception("Unsupported Scalartype");
}
}
template<class Fun>
static typename Fun::result_type call_on_vector(scheduler::lhs_rhs_element element, Fun const & fun){
assert(element.type_family == scheduler::VECTOR_TYPE_FAMILY && bool("Must be called on a vector"));
switch (element.numeric_type)
{
//case scheduler::CHAR_TYPE: return fun(*element.vector_char);
//case scheduler::UCHAR_TYPE: return fun(*element.vector_uchar);
//case scheduler::SHORT_TYPE: return fun(*element.vector_short);
//case scheduler::USHORT_TYPE: return fun(*element.vector_ushort);
case scheduler::INT_TYPE: return fun(*element.vector_int);
case scheduler::UINT_TYPE: return fun(*element.vector_uint);
case scheduler::LONG_TYPE: return fun(*element.vector_long);
case scheduler::ULONG_TYPE: return fun(*element.vector_ulong);
case scheduler::FLOAT_TYPE : return fun(*element.vector_float);
case scheduler::DOUBLE_TYPE : return fun(*element.vector_double);
default : throw generator_not_supported_exception("Unsupported Scalartype");
}
}
template<class Fun>
static typename Fun::result_type call_on_implicit_vector(scheduler::lhs_rhs_element element, Fun const & fun){
assert(element.type_family == scheduler::VECTOR_TYPE_FAMILY && bool("Must be called on a implicit_vector"));
assert(element.subtype == scheduler::IMPLICIT_VECTOR_TYPE && bool("Must be called on a implicit_vector"));
switch (element.numeric_type)
{
//case scheduler::CHAR_TYPE: return fun(*element.implicit_vector_char);
//case scheduler::UCHAR_TYPE: return fun(*element.implicit_vector_uchar);
//case scheduler::SHORT_TYPE: return fun(*element.implicit_vector_short);
//case scheduler::USHORT_TYPE: return fun(*element.implicit_vector_ushort);
case scheduler::INT_TYPE: return fun(*element.implicit_vector_int);
case scheduler::UINT_TYPE: return fun(*element.implicit_vector_uint);
case scheduler::LONG_TYPE: return fun(*element.implicit_vector_long);
case scheduler::ULONG_TYPE: return fun(*element.implicit_vector_ulong);
case scheduler::FLOAT_TYPE : return fun(*element.implicit_vector_float);
case scheduler::DOUBLE_TYPE : return fun(*element.implicit_vector_double);
default : throw generator_not_supported_exception("Unsupported Scalartype");
}
}
template<class Fun>
static typename Fun::result_type call_on_matrix(scheduler::lhs_rhs_element element, Fun const & fun){
assert(element.type_family == scheduler::MATRIX_TYPE_FAMILY && bool("Must be called on a matrix"));
switch (element.numeric_type)
{
//case scheduler::CHAR_TYPE: return fun(*element.matrix_char);
//case scheduler::UCHAR_TYPE: return fun(*element.matrix_uchar);
//case scheduler::SHORT_TYPE: return fun(*element.matrix_short);
//case scheduler::USHORT_TYPE: return fun(*element.matrix_ushort);
case scheduler::INT_TYPE: return fun(*element.matrix_int);
case scheduler::UINT_TYPE: return fun(*element.matrix_uint);
case scheduler::LONG_TYPE: return fun(*element.matrix_long);
case scheduler::ULONG_TYPE: return fun(*element.matrix_ulong);
case scheduler::FLOAT_TYPE : return fun(*element.matrix_float);
case scheduler::DOUBLE_TYPE : return fun(*element.matrix_double);
default : throw generator_not_supported_exception("Unsupported Scalartype");
}
}
template<class Fun>
static typename Fun::result_type call_on_implicit_matrix(scheduler::lhs_rhs_element element, Fun const & fun){
assert(element.subtype == scheduler::IMPLICIT_MATRIX_TYPE && bool("Must be called on a implicit matrix"));
switch (element.numeric_type)
{
//case scheduler::CHAR_TYPE: return fun(*element.implicit_matrix_char);
//case scheduler::UCHAR_TYPE: return fun(*element.implicit_matrix_uchar);
//case scheduler::SHORT_TYPE: return fun(*element.implicit_matrix_short);
//case scheduler::USHORT_TYPE: return fun(*element.implicit_matrix_ushort);
case scheduler::INT_TYPE: return fun(*element.implicit_matrix_int);
case scheduler::UINT_TYPE: return fun(*element.implicit_matrix_uint);
case scheduler::LONG_TYPE: return fun(*element.implicit_matrix_long);
case scheduler::ULONG_TYPE: return fun(*element.implicit_matrix_ulong);
case scheduler::FLOAT_TYPE : return fun(*element.implicit_matrix_float);
case scheduler::DOUBLE_TYPE : return fun(*element.implicit_matrix_double);
default : throw generator_not_supported_exception("Unsupported Scalartype");
}
}
template<class Fun>
static typename Fun::result_type call_on_element(scheduler::lhs_rhs_element const & element, Fun const & fun){
switch (element.type_family)
{
case scheduler::SCALAR_TYPE_FAMILY:
if (element.subtype == scheduler::HOST_SCALAR_TYPE)
return call_on_host_scalar(element, fun);
else
return call_on_scalar(element, fun);
case scheduler::VECTOR_TYPE_FAMILY :
if (element.subtype == scheduler::IMPLICIT_VECTOR_TYPE)
return call_on_implicit_vector(element, fun);
else
return call_on_vector(element, fun);
case scheduler::MATRIX_TYPE_FAMILY:
if (element.subtype == scheduler::IMPLICIT_MATRIX_TYPE)
return call_on_implicit_matrix(element, fun);
else
return call_on_matrix(element,fun);
default:
throw generator_not_supported_exception("Unsupported datastructure type : Not among {Scalar, Vector, Matrix}");
}
}
struct scalartype_size_fun
{
typedef vcl_size_t result_type;
result_type operator()(float const &) const { return sizeof(float); }
result_type operator()(double const &) const { return sizeof(double); }
template<class T> result_type operator()(T const &) const { return sizeof(typename viennacl::result_of::cpu_value_type<T>::type); }
};
struct internal_size_fun
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::internal_size(t); }
};
struct size_fun
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::size(t); }
};
struct stride_fun
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::stride(t); }
};
struct start1_fun
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::start1(t); }
};
struct start2_fun
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::start2(t); }
};
struct leading_stride
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::row_major(t)?viennacl::traits::stride2(t):viennacl::traits::stride1(t); }
};
struct leading_start
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::row_major(t)?viennacl::traits::start2(t):viennacl::traits::start1(t); }
};
struct stride1_fun
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::stride1(t); }
};
struct stride2_fun
{
typedef vcl_size_t result_type;
template<class T> result_type operator()(T const &t) const { return viennacl::traits::stride2(t); }
};
struct handle_fun
{
typedef cl_mem result_type;
template<class T>
result_type operator()(T const &t) const { return viennacl::traits::opencl_handle(t); }
};
struct internal_size1_fun
{
typedef vcl_size_t result_type;
template<class T>
result_type operator()(T const &t) const { return viennacl::traits::internal_size1(t); }
};
struct row_major_fun
{
typedef bool result_type;
template<class T>
result_type operator()(T const &t) const { return viennacl::traits::row_major(t); }
};
struct internal_size2_fun
{
typedef vcl_size_t result_type;
template<class T>
result_type operator()(T const &t) const { return viennacl::traits::internal_size2(t); }
};
struct size1_fun
{
typedef vcl_size_t result_type;
template<class T>
result_type operator()(T const &t) const { return viennacl::traits::size1(t); }
};
struct size2_fun
{
typedef vcl_size_t result_type;
template<class T>
result_type operator()(T const &t) const { return viennacl::traits::size2(t); }
};
template<class T, class U>
struct is_same_type { enum { value = 0 }; };
template<class T>
struct is_same_type<T,T> { enum { value = 1 }; };
inline bool is_reduction(scheduler::statement_node const & node)
{
return node.op.type_family==scheduler::OPERATION_VECTOR_REDUCTION_TYPE_FAMILY
|| node.op.type_family==scheduler::OPERATION_COLUMNS_REDUCTION_TYPE_FAMILY
|| node.op.type_family==scheduler::OPERATION_ROWS_REDUCTION_TYPE_FAMILY
|| node.op.type==scheduler::OPERATION_BINARY_INNER_PROD_TYPE
|| node.op.type==scheduler::OPERATION_BINARY_MAT_VEC_PROD_TYPE;
}
inline bool is_index_reduction(scheduler::op_element const & op)
{
return op.type==scheduler::OPERATION_BINARY_ELEMENT_ARGFMAX_TYPE
|| op.type==scheduler::OPERATION_BINARY_ELEMENT_ARGMAX_TYPE
|| op.type==scheduler::OPERATION_BINARY_ELEMENT_ARGFMIN_TYPE
|| op.type==scheduler::OPERATION_BINARY_ELEMENT_ARGMIN_TYPE;
}
template<class T>
struct type_to_string;
template<> struct type_to_string<unsigned char> { static const char * value() { return "uchar"; } };
template<> struct type_to_string<char> { static const char * value() { return "char"; } };
template<> struct type_to_string<unsigned short> { static const char * value() { return "ushort"; } };
template<> struct type_to_string<short> { static const char * value() { return "short"; } };
template<> struct type_to_string<unsigned int> { static const char * value() { return "uint"; } };
template<> struct type_to_string<int> { static const char * value() { return "int"; } };
template<> struct type_to_string<unsigned long> { static const char * value() { return "ulong"; } };
template<> struct type_to_string<long> { static const char * value() { return "long"; } };
template<> struct type_to_string<float> { static const char * value() { return "float"; } };
template<> struct type_to_string<double> { static const char * value() { return "double"; } };
template<class T>
struct first_letter_of_type;
template<> struct first_letter_of_type<char> { static char value() { return 'c'; } };
template<> struct first_letter_of_type<unsigned char> { static char value() { return 'd'; } };
template<> struct first_letter_of_type<short> { static char value() { return 's'; } };
template<> struct first_letter_of_type<unsigned short> { static char value() { return 't'; } };
template<> struct first_letter_of_type<int> { static char value() { return 'i'; } };
template<> struct first_letter_of_type<unsigned int> { static char value() { return 'j'; } };
template<> struct first_letter_of_type<long> { static char value() { return 'l'; } };
template<> struct first_letter_of_type<unsigned long> { static char value() { return 'm'; } };
template<> struct first_letter_of_type<float> { static char value() { return 'f'; } };
template<> struct first_letter_of_type<double> { static char value() { return 'd'; } };
class kernel_generation_stream : public std::ostream
{
class kgenstream : public std::stringbuf
{
public:
kgenstream(std::ostringstream& osstream,unsigned int const & tab_count) : oss_(osstream), tab_count_(tab_count){ }
int sync() {
for (unsigned int i=0; i<tab_count_;++i)
oss_ << " ";
oss_ << str();
str("");
return !oss_;
}
#if defined(_MSC_VER)
~kgenstream() throw() { pubsync(); }
#else
~kgenstream() { pubsync(); }
#endif
private:
std::ostream& oss_;
unsigned int const & tab_count_;
};
public:
kernel_generation_stream() : std::ostream(new kgenstream(oss,tab_count_)), tab_count_(0){ }
#if defined(_MSC_VER)
~kernel_generation_stream() throw() { delete rdbuf(); }
#else
~kernel_generation_stream(){ delete rdbuf(); }
#endif
std::string str(){ return oss.str(); }
void inc_tab(){ ++tab_count_; }
void dec_tab(){ --tab_count_; }
private:
unsigned int tab_count_;
std::ostringstream oss;
};
inline bool node_leaf(scheduler::op_element const & op)
{
using namespace scheduler;
return op.type==OPERATION_UNARY_NORM_1_TYPE
|| op.type==OPERATION_UNARY_NORM_2_TYPE
|| op.type==OPERATION_UNARY_NORM_INF_TYPE
|| op.type==OPERATION_UNARY_TRANS_TYPE
|| op.type==OPERATION_BINARY_MAT_VEC_PROD_TYPE
|| op.type==OPERATION_BINARY_MAT_MAT_PROD_TYPE
|| op.type==OPERATION_BINARY_INNER_PROD_TYPE
|| op.type==OPERATION_BINARY_MATRIX_DIAG_TYPE
|| op.type==OPERATION_BINARY_VECTOR_DIAG_TYPE
|| op.type==OPERATION_BINARY_MATRIX_ROW_TYPE
|| op.type==OPERATION_BINARY_MATRIX_COLUMN_TYPE
|| op.type_family==OPERATION_VECTOR_REDUCTION_TYPE_FAMILY
|| op.type_family==OPERATION_ROWS_REDUCTION_TYPE_FAMILY
|| op.type_family==OPERATION_COLUMNS_REDUCTION_TYPE_FAMILY;
}
inline bool elementwise_operator(scheduler::op_element const & op)
{
using namespace scheduler;
return op.type== OPERATION_BINARY_ASSIGN_TYPE
|| op.type== OPERATION_BINARY_INPLACE_ADD_TYPE
|| op.type== OPERATION_BINARY_INPLACE_SUB_TYPE
|| op.type== OPERATION_BINARY_ADD_TYPE
|| op.type== OPERATION_BINARY_SUB_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_PROD_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_DIV_TYPE
|| op.type== OPERATION_BINARY_MULT_TYPE
|| op.type== OPERATION_BINARY_DIV_TYPE;
}
inline bool elementwise_function(scheduler::op_element const & op)
{
using namespace scheduler;
return
op.type == OPERATION_UNARY_CAST_CHAR_TYPE
|| op.type == OPERATION_UNARY_CAST_UCHAR_TYPE
|| op.type == OPERATION_UNARY_CAST_SHORT_TYPE
|| op.type == OPERATION_UNARY_CAST_USHORT_TYPE
|| op.type == OPERATION_UNARY_CAST_INT_TYPE
|| op.type == OPERATION_UNARY_CAST_UINT_TYPE
|| op.type == OPERATION_UNARY_CAST_LONG_TYPE
|| op.type == OPERATION_UNARY_CAST_ULONG_TYPE
|| op.type == OPERATION_UNARY_CAST_HALF_TYPE
|| op.type == OPERATION_UNARY_CAST_FLOAT_TYPE
|| op.type == OPERATION_UNARY_CAST_DOUBLE_TYPE
|| op.type== OPERATION_UNARY_ABS_TYPE
|| op.type== OPERATION_UNARY_ACOS_TYPE
|| op.type== OPERATION_UNARY_ASIN_TYPE
|| op.type== OPERATION_UNARY_ATAN_TYPE
|| op.type== OPERATION_UNARY_CEIL_TYPE
|| op.type== OPERATION_UNARY_COS_TYPE
|| op.type== OPERATION_UNARY_COSH_TYPE
|| op.type== OPERATION_UNARY_EXP_TYPE
|| op.type== OPERATION_UNARY_FABS_TYPE
|| op.type== OPERATION_UNARY_FLOOR_TYPE
|| op.type== OPERATION_UNARY_LOG_TYPE
|| op.type== OPERATION_UNARY_LOG10_TYPE
|| op.type== OPERATION_UNARY_SIN_TYPE
|| op.type== OPERATION_UNARY_SINH_TYPE
|| op.type== OPERATION_UNARY_SQRT_TYPE
|| op.type== OPERATION_UNARY_TAN_TYPE
|| op.type== OPERATION_UNARY_TANH_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_POW_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_EQ_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_NEQ_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_GREATER_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_LESS_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_GEQ_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_LEQ_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_FMAX_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_FMIN_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_MAX_TYPE
|| op.type== OPERATION_BINARY_ELEMENT_MIN_TYPE;
}
inline scheduler::lhs_rhs_element & lhs_rhs_element(scheduler::statement const & st, vcl_size_t idx, leaf_t leaf)
{
using namespace tree_parsing;
assert(leaf==LHS_NODE_TYPE || leaf==RHS_NODE_TYPE);
if (leaf==LHS_NODE_TYPE)
return const_cast<scheduler::lhs_rhs_element &>(st.array()[idx].lhs);
return const_cast<scheduler::lhs_rhs_element &>(st.array()[idx].rhs);
}
inline unsigned int size_of(scheduler::statement_node_numeric_type type)
{
using namespace scheduler;
switch (type)
{
case UCHAR_TYPE:
case CHAR_TYPE: return 1;
case USHORT_TYPE:
case SHORT_TYPE:
case HALF_TYPE: return 2;
case UINT_TYPE:
case INT_TYPE:
case FLOAT_TYPE: return 4;
case ULONG_TYPE:
case LONG_TYPE:
case DOUBLE_TYPE: return 8;
default: throw generator_not_supported_exception("Unsupported scalartype");
}
}
inline std::string append_width(std::string const & str, unsigned int width)
{
if (width==1)
return str;
return str + tools::to_string(width);
}
}
}
}
#endif
viennacl::device_specific::utils::first_letter_of_type< unsigned short >::value
static char value()
Definition: utils.hpp:396
viennacl::scheduler::OPERATION_UNARY_NORM_INF_TYPE
Definition: forwards.h:109
viennacl::device_specific::utils::row_major_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:331
viennacl::scheduler::OPERATION_UNARY_ATAN_TYPE
Definition: forwards.h:91
viennacl::device_specific::utils::kernel_generation_stream::~kernel_generation_stream
~kernel_generation_stream()
Definition: utils.hpp:432
forwards.h
This file provides the forward declarations for the OpenCL layer of ViennaCL.
viennacl::device_specific::utils::internal_size_fun
Definition: utils.hpp:259
viennacl::scheduler::INT_TYPE
Definition: forwards.h:294
viennacl::scheduler::OPERATION_BINARY_VECTOR_DIAG_TYPE
Definition: forwards.h:141
viennacl::scheduler::HALF_TYPE
Definition: forwards.h:298
viennacl::device_specific::utils::first_letter_of_type
Definition: utils.hpp:392
viennacl::device_specific::utils::type_to_string< long >::value
static const char * value()
Definition: utils.hpp:386
viennacl::device_specific::utils::size1_fun
Definition: utils.hpp:341
viennacl::device_specific::utils::handle_fun::result_type
cl_mem result_type
Definition: utils.hpp:315
viennacl::scheduler::OPERATION_BINARY_MAT_VEC_PROD_TYPE
Definition: forwards.h:144
viennacl::scheduler::OPERATION_BINARY_ELEMENT_POW_TYPE
Definition: forwards.h:134
viennacl::traits::stride1
result_of::size_type< matrix_base< NumericT > >::type stride1(matrix_base< NumericT > const &s)
Definition: stride.hpp:55
viennacl::device_specific::utils::stride_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:273
viennacl::device_specific::utils::first_letter_of_type< int >::value
static char value()
Definition: utils.hpp:397
viennacl::scheduler::OPERATION_VECTOR_REDUCTION_TYPE_FAMILY
Definition: forwards.h:62
viennacl::scheduler::OPERATION_BINARY_ELEMENT_MAX_TYPE
Definition: forwards.h:137
viennacl::device_specific::generator_not_supported_exception
Exception for the case the generator is unable to deal with the operation.
Definition: forwards.h:163
size.hpp
Generic size and resize functionality for different vector and matrix types.
viennacl::scheduler::OPERATION_UNARY_TRANS_TYPE
Definition: forwards.h:106
viennacl::scheduler::UCHAR_TYPE
Definition: forwards.h:291
viennacl::scheduler::OPERATION_UNARY_COS_TYPE
Definition: forwards.h:93
viennacl::device_specific::utils::internal_size1_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:324
viennacl::device_specific::utils::size1_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:345
viennacl::scheduler::OPERATION_UNARY_FABS_TYPE
Definition: forwards.h:96
viennacl::device_specific::utils::first_letter_of_type< char >::value
static char value()
Definition: utils.hpp:393
viennacl::scheduler::OPERATION_BINARY_ELEMENT_GEQ_TYPE
Definition: forwards.h:131
viennacl::device_specific::utils::scalartype_size_fun
Definition: utils.hpp:251
tools.hpp
Various little tools used here and there in ViennaCL.
viennacl::traits::internal_size1
vcl_size_t internal_size1(matrix_base< NumericT > const &mat)
Helper routine for obtaining the internal number of entries per row of a ViennaCL matrix...
Definition: size.hpp:386
viennacl::device_specific::utils::stride2_fun
Definition: utils.hpp:307
viennacl::device_specific::utils::start1_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:280
viennacl::device_specific::utils::size2_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:352
viennacl::traits::size1
vcl_size_t size1(MatrixType const &mat)
Generic routine for obtaining the number of rows of a matrix (ViennaCL, uBLAS, etc.)
Definition: size.hpp:163
viennacl::device_specific::utils::stride_fun
Definition: utils.hpp:271
viennacl::scheduler::OPERATION_BINARY_ELEMENT_DIV_TYPE
Definition: forwards.h:127
viennacl::scheduler::DOUBLE_TYPE
Definition: forwards.h:300
viennacl::device_specific::utils::internal_size2_fun
Definition: utils.hpp:334
viennacl::scheduler::OPERATION_UNARY_ACOS_TYPE
Definition: forwards.h:89
viennacl::traits::internal_size2
vcl_size_t internal_size2(matrix_base< NumericT > const &mat)
Helper routine for obtaining the internal number of entries per column of a ViennaCL matrix...
Definition: size.hpp:394
viennacl::scheduler::OPERATION_UNARY_EXP_TYPE
Definition: forwards.h:95
viennacl::device_specific::utils::row_major_fun::result_type
bool result_type
Definition: utils.hpp:329
viennacl::device_specific::utils::elementwise_operator
bool elementwise_operator(scheduler::op_element const &op)
Definition: utils.hpp:462
viennacl::device_specific::utils::type_to_string< char >::value
static const char * value()
Definition: utils.hpp:380
viennacl::scheduler::OPERATION_UNARY_CAST_FLOAT_TYPE
Definition: forwards.h:85
viennacl::device_specific::utils::stride1_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:304
viennacl::device_specific::leaf_t
leaf_t
Definition: forwards.h:180
viennacl::scheduler::MATRIX_TYPE_FAMILY
Definition: forwards.h:260
viennacl::scheduler::OPERATION_UNARY_ASIN_TYPE
Definition: forwards.h:90
viennacl::device_specific::utils::type_to_string< double >::value
static const char * value()
Definition: utils.hpp:388
viennacl::scheduler::statement_node::op
op_element op
Definition: forwards.h:481
viennacl::scheduler::SCALAR_TYPE_FAMILY
Definition: forwards.h:254
viennacl::traits::stride
result_of::size_type< viennacl::vector_base< T > >::type stride(viennacl::vector_base< T > const &s)
Definition: stride.hpp:45
viennacl::traits::start1
result_of::size_type< T >::type start1(T const &obj)
Definition: start.hpp:65
viennacl::scheduler::OPERATION_BINARY_MATRIX_ROW_TYPE
Definition: forwards.h:142
viennacl::scheduler::OPERATION_BINARY_ASSIGN_TYPE
Definition: forwards.h:115
viennacl::scheduler::OPERATION_UNARY_CAST_HALF_TYPE
Definition: forwards.h:84
viennacl::scheduler::lhs_rhs_element
A class representing the 'data' for the LHS or RHS operand of the respective node.
Definition: forwards.h:337
viennacl::scheduler::statement::array
container_type const & array() const
Definition: forwards.h:528
viennacl::device_specific::utils::first_letter_of_type< long >::value
static char value()
Definition: utils.hpp:399
viennacl::scheduler::HOST_SCALAR_TYPE
Definition: forwards.h:268
viennacl::traits::internal_size
vcl_size_t internal_size(vector_base< NumericT > const &vec)
Helper routine for obtaining the buffer length of a ViennaCL vector.
Definition: size.hpp:375
viennacl::scheduler::statement_node_numeric_type
statement_node_numeric_type
Encodes the type of a node in the statement tree.
Definition: forwards.h:286
viennacl::device_specific::utils::is_same_type::value
Definition: utils.hpp:356
viennacl::device_specific::utils::node_leaf
bool node_leaf(scheduler::op_element const &op)
Definition: utils.hpp:443
matrix_def.hpp
Forward declaration of dense matrix classes.
viennacl::device_specific::utils::size_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:267
viennacl::device_specific::utils::opencl_source_to_cuda_source
std::string opencl_source_to_cuda_source(std::string const &opencl_src)
Definition: utils.hpp:50
viennacl::scheduler::op_element::type_family
operation_node_type_family type_family
Definition: forwards.h:473
viennacl::scheduler::op_element
Struct for holding the type family as well as the type of an operation (could be addition, subtraction, norm, etc.)
Definition: forwards.h:471
viennacl::scheduler::OPERATION_UNARY_CAST_UCHAR_TYPE
Definition: forwards.h:77
viennacl::traits::size2
result_of::size_type< MatrixType >::type size2(MatrixType const &mat)
Generic routine for obtaining the number of columns of a matrix (ViennaCL, uBLAS, etc...
Definition: size.hpp:201
viennacl::scheduler::OPERATION_COLUMNS_REDUCTION_TYPE_FAMILY
Definition: forwards.h:64
viennacl::scheduler::OPERATION_BINARY_ELEMENT_PROD_TYPE
Definition: forwards.h:126
viennacl::device_specific::utils::scalartype_size_fun::operator()
result_type operator()(float const &) const
Definition: utils.hpp:254
viennacl::device_specific::utils::kernel_generation_stream::str
std::string str()
Definition: utils.hpp:435
viennacl::device_specific::utils::internal_size_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:262
viennacl::device_specific::utils::stride1_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:303
viennacl::scheduler::OPERATION_UNARY_NORM_2_TYPE
Definition: forwards.h:108
viennacl::scheduler::SHORT_TYPE
Definition: forwards.h:292
viennacl::device_specific::utils::size2_fun
Definition: utils.hpp:348
viennacl::scheduler::LONG_TYPE
Definition: forwards.h:296
viennacl::scheduler::OPERATION_BINARY_ELEMENT_FMIN_TYPE
Definition: forwards.h:136
viennacl::traits::size
vcl_size_t size(VectorType const &vec)
Generic routine for obtaining the size of a vector (ViennaCL, uBLAS, etc.)
Definition: size.hpp:239
viennacl::traits::start2
result_of::size_type< T >::type start2(T const &obj)
Definition: start.hpp:84
viennacl::device_specific::utils::kernel_generation_stream
Definition: utils.hpp:404
viennacl::scheduler::OPERATION_UNARY_LOG_TYPE
Definition: forwards.h:98
viennacl::device_specific::utils::elementwise_function
bool elementwise_function(scheduler::op_element const &op)
Definition: utils.hpp:476
viennacl::device_specific::utils::type_to_string< int >::value
static const char * value()
Definition: utils.hpp:384
viennacl::scheduler::OPERATION_BINARY_ELEMENT_GREATER_TYPE
Definition: forwards.h:130
viennacl::device_specific::utils::internal_size_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:261
viennacl::scheduler::OPERATION_UNARY_CAST_LONG_TYPE
Definition: forwards.h:82
viennacl::device_specific::utils::start2_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:285
viennacl::scheduler::OPERATION_BINARY_INNER_PROD_TYPE
Definition: forwards.h:146
viennacl::device_specific::utils::internal_size1_fun
Definition: utils.hpp:320
viennacl::device_specific::utils::type_to_string< unsigned long >::value
static const char * value()
Definition: utils.hpp:385
viennacl::scheduler::OPERATION_UNARY_SQRT_TYPE
Definition: forwards.h:102
viennacl::scheduler::CHAR_TYPE
Definition: forwards.h:290
viennacl::device_specific::utils::row_major_fun
Definition: utils.hpp:327
viennacl::tools::find_and_replace
int find_and_replace(std::string &source, std::string const &find, std::string const &replace)
Replace in a source string a pattern by another.
Definition: tools.hpp:160
viennacl::scheduler::OPERATION_BINARY_INPLACE_ADD_TYPE
Definition: forwards.h:116
viennacl::scheduler::OPERATION_UNARY_SINH_TYPE
Definition: forwards.h:101
viennacl::device_specific::RHS_NODE_TYPE
Definition: forwards.h:184
viennacl::device_specific::LHS_NODE_TYPE
Definition: forwards.h:182
viennacl::scheduler::OPERATION_BINARY_ELEMENT_NEQ_TYPE
Definition: forwards.h:129
viennacl::device_specific::utils::scalartype_size_fun::operator()
result_type operator()(T const &) const
Definition: utils.hpp:256
viennacl::scheduler::OPERATION_UNARY_SIN_TYPE
Definition: forwards.h:100
viennacl::device_specific::utils::type_to_string< short >::value
static const char * value()
Definition: utils.hpp:382
viennacl::device_specific::utils::type_to_string< unsigned short >::value
static const char * value()
Definition: utils.hpp:381
viennacl::device_specific::utils::scalartype_size_fun::operator()
result_type operator()(double const &) const
Definition: utils.hpp:255
viennacl::scheduler::OPERATION_UNARY_CAST_USHORT_TYPE
Definition: forwards.h:79
viennacl::device_specific::utils::start2_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:286
viennacl::scheduler::OPERATION_UNARY_CAST_DOUBLE_TYPE
Definition: forwards.h:86
viennacl::device_specific::utils::leading_stride
Definition: utils.hpp:289
viennacl::device_specific::utils::kernel_generation_stream::dec_tab
void dec_tab()
Definition: utils.hpp:437
viennacl::scheduler::OPERATION_UNARY_ABS_TYPE
Definition: forwards.h:88
viennacl::device_specific::utils::is_reduction
bool is_reduction(scheduler::statement_node const &node)
Definition: utils.hpp:361
forwards.h
Forwards declaration.
row_major.hpp
Determines whether a given expression has a row-major matrix layout.
viennacl::scheduler::OPERATION_UNARY_CAST_UINT_TYPE
Definition: forwards.h:81
viennacl::vcl_size_t
std::size_t vcl_size_t
Definition: forwards.h:75
viennacl::device_specific::utils::leading_stride::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:292
viennacl::scheduler::OPERATION_BINARY_MATRIX_DIAG_TYPE
Definition: forwards.h:140
forwards.h
Provides the datastructures for dealing with a single statement such as 'x = y + z;'.
viennacl::scheduler::OPERATION_UNARY_LOG10_TYPE
Definition: forwards.h:99
viennacl::device_specific::utils::start1_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:279
viennacl::scheduler::OPERATION_UNARY_FLOOR_TYPE
Definition: forwards.h:97
viennacl::scheduler::OPERATION_BINARY_SUB_TYPE
Definition: forwards.h:119
viennacl::device_specific::utils::type_to_string< unsigned int >::value
static const char * value()
Definition: utils.hpp:383
viennacl::scheduler::USHORT_TYPE
Definition: forwards.h:293
viennacl::result_of::cpu_value_type::type
T::ERROR_CANNOT_DEDUCE_CPU_SCALAR_TYPE_FOR_T type
Definition: result_of.hpp:271
viennacl::device_specific::utils::kernel_generation_stream::inc_tab
void inc_tab()
Definition: utils.hpp:436
viennacl::traits::stride2
result_of::size_type< matrix_base< NumericT > >::type stride2(matrix_base< NumericT > const &s)
Definition: stride.hpp:65
viennacl::scheduler::IMPLICIT_VECTOR_TYPE
Definition: forwards.h:272
viennacl::scheduler::OPERATION_UNARY_CAST_INT_TYPE
Definition: forwards.h:80
viennacl::device_specific::utils::kernel_generation_stream::kernel_generation_stream
kernel_generation_stream()
Definition: utils.hpp:428
viennacl::scheduler::OPERATION_BINARY_INPLACE_SUB_TYPE
Definition: forwards.h:117
viennacl::scheduler::OPERATION_BINARY_MATRIX_COLUMN_TYPE
Definition: forwards.h:143
viennacl::device_specific::utils::leading_start::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:298
viennacl::device_specific::utils::size1_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:343
viennacl::scheduler::OPERATION_BINARY_ELEMENT_ARGMAX_TYPE
Definition: forwards.h:124
viennacl::device_specific::utils::start1_fun
Definition: utils.hpp:277
viennacl::scheduler::IMPLICIT_MATRIX_TYPE
Definition: forwards.h:275
viennacl::device_specific::utils::lhs_rhs_element
scheduler::lhs_rhs_element & lhs_rhs_element(scheduler::statement const &st, vcl_size_t idx, leaf_t leaf)
Definition: utils.hpp:525
viennacl::traits::row_major
bool row_major(T const &)
Definition: row_major.hpp:38
viennacl::device_specific::utils::type_to_string< float >::value
static const char * value()
Definition: utils.hpp:387
viennacl::scheduler::OPERATION_UNARY_TAN_TYPE
Definition: forwards.h:103
viennacl::device_specific::utils::internal_size2_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:336
viennacl::device_specific::utils::leading_start
Definition: utils.hpp:295
viennacl::device_specific::utils::first_letter_of_type< short >::value
static char value()
Definition: utils.hpp:395
viennacl::device_specific::utils::handle_fun
Definition: utils.hpp:313
viennacl::device_specific::utils::scalartype_size_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:253
viennacl::scheduler::OPERATION_UNARY_CEIL_TYPE
Definition: forwards.h:92
viennacl::scheduler::op_element::type
operation_node_type type
Definition: forwards.h:474
viennacl::scheduler::OPERATION_BINARY_DIV_TYPE
Definition: forwards.h:121
viennacl::device_specific::utils::stride1_fun
Definition: utils.hpp:301
viennacl::device_specific::utils::is_index_reduction
bool is_index_reduction(scheduler::op_element const &op)
Definition: utils.hpp:370
viennacl::device_specific::utils::first_letter_of_type< unsigned long >::value
static char value()
Definition: utils.hpp:400
viennacl::scheduler::OPERATION_BINARY_ELEMENT_FMAX_TYPE
Definition: forwards.h:135
viennacl::scheduler::UINT_TYPE
Definition: forwards.h:295
viennacl::scheduler::OPERATION_BINARY_ELEMENT_ARGFMAX_TYPE
Definition: forwards.h:122
viennacl::scheduler::OPERATION_UNARY_TANH_TYPE
Definition: forwards.h:104
viennacl::scheduler::FLOAT_TYPE
Definition: forwards.h:299
viennacl::scheduler::OPERATION_UNARY_CAST_SHORT_TYPE
Definition: forwards.h:78
viennacl::device_specific::utils::first_letter_of_type< double >::value
static char value()
Definition: utils.hpp:402
viennacl::scheduler::statement
The main class for representing a statement such as x = inner_prod(y,z); at runtime.
Definition: forwards.h:502
viennacl::scheduler::OPERATION_BINARY_ADD_TYPE
Definition: forwards.h:118
viennacl::device_specific::utils::is_same_type
Definition: utils.hpp:356
vector_def.hpp
Forward declarations of the implicit_vector_base, vector_base class.
viennacl::device_specific::utils::type_to_string
Definition: utils.hpp:378
viennacl::device_specific::utils::internal_size2_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:338
viennacl::scheduler::OPERATION_BINARY_ELEMENT_ARGMIN_TYPE
Definition: forwards.h:125
viennacl::scheduler::OPERATION_BINARY_ELEMENT_ARGFMIN_TYPE
Definition: forwards.h:123
handle.hpp
Extracts the underlying OpenCL handle from a vector, a matrix, an expression etc. ...
viennacl::device_specific::utils::size_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:268
viennacl::scheduler::OPERATION_BINARY_ELEMENT_MIN_TYPE
Definition: forwards.h:138
viennacl::scheduler::ULONG_TYPE
Definition: forwards.h:297
viennacl::tools::to_string
std::string to_string(T const t)
Definition: tools.hpp:304
viennacl::scheduler::OPERATION_ROWS_REDUCTION_TYPE_FAMILY
Definition: forwards.h:63
viennacl::scheduler::OPERATION_UNARY_CAST_ULONG_TYPE
Definition: forwards.h:83
viennacl::device_specific::utils::stride2_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:310
viennacl::device_specific::utils::type_to_string< unsigned char >::value
static const char * value()
Definition: utils.hpp:379
viennacl::device_specific::utils::leading_start::result_type
vcl_size_t result_type
Definition: utils.hpp:297
viennacl::scheduler::OPERATION_UNARY_COSH_TYPE
Definition: forwards.h:94
viennacl::scheduler::VECTOR_TYPE_FAMILY
Definition: forwards.h:257
viennacl::scheduler::OPERATION_BINARY_ELEMENT_EQ_TYPE
Definition: forwards.h:128
viennacl::device_specific::utils::stride2_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:309
viennacl::scheduler::OPERATION_UNARY_CAST_CHAR_TYPE
Definition: forwards.h:76
viennacl::scheduler::statement_node
Main datastructure for an node in the statement tree.
Definition: forwards.h:478
viennacl::device_specific::utils::first_letter_of_type< unsigned char >::value
static char value()
Definition: utils.hpp:394
viennacl::device_specific::utils::size_of
unsigned int size_of(scheduler::statement_node_numeric_type type)
Definition: utils.hpp:534
viennacl::scheduler::OPERATION_BINARY_MAT_MAT_PROD_TYPE
Definition: forwards.h:145
viennacl::device_specific::utils::internal_size1_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:322
viennacl::device_specific::utils::handle_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:317
viennacl::device_specific::utils::size_fun
Definition: utils.hpp:265
viennacl::device_specific::utils::first_letter_of_type< unsigned int >::value
static char value()
Definition: utils.hpp:398
viennacl::device_specific::utils::stride_fun::operator()
result_type operator()(T const &t) const
Definition: utils.hpp:274
viennacl::scheduler::OPERATION_BINARY_ELEMENT_LEQ_TYPE
Definition: forwards.h:133
viennacl::scheduler::OPERATION_UNARY_NORM_1_TYPE
Definition: forwards.h:107
viennacl::device_specific::utils::append_width
std::string append_width(std::string const &str, unsigned int width)
Definition: utils.hpp:558
viennacl::device_specific::utils::first_letter_of_type< float >::value
static char value()
Definition: utils.hpp:401
viennacl::device_specific::utils::size2_fun::result_type
vcl_size_t result_type
Definition: utils.hpp:350
viennacl::device_specific::utils::leading_stride::result_type
vcl_size_t result_type
Definition: utils.hpp:291
viennacl::scheduler::OPERATION_BINARY_ELEMENT_LESS_TYPE
Definition: forwards.h:132
viennacl::device_specific::utils::start2_fun
Definition: utils.hpp:283
viennacl::scheduler::OPERATION_BINARY_MULT_TYPE
Definition: forwards.h:120