All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
template_base.hpp
Go to the documentation of this file.
#ifndef VIENNACL_DEVICE_SPECIFIC_TEMPLATES_TEMPLATE_BASE_
#define VIENNACL_DEVICE_SPECIFIC_TEMPLATES_TEMPLATE_BASE_
/* =========================================================================
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 <list>
#include <set>
#include "viennacl/ocl/kernel.hpp"
#include "viennacl/ocl/device.hpp"
#include "viennacl/ocl/device_utils.hpp"
#include "viennacl/scheduler/forwards.h"
#include "viennacl/scheduler/io.hpp"
#include "viennacl/device_specific/lazy_program_compiler.hpp"
#include "viennacl/device_specific/mapped_objects.hpp"
#include "viennacl/device_specific/tree_parsing.hpp"
#include "viennacl/device_specific/utils.hpp"
namespace viennacl
{
namespace device_specific
{
enum fetching_policy_type
{
FETCH_FROM_LOCAL,
FETCH_FROM_GLOBAL_STRIDED,
FETCH_FROM_GLOBAL_CONTIGUOUS
};
class template_base
{
public:
struct parameters_type
{
parameters_type(unsigned int _simd_width, unsigned int _local_size_1, unsigned int _local_size_2, unsigned int _num_kernels) : simd_width(_simd_width), local_size_0(_local_size_1), local_size_1(_local_size_2), num_kernels(_num_kernels){ }
unsigned int simd_width;
unsigned int local_size_0;
unsigned int local_size_1;
unsigned int num_kernels;
};
private:
class map_functor : public tree_parsing::traversal_functor
{
scheduler::statement_node_numeric_type numeric_type(scheduler::statement const * statement, vcl_size_t root_idx) const
{
scheduler::statement_node const * root_node = &statement->array()[root_idx];
while (root_node->lhs.numeric_type==scheduler::INVALID_NUMERIC_TYPE)
root_node = &statement->array()[root_node->lhs.node_index];
return root_node->lhs.numeric_type;
}
public:
typedef tools::shared_ptr<mapped_object> result_type;
map_functor(symbolic_binder & binder, mapping_type & mapping) : binder_(binder), mapping_(mapping){ }
template<class T>
result_type binary_leaf(scheduler::statement const * statement, vcl_size_t root_idx, mapping_type const * mapping) const
{
return result_type(new T(utils::numeric_type_to_string(numeric_type(statement,root_idx)), binder_.get(NULL), mapped_object::node_info(mapping, statement, root_idx)));
}
template<class NumericT>
result_type operator()(NumericT const & /*scalar*/) const
{
return result_type(new mapped_host_scalar(utils::type_to_string<NumericT>::value(), binder_.get(NULL)));
}
template<class NumericT>
result_type operator()(scalar<NumericT> const & scal) const
{
return result_type(new mapped_scalar(utils::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(scal))));
}
template<class NumericT>
result_type operator()(vector_base<NumericT> const & vec) const
{
return result_type(new mapped_vector(utils::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(vec))));
}
template<class NumericT>
result_type operator()(implicit_vector_base<NumericT> const & /*vec*/) const
{
return result_type(new mapped_implicit_vector(utils::type_to_string<NumericT>::value(), binder_.get(NULL)));
}
template<class NumericT>
result_type operator()(matrix_base<NumericT> const & mat) const
{
return result_type(new mapped_matrix(utils::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(mat)),
viennacl::traits::row_major(mat)));
}
template<class NumericT>
result_type operator()(implicit_matrix_base<NumericT> const & /*mat*/) const
{
return result_type(new mapped_implicit_matrix(utils::type_to_string<NumericT>::value(), binder_.get(NULL)));
}
void operator()(scheduler::statement const & statement, vcl_size_t root_idx, leaf_t leaf_t) const {
mapping_type::key_type key(root_idx, leaf_t);
scheduler::statement_node const & root_node = statement.array()[root_idx];
if (leaf_t == LHS_NODE_TYPE && root_node.lhs.type_family != scheduler::COMPOSITE_OPERATION_FAMILY)
mapping_.insert(mapping_type::value_type(key, utils::call_on_element(root_node.lhs, *this)));
else if (leaf_t == RHS_NODE_TYPE && root_node.rhs.type_family != scheduler::COMPOSITE_OPERATION_FAMILY)
mapping_.insert(mapping_type::value_type(key, utils::call_on_element(root_node.rhs, *this)));
else if ( leaf_t== PARENT_NODE_TYPE)
{
if (root_node.op.type==scheduler::OPERATION_BINARY_VECTOR_DIAG_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_vector_diag>(&statement, root_idx, &mapping_)));
else if (root_node.op.type==scheduler::OPERATION_BINARY_MATRIX_DIAG_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_diag>(&statement, root_idx, &mapping_)));
else if (root_node.op.type==scheduler::OPERATION_BINARY_MATRIX_ROW_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_row>(&statement, root_idx, &mapping_)));
else if (root_node.op.type==scheduler::OPERATION_BINARY_MATRIX_COLUMN_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_column>(&statement, root_idx, &mapping_)));
else if (is_scalar_reduction(root_node))
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_scalar_reduction>(&statement, root_idx, &mapping_)));
else if (is_vector_reduction(root_node))
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_row_wise_reduction>(&statement, root_idx, &mapping_)));
else if (root_node.op.type == scheduler::OPERATION_BINARY_MAT_MAT_PROD_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_product>(&statement, root_idx, &mapping_)));
else if (root_node.op.type == scheduler::OPERATION_UNARY_TRANS_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_trans>(&statement, root_idx, &mapping_)));
}
}
private:
symbolic_binder & binder_;
mapping_type & mapping_;
};
class prototype_generation_traversal : public tree_parsing::traversal_functor
{
private:
std::set<std::string> & already_generated_;
std::string & str_;
mapping_type const & mapping_;
std::map<std::string, unsigned int> const & widths_;
public:
prototype_generation_traversal(std::set<std::string> & already_generated, std::string & str, mapping_type const & mapping, std::map<std::string, unsigned int> const & widths) :
already_generated_(already_generated), str_(str), mapping_(mapping), widths_(widths){ }
void operator()(scheduler::statement const & statement, vcl_size_t root_idx, leaf_t leaf) const
{
scheduler::statement_node const & root_node = statement.array()[root_idx];
if ( (leaf==LHS_NODE_TYPE && root_node.lhs.type_family!=scheduler::COMPOSITE_OPERATION_FAMILY)
||(leaf==RHS_NODE_TYPE && root_node.rhs.type_family!=scheduler::COMPOSITE_OPERATION_FAMILY) )
{
mapped_object * obj = at(mapping_, std::make_pair(root_idx,leaf)).get();
if(widths_.find(obj->name())!=widths_.end())
obj->append_kernel_arguments(already_generated_, str_, at(widths_, obj->name()));
else
obj->append_kernel_arguments(already_generated_, str_, 1);
}
}
};
class set_arguments_functor : public tree_parsing::traversal_functor
{
public:
typedef void result_type;
set_arguments_functor(symbolic_binder & binder, unsigned int & current_arg, viennacl::ocl::kernel & kernel) : binder_(binder), current_arg_(current_arg), kernel_(kernel){ }
template<class NumericT>
result_type operator()(NumericT const & scal) const {
typedef typename viennacl::result_of::cl_type<NumericT>::type cl_scalartype;
kernel_.arg(current_arg_++, cl_scalartype(scal));
}
template<class NumericT>
result_type operator()(scalar<NumericT> const & scal) const {
if (binder_.bind(&viennacl::traits::handle(scal)))
kernel_.arg(current_arg_++, scal.handle().opencl_handle());
}
template<class NumericT>
result_type operator()(vector_base<NumericT> const & vec) const {
if (binder_.bind(&viennacl::traits::handle(vec)))
{
kernel_.arg(current_arg_++, vec.handle().opencl_handle());
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::start(vec)));
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::stride(vec)));
}
}
template<class NumericT>
result_type operator()(implicit_vector_base<NumericT> const & vec) const
{
typedef typename viennacl::result_of::cl_type<NumericT>::type cl_scalartype;
kernel_.arg(current_arg_++, cl_scalartype(vec.value()));
if (vec.has_index())
kernel_.arg(current_arg_++, cl_uint(vec.index()));
}
template<class NumericT>
result_type operator()(matrix_base<NumericT> const & mat) const
{
if (binder_.bind(&viennacl::traits::handle(mat)))
{
kernel_.arg(current_arg_++, mat.handle().opencl_handle());
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::ld(mat)));
if (mat.row_major())
{
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::start2(mat)));
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::start1(mat)));
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::stride2(mat)));
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::stride1(mat)));
}
else
{
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::start1(mat)));
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::start2(mat)));
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::stride1(mat)));
kernel_.arg(current_arg_++, cl_uint(viennacl::traits::stride2(mat)));
}
}
}
template<class NumericT>
result_type operator()(implicit_matrix_base<NumericT> const & mat) const
{
kernel_.arg(current_arg_++, typename viennacl::result_of::cl_type<NumericT>::type(mat.value()));
}
void operator()(scheduler::statement const & statement, vcl_size_t root_idx, leaf_t leaf_t) const
{
scheduler::statement_node const & root_node = statement.array()[root_idx];
if (leaf_t==LHS_NODE_TYPE && root_node.lhs.type_family != scheduler::COMPOSITE_OPERATION_FAMILY)
utils::call_on_element(root_node.lhs, *this);
else if (leaf_t==RHS_NODE_TYPE && root_node.rhs.type_family != scheduler::COMPOSITE_OPERATION_FAMILY)
utils::call_on_element(root_node.rhs, *this);
}
private:
symbolic_binder & binder_;
unsigned int & current_arg_;
viennacl::ocl::kernel & kernel_;
};
protected:
static void generate_prototype(utils::kernel_generation_stream & stream, std::string const & name, std::string const & first_arguments, std::vector<mapping_type> const & mappings, statements_container const &statements,
std::map<std::string, unsigned int> const & widths)
{
statements_container::data_type::const_iterator sit;
std::vector<mapping_type>::const_iterator mit;
std::set<std::string> already_generated;
std::string arguments = first_arguments;
for (mit = mappings.begin(), sit = statements.data().begin(); sit != statements.data().end(); ++sit, ++mit)
tree_parsing::traverse(*sit, sit->root(), prototype_generation_traversal(already_generated, arguments, *mit, widths), true);
arguments.erase(arguments.size()-1); //Last comma pruned
stream << "__kernel " << "void " << name << "(" << arguments << ")" << std::endl;
}
static void generate_prototype(utils::kernel_generation_stream & stream, std::string const & name, std::string const & first_arguments, std::vector<mapping_type> const & mappings, statements_container const & statements)
{
generate_prototype(stream, name, first_arguments, mappings, statements, std::map<std::string, unsigned int>());
}
void set_arguments(statements_container const & statements, viennacl::ocl::kernel & kernel, unsigned int & current_arg)
{
tools::shared_ptr<symbolic_binder> binder = make_binder(binding_policy_);
for (statements_container::data_type::const_iterator itt = statements.data().begin(); itt != statements.data().end(); ++itt)
tree_parsing::traverse(*itt, itt->root(), set_arguments_functor(*binder,current_arg,kernel), true);
}
class invalid_template_exception : public std::exception
{
public:
invalid_template_exception() : message_() {}
invalid_template_exception(std::string message) :
message_("ViennaCL: Internal error: The generator cannot apply the given template to the given statement: " + message + "\n"
"If you are using a builtin template, please report on viennacl-support@lists.sourceforge.net! We will provide a fix as soon as possible\n"
"If you are using your own template, please try using other parameters") {}
virtual const char* what() const throw() { return message_.c_str(); }
virtual ~invalid_template_exception() throw() {}
private:
std::string message_;
};
static void fetching_loop_info(fetching_policy_type policy, std::string const & bound, utils::kernel_generation_stream & stream, std::string & init, std::string & upper_bound, std::string & inc, std::string const & domain_id, std::string const & domain_size)
{
if (policy==FETCH_FROM_GLOBAL_STRIDED)
{
init = domain_id;
upper_bound = bound;
inc = domain_size;
}
else if (policy==FETCH_FROM_GLOBAL_CONTIGUOUS)
{
std::string chunk_size = "chunk_size";
std::string chunk_start = "chunk_start";
std::string chunk_end = "chunk_end";
stream << "unsigned int " << chunk_size << " = (" << bound << "+" << domain_size << "-1)/" << domain_size << ";" << std::endl;
stream << "unsigned int " << chunk_start << " =" << domain_id << "*" << chunk_size << ";" << std::endl;
stream << "unsigned int " << chunk_end << " = min(" << chunk_start << "+" << chunk_size << ", " << bound << ");" << std::endl;
init = chunk_start;
upper_bound = chunk_end;
inc = "1";
}
}
static bool is_node_trans(scheduler::statement::container_type const & array, vcl_size_t root_idx, leaf_t leaf_type)
{
bool res = false;
scheduler::lhs_rhs_element scheduler::statement_node::*ptr;
if (leaf_type==LHS_NODE_TYPE)
ptr = &scheduler::statement_node::lhs;
else
ptr = &scheduler::statement_node::rhs;
scheduler::statement_node const * node = &array[root_idx];
while ((node->*ptr).type_family==scheduler::COMPOSITE_OPERATION_FAMILY)
{
if (array[(node->*ptr).node_index].op.type==scheduler::OPERATION_UNARY_TRANS_TYPE)
res = !res;
node = &array[(node->*ptr).node_index];
}
return res;
}
protected:
static std::string append_simd_suffix(std::string const & str, unsigned int i)
{
assert(i < 16);
static char suffixes[] = {'0','1','2','3','4','5','6','7','8','9',
'a','b','c','d','e','f'};
return str + tools::to_string(suffixes[i]);
}
static bool is_striding_operator(scheduler::statement_node const & node)
{
return node.op.type==scheduler::OPERATION_BINARY_MATRIX_COLUMN_TYPE
|| node.op.type==scheduler::OPERATION_BINARY_MATRIX_ROW_TYPE
|| node.op.type==scheduler::OPERATION_BINARY_MATRIX_DIAG_TYPE;
}
static bool has_strided_access(statements_container const & statements)
{
for (statements_container::data_type::const_iterator it = statements.data().begin(); it != statements.data().end(); ++it)
{
//checks for vectors
std::vector<scheduler::lhs_rhs_element> vectors;
tree_parsing::traverse(*it, it->root(), tree_parsing::filter_elements(scheduler::DENSE_VECTOR_TYPE, vectors), true);
for (std::vector<scheduler::lhs_rhs_element>::iterator itt = vectors.begin(); itt != vectors.end(); ++itt)
if (utils::call_on_vector(*itt, utils::stride_fun())>1)
return true;
//checks for matrix
std::vector<scheduler::lhs_rhs_element> matrices;
tree_parsing::traverse(*it, it->root(), tree_parsing::filter_elements(scheduler::DENSE_MATRIX_TYPE, matrices), true);
for (std::vector<scheduler::lhs_rhs_element>::iterator itt = matrices.begin(); itt != matrices.end(); ++itt)
if (utils::call_on_matrix(*itt, utils::stride1_fun())>1 || utils::call_on_matrix(*itt, utils::stride2_fun())>2)
return true;
std::vector<vcl_size_t> striding_operators;
tree_parsing::traverse(*it, it->root(), tree_parsing::filter(&is_striding_operator, striding_operators), false);
if(striding_operators.size() > 0)
return true;
}
return false;
}
static vcl_size_t vector_size(scheduler::statement_node const & node, bool up_to_internal_size)
{
using namespace scheduler;
using namespace utils;
if (node.op.type==OPERATION_BINARY_MATRIX_DIAG_TYPE)
{
vcl_size_t size1 = up_to_internal_size?call_on_matrix(node.lhs, internal_size1_fun()):call_on_matrix(node.lhs, size1_fun());
vcl_size_t size2 = up_to_internal_size?call_on_matrix(node.lhs, internal_size2_fun()):call_on_matrix(node.lhs, size2_fun());
return std::min<vcl_size_t>(size1, size2);
}
else if (node.op.type==OPERATION_BINARY_MATRIX_ROW_TYPE)
return up_to_internal_size?call_on_matrix(node.lhs, internal_size2_fun()):call_on_matrix(node.lhs, size2_fun());
else if (node.op.type==OPERATION_BINARY_MATRIX_COLUMN_TYPE)
return up_to_internal_size?call_on_matrix(node.lhs, internal_size1_fun()):call_on_matrix(node.lhs, size1_fun());
else
return up_to_internal_size?call_on_vector(node.lhs, internal_size_fun()):call_on_vector(node.lhs, size_fun());
}
//NB : templates are not used here because declaring a functor out of the generate() functions would be harder to read
struct loop_body_base
{
virtual void operator()(utils::kernel_generation_stream & stream, unsigned int simd_width) const = 0;
virtual ~loop_body_base() {}
};
static void element_wise_loop_1D(utils::kernel_generation_stream & stream, loop_body_base const & loop_body,
fetching_policy_type fetch, unsigned int simd_width, std::string const & i, std::string const & bound, std::string const & domain_id, std::string const & domain_size)
{
std::string strwidth = tools::to_string(simd_width);
std::string boundround = bound + "/" + strwidth;
std::string init, upper_bound, inc;
fetching_loop_info(fetch, boundround, stream, init, upper_bound, inc, domain_id, domain_size);
stream << "for(unsigned int " << i << " = " << init << "; " << i << " < " << upper_bound << "; " << i << " += " << inc << ")" << std::endl;
stream << "{" << std::endl;
stream.inc_tab();
loop_body(stream, simd_width);
stream.dec_tab();
stream << "}" << std::endl;
if (simd_width>1)
{
stream << "for(unsigned int " << i << " = " << boundround << "*" << strwidth << " + " << domain_id << "; " << i << " < " << bound << "; " << i << " += " + domain_size + ")" << std::endl;
stream << "{" << std::endl;
stream.inc_tab();
loop_body(stream, 1);
stream.dec_tab();
stream << "}" << std::endl;
}
}
static std::string vstore(unsigned int simd_width, std::string const & value, std::string const & offset, std::string const & ptr)
{
if (simd_width==1)
return "(" + ptr + ")[" + offset + "] = " + value;
else
return utils::append_width("vstore", simd_width) + "(" + value + ", " + offset + ", " + ptr + ")";
}
static std::string vload(unsigned int simd_width, std::string const & offset, std::string const & ptr)
{
if (simd_width==1)
return "(" + ptr + ")[" + offset + "]";
else
return utils::append_width("vload", simd_width) + "(" + offset + ", " + ptr + ")";
}
private:
virtual std::vector<std::string> generate_impl(std::string const & kernel_prefix, statements_container const & statements, std::vector<mapping_type> const & mapping) const = 0;
public:
template_base(binding_policy_t binding_policy) : binding_policy_(binding_policy) {}
virtual ~template_base(){ }
std::vector<std::string> generate(std::string const & kernel_prefix, statements_container const & statements, viennacl::ocl::device const & device)
{
statements_container::data_type::const_iterator sit;
std::vector<mapping_type>::iterator mit;
if(int err = check_invalid(statements, device))
throw generator_not_supported_exception("The supplied parameters for this template are invalid : err " + tools::to_string(err));
//Create mapping
std::vector<mapping_type> mappings(statements.data().size());
tools::shared_ptr<symbolic_binder> binder = make_binder(binding_policy_);
for (mit = mappings.begin(), sit = statements.data().begin(); sit != statements.data().end(); ++sit, ++mit)
tree_parsing::traverse(*sit, sit->root(), map_functor(*binder,*mit), true);
return generate_impl(kernel_prefix, statements, mappings);
}
virtual int check_invalid(statements_container const & statements, viennacl::ocl::device const & device) const = 0;
virtual void enqueue(std::string const & kernel_prefix, std::vector<lazy_program_compiler> & programs, statements_container const & statements) = 0;
virtual tools::shared_ptr<template_base> clone() const = 0;
private:
binding_policy_t binding_policy_;
};
template<class TemplateType, class ParametersType>
class template_base_impl : public template_base
{
private:
virtual int check_invalid_impl(viennacl::ocl::device const & /*dev*/) const { return TEMPLATE_VALID; }
virtual unsigned int n_lmem_elements() const { return 0; }
public:
typedef ParametersType parameters_type;
template_base_impl(parameters_type const & parameters, binding_policy_t binding_policy) : template_base(binding_policy), p_(parameters){ }
parameters_type const & parameters() const
{
return p_;
}
tools::shared_ptr<template_base> clone() const
{
return tools::shared_ptr<template_base>(new TemplateType(*dynamic_cast<TemplateType const *>(this)));
}
int check_invalid(statements_container const & statements, viennacl::ocl::device const & device) const
{
using namespace viennacl::tools;
scheduler::statement const & statement = statements.data().front();
unsigned int scalartype_size = utils::size_of(lhs_most(statement.array(), statement.root()).lhs.numeric_type);
//Query device informations
vcl_size_t lmem_available = static_cast<vcl_size_t>(device.local_mem_size());
vcl_size_t lmem_usage = scalartype_size*n_lmem_elements();
if (lmem_usage>lmem_available)
return TEMPLATE_LOCAL_MEMORY_OVERFLOW;
//Invalid work group size
vcl_size_t max_workgroup_size = device.max_work_group_size();
std::vector<vcl_size_t> max_work_item_sizes = device.max_work_item_sizes();
if (p_.local_size_0*p_.local_size_1 > max_workgroup_size)
return TEMPLATE_WORK_GROUP_SIZE_OVERFLOW;
if (p_.local_size_0 > max_work_item_sizes[0])
return TEMPLATE_LOCAL_SIZE_0_OVERFLOW;
if (p_.local_size_1 > max_work_item_sizes[1])
return TEMPLATE_LOCAL_SIZE_1_OVERFLOW;
//Advice from the Intel guide
unsigned int warp_size = 8;
if (device.type()==CL_DEVICE_TYPE_GPU)
{
//Advice from the nvidia guide
warp_size = 32;
//Advice from the AMD guide
if (device.vendor_id()==4098)
warp_size = 64;
}
if (((p_.local_size_0*p_.local_size_1)%warp_size)>0)
return TEMPLATE_LOCAL_SIZE_NOT_WARP_MULTIPLE;
//Invalid SIMD Width
if (p_.simd_width!=1 && p_.simd_width!=2 &&
p_.simd_width!=4 && p_.simd_width!=8 &&
p_.simd_width!=16)
return TEMPLATE_INVALID_SIMD_WIDTH;
return check_invalid_impl(device);
}
protected:
parameters_type p_;
};
}
}
#endif
viennacl::device_specific::template_base::parameters_type::local_size_0
unsigned int local_size_0
Definition: template_base.hpp:62
viennacl::device_specific::template_base_impl::clone
tools::shared_ptr< template_base > clone() const
Definition: template_base.hpp:537
viennacl::scheduler::OPERATION_BINARY_VECTOR_DIAG_TYPE
Definition: forwards.h:141
viennacl::device_specific::template_base::check_invalid
virtual int check_invalid(statements_container const &statements, viennacl::ocl::device const &device) const =0
returns whether or not the profile has undefined behavior on particular device
viennacl::device_specific::template_base::loop_body_base::~loop_body_base
virtual ~loop_body_base()
Definition: template_base.hpp:436
viennacl::device_specific::tree_parsing::filter_elements
Definition: tree_parsing.hpp:102
device.hpp
Represents an OpenCL device within ViennaCL.
viennacl::traits::stride1
result_of::size_type< matrix_base< NumericT > >::type stride1(matrix_base< NumericT > const &s)
Definition: stride.hpp:55
viennacl::device_specific::generator_not_supported_exception
Exception for the case the generator is unable to deal with the operation.
Definition: forwards.h:163
viennacl::device_specific::template_base::set_arguments
void set_arguments(statements_container const &statements, viennacl::ocl::kernel &kernel, unsigned int &current_arg)
Definition: template_base.hpp:309
viennacl::scheduler::OPERATION_UNARY_TRANS_TYPE
Definition: forwards.h:106
viennacl::device_specific::symbolic_binder
Definition: forwards.h:216
viennacl::ocl::kernel
Represents an OpenCL kernel within ViennaCL.
Definition: kernel.hpp:58
viennacl::device_specific::template_base::invalid_template_exception::invalid_template_exception
invalid_template_exception()
Definition: template_base.hpp:319
viennacl::device_specific::matrix_product_parameters
Definition: matrix_product_template.hpp:47
viennacl::device_specific::utils::stride2_fun
Definition: utils.hpp:307
viennacl::scheduler::DENSE_VECTOR_TYPE
Definition: forwards.h:271
viennacl::device_specific::template_base::vstore
static std::string vstore(unsigned int simd_width, std::string const &value, std::string const &offset, std::string const &ptr)
Definition: template_base.hpp:465
viennacl::device_specific::template_base_impl::check_invalid
int check_invalid(statements_container const &statements, viennacl::ocl::device const &device) const
returns whether or not the profile has undefined behavior on particular device
Definition: template_base.hpp:543
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
io.hpp
Some helper routines for reading/writing/printing scheduler expressions.
viennacl::device_specific::tree_parsing::traverse
void traverse(scheduler::statement const &statement, vcl_size_t root_idx, Fun const &fun, bool inspect)
Recursively execute a functor on a statement.
Definition: tree_parsing.hpp:53
viennacl::device_specific::template_base::is_node_trans
static bool is_node_trans(scheduler::statement::container_type const &array, vcl_size_t root_idx, leaf_t leaf_type)
Definition: template_base.hpp:353
viennacl::device_specific::template_base::invalid_template_exception
Definition: template_base.hpp:316
viennacl::device_specific::template_base_impl::parameters
parameters_type const & parameters() const
Definition: template_base.hpp:532
viennacl::device_specific::leaf_t
leaf_t
Definition: forwards.h:180
viennacl::device_specific::template_base::is_striding_operator
static bool is_striding_operator(scheduler::statement_node const &node)
Definition: template_base.hpp:381
viennacl::ocl::device
A class representing a compute device (e.g. a GPU)
Definition: device.hpp:49
viennacl::scheduler::statement_node::op
op_element op
Definition: forwards.h:481
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::device_specific::template_base
Definition: template_base.hpp:54
viennacl::ocl::device::local_mem_size
cl_ulong local_mem_size() const
Size of local memory arena in bytes. The minimum value is 32 KB.
Definition: device.hpp:360
viennacl::scheduler::OPERATION_BINARY_MATRIX_ROW_TYPE
Definition: forwards.h:142
viennacl::ocl::device::type
cl_device_type type() const
The OpenCL device type.
Definition: device.hpp:893
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::result_of::cl_type::type
T type
Definition: result_of.hpp:590
viennacl::scheduler::statement::array
container_type const & array() const
Definition: forwards.h:528
viennacl::ocl::device::vendor_id
cl_uint vendor_id() const
A unique device vendor identifier. An example of a unique device identifier could be the PCIe ID...
Definition: device.hpp:917
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::statements_container::data
std::list< scheduler::statement > const & data() const
Definition: forwards.h:282
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::device_specific::template_base::append_simd_suffix
static std::string append_simd_suffix(std::string const &str, unsigned int i)
Definition: template_base.hpp:373
viennacl::device_specific::PARENT_NODE_TYPE
Definition: forwards.h:183
viennacl::device_specific::template_base::parameters_type::num_kernels
unsigned int num_kernels
Definition: template_base.hpp:64
NumericT
float NumericT
Definition: bisect.cpp:40
viennacl::scheduler::COMPOSITE_OPERATION_FAMILY
Definition: forwards.h:251
viennacl::device_specific::template_base::clone
virtual tools::shared_ptr< template_base > clone() const =0
viennacl::device_specific::lhs_most
scheduler::statement_node const & lhs_most(scheduler::statement::container_type const &array, vcl_size_t root)
Definition: forwards.h:87
viennacl::device_specific::is_scalar_reduction
bool is_scalar_reduction(scheduler::statement_node const &node)
Definition: forwards.h:75
viennacl::scheduler::statement::container_type
std::vector< value_type > container_type
Definition: forwards.h:507
viennacl::device_specific::FETCH_FROM_GLOBAL_STRIDED
Definition: template_base.hpp:50
viennacl::scheduler::DENSE_MATRIX_TYPE
Definition: forwards.h:274
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::traits::ld
vcl_size_t ld(matrix_base< NumericT > const &mat)
Helper routine for obtaining the internal number of entries per row of a ViennaCL matrix...
Definition: size.hpp:398
viennacl::device_specific::template_base::parameters_type::local_size_1
unsigned int local_size_1
Definition: template_base.hpp:63
mapped_objects.hpp
Map ViennaCL objects to generator wrappers.
viennacl::device_specific::template_base::has_strided_access
static bool has_strided_access(statements_container const &statements)
Definition: template_base.hpp:388
viennacl::device_specific::template_base::fetching_loop_info
static void fetching_loop_info(fetching_policy_type policy, std::string const &bound, utils::kernel_generation_stream &stream, std::string &init, std::string &upper_bound, std::string &inc, std::string const &domain_id, std::string const &domain_size)
Definition: template_base.hpp:330
viennacl::scheduler::INVALID_NUMERIC_TYPE
Definition: forwards.h:288
viennacl::device_specific::RHS_NODE_TYPE
Definition: forwards.h:184
viennacl::device_specific::LHS_NODE_TYPE
Definition: forwards.h:182
lazy_program_compiler.hpp
Helper for compiling a program lazily.
viennacl::scheduler::lhs_rhs_element::numeric_type
statement_node_numeric_type numeric_type
Definition: forwards.h:341
viennacl::traits::start
result_of::size_type< T >::type start(T const &obj)
Definition: start.hpp:44
viennacl::scheduler::statement_node::rhs
lhs_rhs_element rhs
Definition: forwards.h:482
device_utils.hpp
Various utility implementations for dispatching with respect to the different devices available on th...
viennacl::device_specific::template_base::invalid_template_exception::invalid_template_exception
invalid_template_exception(std::string message)
Definition: template_base.hpp:320
viennacl::device_specific::utils::kernel_generation_stream::dec_tab
void dec_tab()
Definition: utils.hpp:437
viennacl::tools::shared_ptr
A shared pointer class similar to boost::shared_ptr. Reimplemented in order to avoid a Boost-dependen...
Definition: shared_ptr.hpp:83
viennacl::device_specific::template_base::loop_body_base
Definition: template_base.hpp:433
viennacl::scheduler::statement_node::lhs
lhs_rhs_element lhs
Definition: forwards.h:480
viennacl::vcl_size_t
std::size_t vcl_size_t
Definition: forwards.h:75
viennacl::device_specific::template_base::invalid_template_exception::~invalid_template_exception
virtual ~invalid_template_exception()
Definition: template_base.hpp:325
viennacl::device_specific::tree_parsing::traversal_functor
base functor class for traversing a statement
Definition: tree_parsing.hpp:44
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::device_specific::template_base::vector_size
static vcl_size_t vector_size(scheduler::statement_node const &node, bool up_to_internal_size)
Definition: template_base.hpp:414
viennacl::device_specific::utils::kernel_generation_stream::inc_tab
void inc_tab()
Definition: utils.hpp:436
viennacl::device_specific::template_base::generate_prototype
static void generate_prototype(utils::kernel_generation_stream &stream, std::string const &name, std::string const &first_arguments, std::vector< mapping_type > const &mappings, statements_container const &statements, std::map< std::string, unsigned int > const &widths)
Definition: template_base.hpp:290
viennacl::device_specific::template_base::invalid_template_exception::what
virtual const char * what() const
Definition: template_base.hpp:324
viennacl::traits::stride2
result_of::size_type< matrix_base< NumericT > >::type stride2(matrix_base< NumericT > const &s)
Definition: stride.hpp:65
viennacl::device_specific::template_base::generate_prototype
static void generate_prototype(utils::kernel_generation_stream &stream, std::string const &name, std::string const &first_arguments, std::vector< mapping_type > const &mappings, statements_container const &statements)
Definition: template_base.hpp:304
viennacl::device_specific::mapping_type
std::map< mapping_key, tools::shared_ptr< mapped_object > > mapping_type
Definition: forwards.h:191
viennacl::device_specific::template_base::enqueue
virtual void enqueue(std::string const &kernel_prefix, std::vector< lazy_program_compiler > &programs, statements_container const &statements)=0
viennacl::device_specific::tree_parsing::filter
Definition: tree_parsing.hpp:84
viennacl::scheduler::OPERATION_BINARY_MATRIX_COLUMN_TYPE
Definition: forwards.h:143
viennacl::device_specific::fetching_policy_type
fetching_policy_type
Definition: template_base.hpp:47
tree_parsing.hpp
Code for parsing the expression trees.
viennacl::device_specific::template_base::generate
std::vector< std::string > generate(std::string const &kernel_prefix, statements_container const &statements, viennacl::ocl::device const &device)
Definition: template_base.hpp:490
kernel.hpp
Representation of an OpenCL kernel in ViennaCL.
utils.hpp
Internal utils.
viennacl::device_specific::template_base_impl::template_base_impl
template_base_impl(parameters_type const &parameters, binding_policy_t binding_policy)
The constructor.
Definition: template_base.hpp:530
viennacl::device_specific::template_base_impl
Definition: template_base.hpp:519
viennacl::traits::row_major
bool row_major(T const &)
Definition: row_major.hpp:38
viennacl::device_specific::FETCH_FROM_LOCAL
Definition: template_base.hpp:49
viennacl::device_specific::binding_policy_t
binding_policy_t
Definition: forwards.h:246
viennacl::scheduler::op_element::type
operation_node_type type
Definition: forwards.h:474
viennacl::device_specific::utils::stride1_fun
Definition: utils.hpp:301
viennacl::device_specific::template_base::element_wise_loop_1D
static void element_wise_loop_1D(utils::kernel_generation_stream &stream, loop_body_base const &loop_body, fetching_policy_type fetch, unsigned int simd_width, std::string const &i, std::string const &bound, std::string const &domain_id, std::string const &domain_size)
Definition: template_base.hpp:439
viennacl::device_specific::template_base::vload
static std::string vload(unsigned int simd_width, std::string const &offset, std::string const &ptr)
Definition: template_base.hpp:473
viennacl::device_specific::template_base::~template_base
virtual ~template_base()
Definition: template_base.hpp:488
viennacl::scheduler::statement
The main class for representing a statement such as x = inner_prod(y,z); at runtime.
Definition: forwards.h:502
viennacl::device_specific::template_base::parameters_type::simd_width
unsigned int simd_width
Definition: template_base.hpp:61
viennacl::scheduler::lhs_rhs_element::node_index
vcl_size_t node_index
Definition: forwards.h:346
viennacl::traits::handle
viennacl::backend::mem_handle & handle(T &obj)
Returns the generic memory handle of an object. Non-const version.
Definition: handle.hpp:41
viennacl::device_specific::is_vector_reduction
bool is_vector_reduction(scheduler::statement_node const &node)
Definition: forwards.h:80
viennacl::device_specific::at
ValueT const & at(std::map< KeyT, ValueT > const &map, KeyT const &key)
Emulation of C++11's .at() member for std::map<>, const-version.
Definition: forwards.h:142
viennacl::device_specific::template_base::template_base
template_base(binding_policy_t binding_policy)
Definition: template_base.hpp:486
viennacl::tools::to_string
std::string to_string(T const t)
Definition: tools.hpp:304
viennacl::device_specific::FETCH_FROM_GLOBAL_CONTIGUOUS
Definition: template_base.hpp:51
viennacl::device_specific::template_base::parameters_type
Definition: template_base.hpp:57
viennacl::scheduler::statement_node
Main datastructure for an node in the statement tree.
Definition: forwards.h:478
viennacl::ocl::device::max_work_group_size
size_t max_work_group_size() const
Maximum number of work-items in a work-group executing a kernel using the data parallel execution mod...
Definition: device.hpp:483
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::statements_container
Definition: forwards.h:262
viennacl::device_specific::make_binder
tools::shared_ptr< symbolic_binder > make_binder(binding_policy_t policy)
Definition: forwards.h:251
viennacl::device_specific::template_base::parameters_type::parameters_type
parameters_type(unsigned int _simd_width, unsigned int _local_size_1, unsigned int _local_size_2, unsigned int _num_kernels)
Definition: template_base.hpp:59
viennacl::ocl::device::max_work_item_sizes
std::vector< size_t > max_work_item_sizes() const
Maximum number of work-items that can be specified in each dimension of the work-group.
Definition: device.hpp:510
viennacl::device_specific::utils::append_width
std::string append_width(std::string const &str, unsigned int width)
Definition: utils.hpp:558
viennacl::device_specific::template_base_impl::parameters_type
ParametersType parameters_type
Definition: template_base.hpp:527
viennacl::device_specific::template_base_impl::p_
parameters_type p_
Definition: template_base.hpp:590