gaspersic/triton
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Updated the file hierarchy

Browse Source
This commit is contained in:
Philippe Tillet
2014-09-01 19:37:05 -04:00
parent af7c3e3a75
commit e7ff3ba3c0
14 changed files with 207 additions and 178 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
atidlas/execute.hpp
View File

@@ -11,8 +11,8 @@
#include "atidlas/forwards.h"
#include "atidlas/templates/template_base.hpp"
#include "atidlas/tree_parsing.hpp"
#include "atidlas/execution_handler.hpp"
#include "atidlas/tools/tree_parsing.hpp"
#include "atidlas/tools/execution_handler.hpp"
namespace atidlas
{
@@ -20,7 +20,7 @@ namespace atidlas
inline void execute(template_base const & T, statements_container const & statements, viennacl::ocl::context & ctx = viennacl::ocl::current_context(), bool force_compilation = false)
{
//Generate program name
std::string program_name = tree_parsing::statements_representation(statements, BIND_TO_HANDLE);
std::string program_name = tools::statements_representation(statements, BIND_TO_HANDLE);
execution_handler handler(program_name, ctx, ctx.current_device(), force_compilation);
handler.add(program_name, T, statements);
handler.execute(program_name, statements);

6
atidlas/forwards.h
View File

@@ -137,7 +137,7 @@ private:
std::string message_;
};
namespace utils
namespace tools
{
class kernel_generation_stream;
}
@@ -157,13 +157,13 @@ typedef std::pair<atidlas_int_t, leaf_t> mapping_key;
typedef std::map<mapping_key, tools::shared_ptr<mapped_object> > mapping_type;
namespace tree_parsing
namespace tools
{
template<class Fun>
inline void traverse(viennacl::scheduler::statement const & statement, atidlas_int_t root_idx, Fun const & fun, bool inspect);
inline void process(utils::kernel_generation_stream & stream, leaf_t leaf, std::multimap<std::string, std::string> const & accessors,
inline void process(tools::kernel_generation_stream & stream, leaf_t leaf, std::multimap<std::string, std::string> const & accessors,
viennacl::scheduler::statement const & statement, size_t root_idx, mapping_type const & mapping, std::set<std::string> & already_processed);
inline std::string evaluate(leaf_t leaf, std::map<std::string, std::string> const & accessors, viennacl::scheduler::statement const & statement, atidlas_int_t root_idx,mapping_type const & mapping);
}

28
atidlas/mapped_objects.hpp
View File

@@ -7,7 +7,7 @@
#include "atidlas/forwards.h"
#include "atidlas/tools/find_and_replace.hpp"
#include "atidlas/utils.hpp"
#include "atidlas/tools/misc.hpp"
namespace atidlas
{
@@ -118,15 +118,15 @@ class binary_leaf
public:
binary_leaf(mapped_object::node_info info) : info_(info){ }
void process_recursive(utils::kernel_generation_stream & stream, leaf_t leaf, std::multimap<std::string, std::string> const & accessors)
void process_recursive(tools::kernel_generation_stream & stream, leaf_t leaf, std::multimap<std::string, std::string> const & accessors)
{
std::set<std::string> already_fetched;
tree_parsing::process(stream, leaf, accessors, *info_.statement, info_.root_idx, *info_.mapping, already_fetched);
tools::process(stream, leaf, accessors, *info_.statement, info_.root_idx, *info_.mapping, already_fetched);
}
std::string evaluate_recursive(leaf_t leaf, std::map<std::string, std::string> const & accessors)
{
return tree_parsing::evaluate(leaf, accessors, *info_.statement, info_.root_idx, *info_.mapping);
return tools::evaluate(leaf, accessors, *info_.statement, info_.root_idx, *info_.mapping);
}
protected:
@@ -155,7 +155,7 @@ public:
atidlas_int_t root_idx() const { return info_.root_idx; }
viennacl::scheduler::statement const & statement() const { return *info_.statement; }
viennacl::scheduler::statement_node root_node() const { return statement().array()[root_idx()]; }
bool is_index_reduction() const { return utils::is_index_reduction(info_.statement->array()[info_.root_idx].op); }
bool is_index_reduction() const { return tools::is_index_reduction(info_.statement->array()[info_.root_idx].op); }
viennacl::scheduler::op_element root_op() const
{
@@ -302,9 +302,9 @@ private:
void postprocess(std::string &res) const
{
std::map<std::string, std::string> accessors;
tools::find_and_replace(res, "#diag_offset", tree_parsing::evaluate(RHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping));
tools::find_and_replace(res, "#diag_offset", tools::evaluate(RHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping));
accessors["vector"] = res;
res = tree_parsing::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
res = tools::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
}
public:
@@ -323,7 +323,7 @@ private:
{
std::map<std::string, std::string> accessors;
accessors["matrix"] = res;
res = tree_parsing::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
res = tools::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
}
public:
@@ -340,9 +340,9 @@ private:
void postprocess(std::string &res) const
{
std::map<std::string, std::string> accessors;
tools::find_and_replace(res, "#row", tree_parsing::evaluate(RHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping));
tools::find_and_replace(res, "#row", tools::evaluate(RHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping));
accessors["matrix"] = res;
res = tree_parsing::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
res = tools::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
}
public:
@@ -361,9 +361,9 @@ private:
void postprocess(std::string &res) const
{
std::map<std::string, std::string> accessors;
tools::find_and_replace(res, "#column", tree_parsing::evaluate(RHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping));
tools::find_and_replace(res, "#column", tools::evaluate(RHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping));
accessors["matrix"] = res;
res = tree_parsing::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
res = tools::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
}
public:
@@ -381,9 +381,9 @@ private:
void postprocess(std::string &res) const
{
std::map<std::string, std::string> accessors;
tools::find_and_replace(res, "#diag_offset", tree_parsing::evaluate(RHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping));
tools::find_and_replace(res, "#diag_offset", tools::evaluate(RHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping));
accessors["matrix"] = res;
res = tree_parsing::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
res = tools::evaluate(LHS_NODE_TYPE, accessors, *info_.statement, info_.root_idx, *info_.mapping);
}
public:

21
atidlas/templates/matrix_axpy.hpp
View File

@@ -4,9 +4,6 @@
#include <vector>
#include "atidlas/mapped_objects.hpp"
#include "atidlas/tree_parsing.hpp"
#include "atidlas/utils.hpp"
#include "atidlas/templates/template_base.hpp"
#include "viennacl/scheduler/forwards.h"
@@ -40,7 +37,7 @@ private:
std::string generate_impl(std::string const & kernel_prefix, statements_container const & statements, std::vector<mapping_type> const & mappings, unsigned int simd_width) const
{
utils::kernel_generation_stream stream;
tools::kernel_generation_stream stream;
std::string init0, upper_bound0, inc0, init1, upper_bound1, inc1;
@@ -49,7 +46,7 @@ private:
stream << "{" << std::endl;
stream.inc_tab();
tree_parsing::process(stream, PARENT_NODE_TYPE, utils::create_process_accessors("scalar", "#scalartype #namereg = *#pointer;")
tools::process(stream, PARENT_NODE_TYPE, tools::create_process_accessors("scalar", "#scalartype #namereg = *#pointer;")
("matrix", "#pointer += $OFFSET{#start1, #start2};")
("vector", "#pointer += #start;"), statements, mappings);
@@ -62,16 +59,16 @@ private:
stream << "{" << std::endl;
stream.inc_tab();
tree_parsing::process(stream, PARENT_NODE_TYPE, utils::create_process_accessors("matrix", utils::append_width("#scalartype",simd_width) + " #namereg = #pointer[$OFFSET{i*#stride1,j*#stride2}];")
tools::process(stream, PARENT_NODE_TYPE, tools::create_process_accessors("matrix", tools::append_width("#scalartype",simd_width) + " #namereg = #pointer[$OFFSET{i*#stride1,j*#stride2}];")
("vector_diag", "#scalartype #namereg = ((i + ((#diag_offset<0)?#diag_offset:0))!=(j-((#diag_offset>0)?#diag_offset:0)))?0:#pointer[min(i*#stride, j*#stride)];")
, statements, mappings);
tree_parsing::evaluate(stream, PARENT_NODE_TYPE, utils::create_evaluate_accessors("matrix", "#namereg")
tools::evaluate(stream, PARENT_NODE_TYPE, tools::create_evaluate_accessors("matrix", "#namereg")
("vector_diag", "#namereg")
("scalar", "#namereg")
, statements, mappings);
tree_parsing::process(stream, LHS_NODE_TYPE, utils::create_process_accessors("matrix", "#pointer[$OFFSET{i*#stride1,j*#stride2}] = #namereg;")
tools::process(stream, LHS_NODE_TYPE, tools::create_process_accessors("matrix", "#pointer[$OFFSET{i*#stride1,j*#stride2}] = #namereg;")
, statements, mappings);
stream.dec_tab();
@@ -113,13 +110,13 @@ public:
unsigned int current_arg = 0;
if (up_to_internal_size_)
{
kernel.arg(current_arg++, cl_uint(utils::call_on_matrix(root.lhs, utils::internal_size1_fun())));
kernel.arg(current_arg++, cl_uint(utils::call_on_matrix(root.lhs, utils::internal_size2_fun())));
kernel.arg(current_arg++, cl_uint(tools::call_on_matrix(root.lhs, tools::internal_size1_fun())));
kernel.arg(current_arg++, cl_uint(tools::call_on_matrix(root.lhs, tools::internal_size2_fun())));
}
else
{
kernel.arg(current_arg++, cl_uint(utils::call_on_matrix(root.lhs, utils::size1_fun())));
kernel.arg(current_arg++, cl_uint(utils::call_on_matrix(root.lhs, utils::size2_fun())));
kernel.arg(current_arg++, cl_uint(tools::call_on_matrix(root.lhs, tools::size1_fun())));
kernel.arg(current_arg++, cl_uint(tools::call_on_matrix(root.lhs, tools::size2_fun())));
}
set_arguments(statements, kernel, current_arg);

33
atidlas/templates/matrix_product.hpp
View File

@@ -11,9 +11,6 @@
#include "viennacl/tools/tools.hpp"
#include "atidlas/templates/template_base.hpp"
#include "atidlas/mapped_objects.hpp"
#include "atidlas/utils.hpp"
#include "atidlas/tree_parsing.hpp"
#include "atidlas/tools/align.hpp"
namespace atidlas
@@ -113,7 +110,7 @@ private:
atidlas_int_t & A_idx, leaf_t & A_leaf, bool& A_trans, atidlas_int_t & B_idx, leaf_t & B_leaf, bool& B_trans,
atidlas_int_t & beta_idx, leaf_t & beta_leaf)
{
using namespace tree_parsing;
using namespace tools;
using namespace viennacl::scheduler;
viennacl::scheduler::statement::container_type const & array = s.array();
@@ -159,7 +156,7 @@ private:
beta_leaf = RHS_NODE_TYPE;
}
void handle_bounds(bool fallback, utils::kernel_generation_stream & stream, std::string const & inbounds, std::string const & do_if, std::string do_else) const
void handle_bounds(bool fallback, tools::kernel_generation_stream & stream, std::string const & inbounds, std::string const & do_if, std::string do_else) const
{
if (fallback)
{
@@ -195,7 +192,7 @@ private:
//////////////////
/// INIT
/// //////////////
utils::kernel_generation_stream stream;
tools::kernel_generation_stream stream;
viennacl::scheduler::statement const & st = statements.data().front();
mapping_type const & mapping = mappings.front();
@@ -214,8 +211,8 @@ private:
/// DECLARATIONS
/// //////////////
stream << " __attribute__((reqd_work_group_size(" << p.local_size_0 << "," << p.local_size_1 << ",1)))" << std::endl;
stream << "__kernel void " << kernel_prefix << "(unsigned int M, unsigned int N, unsigned int K, " << generate_arguments(mappings, utils::create_process_accessors(A->name(), matrix_arguments(utils::append_width("#scalartype", p.simd_width)))
(B->name(), matrix_arguments(utils::append_width("#scalartype", p.simd_width)))
stream << "__kernel void " << kernel_prefix << "(unsigned int M, unsigned int N, unsigned int K, " << generate_arguments(mappings, tools::create_process_accessors(A->name(), matrix_arguments(tools::append_width("#scalartype", p.simd_width)))
(B->name(), matrix_arguments(tools::append_width("#scalartype", p.simd_width)))
("matrix", matrix_arguments("#scalartype"))
("host_scalar", "#scalartype #name,"),statements) << ")" << std::endl;
stream << "{" << std::endl;
@@ -227,7 +224,7 @@ private:
stream << B->process("#start1/= " + to_string(p.simd_width) + ";") << std::endl;
stream << B->process("#ld /= " + to_string(p.simd_width) + ";") << std::endl;
}
tree_parsing::process(stream, PARENT_NODE_TYPE, utils::create_process_accessors("matrix", "#pointer += $OFFSET{#start1, #start2};")
tools::process(stream, PARENT_NODE_TYPE, tools::create_process_accessors("matrix", "#pointer += $OFFSET{#start1, #start2};")
("matrix", "#ld *= #nldstride;"), statements, mappings);
///Result Values
@@ -235,11 +232,11 @@ private:
if (p.A_fetching_policy==FETCH_FROM_LOCAL)
stream << A->process("#scalartype rA[" + to_string(p.kS) + "][" + to_string(p.mS) + "];") << std::endl;
else
stream << A->process(utils::append_width("#scalartype",p.simd_width) + " rA[" + to_string(p.kS) + "][" + to_string(p.mS/p.simd_width) + "];") << std::endl;
stream << A->process(tools::append_width("#scalartype",p.simd_width) + " rA[" + to_string(p.kS) + "][" + to_string(p.mS/p.simd_width) + "];") << std::endl;
if (p.B_fetching_policy==FETCH_FROM_LOCAL)
stream << B->process("#scalartype rB[" + to_string(p.kS) + "][" + to_string(p.nS) + "];");
else
stream << B->process(utils::append_width("#scalartype",p.simd_width) + " rB[" + to_string(p.kS) + "][" + to_string(p.nS/p.simd_width) + "];") << std::endl;
stream << B->process(tools::append_width("#scalartype",p.simd_width) + " rB[" + to_string(p.kS) + "][" + to_string(p.nS/p.simd_width) + "];") << std::endl;
if (p.A_fetching_policy==FETCH_FROM_LOCAL)
@@ -711,7 +708,7 @@ private:
viennacl::scheduler::statement & statement, viennacl::scheduler::lhs_rhs_element & A, viennacl::scheduler::lhs_rhs_element & B, viennacl::scheduler::lhs_rhs_element & C, viennacl::scheduler::lhs_rhs_element & beta,
NumericT beta_value, std::vector<lazy_program_compiler> & programs, std::string const & kernel_prefix)
{
using namespace utils;
using namespace tools;
atidlas_int_t ldstrideA = call_on_matrix(A, leading_stride_fun());
atidlas_int_t ldstrideB = call_on_matrix(B, leading_stride_fun());
atidlas_int_t ldstrideC = call_on_matrix(C, leading_stride_fun());
@@ -762,8 +759,8 @@ public:
virtual void enqueue(std::string const & kernel_prefix, std::vector<lazy_program_compiler> & programs, statements_container const & statements)
{
using namespace utils;
using namespace tree_parsing;
using namespace tools;
using namespace tools;
viennacl::scheduler::statement const & st = statements.data().front();
bool A_trans, B_trans;
@@ -772,10 +769,10 @@ public:
parse(st, C_idx, C_leaf, alpha_idx, alpha_leaf, A_idx, A_leaf, A_trans, B_idx, B_leaf, B_trans, beta_idx, beta_leaf);
viennacl::scheduler::statement stcopy = st;
viennacl::scheduler::lhs_rhs_element& A = utils::lhs_rhs_element(stcopy, A_idx, A_leaf);
viennacl::scheduler::lhs_rhs_element& B = utils::lhs_rhs_element(stcopy, B_idx, B_leaf);
viennacl::scheduler::lhs_rhs_element& C = utils::lhs_rhs_element(stcopy, C_idx, C_leaf);
viennacl::scheduler::lhs_rhs_element& beta = utils::lhs_rhs_element(stcopy, beta_idx, beta_leaf);
viennacl::scheduler::lhs_rhs_element& A = tools::lhs_rhs_element(stcopy, A_idx, A_leaf);
viennacl::scheduler::lhs_rhs_element& B = tools::lhs_rhs_element(stcopy, B_idx, B_leaf);
viennacl::scheduler::lhs_rhs_element& C = tools::lhs_rhs_element(stcopy, C_idx, C_leaf);
viennacl::scheduler::lhs_rhs_element& beta = tools::lhs_rhs_element(stcopy, beta_idx, beta_leaf);
if (C.numeric_type==viennacl::scheduler::FLOAT_TYPE)
enqueue_impl<float>(&viennacl::scheduler::lhs_rhs_element::matrix_float, stcopy, A, B, C, beta, beta.host_float, programs, kernel_prefix);

22
atidlas/templates/reduction.hpp
View File

@@ -8,8 +8,6 @@
#include "viennacl/scheduler/forwards.h"
#include "viennacl/tools/tools.hpp"
#include "atidlas/tree_parsing.hpp"
#include "atidlas/utils.hpp"
#include "atidlas/templates/template_base.hpp"
namespace atidlas
@@ -41,7 +39,7 @@ private:
return TEMPLATE_VALID;
}
inline void reduce_1d_local_memory(utils::kernel_generation_stream & stream, unsigned int size, std::vector<mapped_scalar_reduction*> exprs,
inline void reduce_1d_local_memory(tools::kernel_generation_stream & stream, unsigned int size, std::vector<mapped_scalar_reduction*> exprs,
std::string const & buf_str, std::string const & buf_value_str) const
{
stream << "#pragma unroll" << std::endl;
@@ -68,7 +66,7 @@ private:
std::string generate_impl(std::string const & kernel_prefix, statements_container const & statements, std::vector<mapping_type> const & mappings, unsigned int simd_width) const
{
utils::kernel_generation_stream stream;
tools::kernel_generation_stream stream;
std::vector<mapped_scalar_reduction*> exprs;
for (std::vector<mapping_type>::const_iterator it = mappings.begin(); it != mappings.end(); ++it)
@@ -80,7 +78,7 @@ private:
std::string arguments = generate_value_kernel_argument("unsigned int", "N");
for (unsigned int k = 0; k < N; ++k)
{
std::string numeric_type = utils::numeric_type_to_string(lhs_most(exprs[k]->statement().array(),
std::string numeric_type = tools::numeric_type_to_string(lhs_most(exprs[k]->statement().array(),
exprs[k]->statement().root()).lhs.numeric_type);
if (exprs[k]->is_index_reduction())
{
@@ -101,7 +99,7 @@ private:
stream.inc_tab();
stream << "unsigned int lid = get_local_id(0);" << std::endl;
tree_parsing::process(stream, PARENT_NODE_TYPE, utils::create_process_accessors("scalar", "#scalartype #namereg = *#pointer;")
tools::process(stream, PARENT_NODE_TYPE, tools::create_process_accessors("scalar", "#scalartype #namereg = *#pointer;")
("vector", "#pointer += #start;"), statements, mappings);
for (unsigned int k = 0; k < N; ++k)
@@ -125,12 +123,12 @@ private:
public:
loop_body(std::vector<mapped_scalar_reduction*> const & _exprs) : exprs(_exprs){ }
void operator()(utils::kernel_generation_stream & stream, unsigned int simd_width) const
void operator()(tools::kernel_generation_stream & stream, unsigned int simd_width) const
{
std::string i = (simd_width==1)?"i*#stride":"i";
//Fetch vector entry
for (std::vector<mapped_scalar_reduction*>::const_iterator it = exprs.begin(); it != exprs.end(); ++it)
(*it)->process_recursive(stream, PARENT_NODE_TYPE, utils::create_process_accessors("vector", utils::append_width("#scalartype",simd_width) + " #namereg = " + vload(simd_width,i,"#pointer")+";")
(*it)->process_recursive(stream, PARENT_NODE_TYPE, tools::create_process_accessors("vector", tools::append_width("#scalartype",simd_width) + " #namereg = " + vload(simd_width,i,"#pointer")+";")
("matrix_row", "#scalartype #namereg = #pointer[$OFFSET{#row*#stride1, i*#stride2}];")
("matrix_column", "#scalartype #namereg = #pointer[$OFFSET{i*#stride1,#column*#stride2}];")
("matrix_diag", "#scalartype #namereg = #pointer[#diag_offset<0?$OFFSET{(i - #diag_offset)*#stride1, i*#stride2}:$OFFSET{i*#stride1, (i + #diag_offset)*#stride2}];"));
@@ -256,7 +254,7 @@ private:
accessors["scalar_reduction"] = "#name_buf[0]";
accessors["scalar"] = "*#pointer";
accessors["vector"] = "#pointer[#start]";
tree_parsing::evaluate(stream, PARENT_NODE_TYPE, accessors, statements, mappings);
tools::evaluate(stream, PARENT_NODE_TYPE, accessors, statements, mappings);
stream.dec_tab();
stream << "}" << std::endl;
@@ -282,14 +280,14 @@ public:
cl_uint size = 0;
for (statements_container::data_type::const_iterator it = statements.data().begin(); it != statements.data().end(); ++it)
{
std::vector<size_t> reductions_idx = tree_parsing::filter_nodes(&utils::is_reduction, *it, false);
std::vector<size_t> reductions_idx = tools::filter_nodes(&tools::is_reduction, *it, false);
size = static_cast<cl_uint>(vector_size(lhs_most(it->array(), reductions_idx[0]), false));
for (std::vector<size_t>::iterator itt = reductions_idx.begin(); itt != reductions_idx.end(); ++itt)
reductions.push_back(&it->array()[*itt]);
}
viennacl::scheduler::statement const & statement = statements.data().front();
unsigned int scalartype_size = utils::size_of(lhs_most(statement.array(), statement.root()).lhs.numeric_type);
unsigned int scalartype_size = tools::size_of(lhs_most(statement.array(), statement.root()).lhs.numeric_type);
viennacl::ocl::kernel * kernels[2];
if (has_strided_access(statements) && p_.simd_width > 1)
@@ -317,7 +315,7 @@ public:
unsigned int j = 0;
for (std::vector<viennacl::scheduler::statement_node const *>::const_iterator it = reductions.begin(); it != reductions.end(); ++it)
{
if (utils::is_index_reduction((*it)->op))
if (tools::is_index_reduction((*it)->op))
{
if (tmpidx_.size() <= j)
tmpidx_.push_back(kernels[k]->context().create_memory(CL_MEM_READ_WRITE, p_.num_groups*4));

25
atidlas/templates/row_wise_reduction.hpp
View File

@@ -8,9 +8,6 @@
#include "viennacl/tools/tools.hpp"
#include "viennacl/scheduler/io.hpp"
#include "atidlas/mapped_objects.hpp"
#include "atidlas/tree_parsing.hpp"
#include "atidlas/utils.hpp"
#include "atidlas/templates/template_base.hpp"
namespace atidlas
@@ -44,7 +41,7 @@ private:
static void parse(viennacl::scheduler::statement const & statement, std::vector<size_t> & idx, bool & is_trans, viennacl::scheduler::lhs_rhs_element & matrix)
{
idx = tree_parsing::filter_nodes(&utils::is_reduction, statement, false);
idx = tools::filter_nodes(&tools::is_reduction, statement, false);
is_trans = is_node_trans(statement.array(), idx[0], LHS_NODE_TYPE);
matrix = lhs_most(statement.array(), idx[0]).lhs;
}
@@ -57,15 +54,15 @@ private:
unsigned int lsize1 = p_.local_size_1+1;
std::string lsize1str = to_string(lsize1);
utils::kernel_generation_stream stream;
tools::kernel_generation_stream stream;
stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl;
stream << "__kernel void " << kernel_prefix << "(unsigned int M, unsigned int N, " << generate_arguments("#scalartype", mappings, statements) << ")" << std::endl;
stream << "{" << std::endl;
stream.inc_tab();
tree_parsing::process(stream, PARENT_NODE_TYPE,
utils::create_process_accessors("scalar", "#scalartype #namereg = *#pointer;")
tools::process(stream, PARENT_NODE_TYPE,
tools::create_process_accessors("scalar", "#scalartype #namereg = *#pointer;")
("matrix", "#pointer += #start1 + #start2*#ld;")
("matrix", "#ld *= #nldstride;")
("vector", "#pointer += #start;"), statements, mappings);
@@ -91,9 +88,9 @@ private:
public:
loop_body(std::vector<mapped_row_wise_reduction*> const & _exprs, bool _is_trans) : exprs(_exprs), is_trans(_is_trans){ }
void operator()(utils::kernel_generation_stream & stream, unsigned int simd_width) const
void operator()(tools::kernel_generation_stream & stream, unsigned int simd_width) const
{
std::string data_type = utils::append_width("#scalartype",simd_width);
std::string data_type = tools::append_width("#scalartype",simd_width);
for (std::vector<mapped_row_wise_reduction*>::const_iterator it = exprs.begin(); it != exprs.end(); ++it)
{
@@ -181,7 +178,7 @@ private:
std::map<std::string, std::string> accessors;
accessors["row_wise_reduction"] = "#name_buf[lid0*" + lsize1str + "]";
accessors["vector"] = "#pointer[r*#stride]";
tree_parsing::evaluate(stream, PARENT_NODE_TYPE, accessors, statements, mappings);
tools::evaluate(stream, PARENT_NODE_TYPE, accessors, statements, mappings);
stream.dec_tab();
stream << "}" << std::endl;
@@ -250,13 +247,13 @@ public:
unsigned int current_arg = 0;
if (is_trans)
{
kernel->arg(current_arg++, cl_uint(utils::call_on_matrix(A, utils::size2_fun())));
kernel->arg(current_arg++, cl_uint(utils::call_on_matrix(A, utils::size1_fun())));
kernel->arg(current_arg++, cl_uint(tools::call_on_matrix(A, tools::size2_fun())));
kernel->arg(current_arg++, cl_uint(tools::call_on_matrix(A, tools::size1_fun())));
}
else
{
kernel->arg(current_arg++, cl_uint(utils::call_on_matrix(A, utils::size1_fun())));
kernel->arg(current_arg++, cl_uint(utils::call_on_matrix(A, utils::size2_fun())));
kernel->arg(current_arg++, cl_uint(tools::call_on_matrix(A, tools::size1_fun())));
kernel->arg(current_arg++, cl_uint(tools::call_on_matrix(A, tools::size2_fun())));
}

88
atidlas/templates/template_base.hpp
View File

@@ -12,10 +12,10 @@
#include "viennacl/scheduler/forwards.h"
#include "viennacl/scheduler/io.hpp"
#include "atidlas/lazy_program_compiler.hpp"
#include "atidlas/tools/lazy_program_compiler.hpp"
#include "atidlas/tools/tree_parsing.hpp"
#include "atidlas/tools/misc.hpp"
#include "atidlas/mapped_objects.hpp"
#include "atidlas/tree_parsing.hpp"
#include "atidlas/utils.hpp"
namespace atidlas
{
@@ -42,7 +42,7 @@ public:
private:
/** @brief Functor to map the statements to the types defined in mapped_objects.hpp */
class map_functor : public tree_parsing::traversal_functor
class map_functor : public tools::traversal_functor
{
viennacl::scheduler::statement_node_numeric_type numeric_type(viennacl::scheduler::statement const * statement, atidlas_int_t root_idx) const
@@ -62,48 +62,48 @@ private:
template<class T>
result_type binary_leaf(viennacl::scheduler::statement const * statement, atidlas_int_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)));
return result_type(new T(tools::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)));
return result_type(new mapped_host_scalar(tools::type_to_string<NumericT>::value(), binder_.get(NULL)));
}
/** @brief Scalar mapping */
template<class NumericT>
result_type operator()(viennacl::scalar<NumericT> const & scal) const
{
return result_type(new mapped_scalar(utils::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(scal))));
return result_type(new mapped_scalar(tools::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(scal))));
}
/** @brief Vector mapping */
template<class NumericT>
result_type operator()(viennacl::vector_base<NumericT> const & vec) const
{
return result_type(new mapped_vector(utils::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(vec))));
return result_type(new mapped_vector(tools::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(vec))));
}
/** @brief Implicit vector mapping */
template<class NumericT>
result_type operator()(viennacl::implicit_vector_base<NumericT> const & /*vec*/) const
{
return result_type(new mapped_implicit_vector(utils::type_to_string<NumericT>::value(), binder_.get(NULL)));
return result_type(new mapped_implicit_vector(tools::type_to_string<NumericT>::value(), binder_.get(NULL)));
}
/** @brief Matrix mapping */
template<class NumericT>
result_type operator()(viennacl::matrix_base<NumericT> const & mat) const
{
return result_type(new mapped_matrix(utils::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(mat))));
return result_type(new mapped_matrix(tools::type_to_string<NumericT>::value(), binder_.get(&viennacl::traits::handle(mat))));
}
/** @brief Implicit matrix mapping */
template<class NumericT>
result_type operator()(viennacl::implicit_matrix_base<NumericT> const & /*mat*/) const
{
return result_type(new mapped_implicit_matrix(utils::type_to_string<NumericT>::value(), binder_.get(NULL)));
return result_type(new mapped_implicit_matrix(tools::type_to_string<NumericT>::value(), binder_.get(NULL)));
}
/** @brief Traversal functor */
@@ -112,9 +112,9 @@ private:
viennacl::scheduler::statement_node const & root_node = statement.array()[root_idx];
if (leaf_t == LHS_NODE_TYPE && root_node.lhs.type_family != viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
mapping_.insert(mapping_type::value_type(key, utils::call_on_element(root_node.lhs, *this)));
mapping_.insert(mapping_type::value_type(key, tools::call_on_element(root_node.lhs, *this)));
else if (leaf_t == RHS_NODE_TYPE && root_node.rhs.type_family != viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
mapping_.insert(mapping_type::value_type(key, utils::call_on_element(root_node.rhs, *this)));
mapping_.insert(mapping_type::value_type(key, tools::call_on_element(root_node.rhs, *this)));
else if ( leaf_t== PARENT_NODE_TYPE)
{
if (root_node.op.type==viennacl::scheduler::OPERATION_BINARY_VECTOR_DIAG_TYPE)
@@ -142,7 +142,7 @@ private:
};
/** @brief functor for setting the arguments of a kernel */
class set_arguments_functor : public tree_parsing::traversal_functor
class set_arguments_functor : public tools::traversal_functor
{
public:
typedef void result_type;
@@ -213,9 +213,9 @@ private:
{
viennacl::scheduler::statement_node const & root_node = statement.array()[root_idx];
if (leaf_t==LHS_NODE_TYPE && root_node.lhs.type_family != viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
utils::call_on_element(root_node.lhs, *this);
tools::call_on_element(root_node.lhs, *this);
else if (leaf_t==RHS_NODE_TYPE && root_node.rhs.type_family != viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
utils::call_on_element(root_node.rhs, *this);
tools::call_on_element(root_node.rhs, *this);
}
private:
@@ -226,15 +226,15 @@ private:
protected:
static inline void compute_reduction(utils::kernel_generation_stream & os, std::string acc, std::string cur, viennacl::scheduler::op_element const & op)
static inline void compute_reduction(tools::kernel_generation_stream & os, std::string acc, std::string cur, viennacl::scheduler::op_element const & op)
{
if (utils::elementwise_function(op))
os << acc << "=" << tree_parsing::evaluate(op.type) << "(" << acc << "," << cur << ");" << std::endl;
if (tools::elementwise_function(op))
os << acc << "=" << tools::evaluate(op.type) << "(" << acc << "," << cur << ");" << std::endl;
else
os << acc << "= (" << acc << ")" << tree_parsing::evaluate(op.type) << "(" << cur << ");" << std::endl;
os << acc << "= (" << acc << ")" << tools::evaluate(op.type) << "(" << cur << ");" << std::endl;
}
static inline void compute_index_reduction(utils::kernel_generation_stream & os, std::string acc, std::string cur, std::string const & acc_value, std::string const & cur_value, viennacl::scheduler::op_element const & op)
static inline void compute_index_reduction(tools::kernel_generation_stream & os, std::string acc, std::string cur, std::string const & acc_value, std::string const & cur_value, viennacl::scheduler::op_element const & op)
{
// os << acc << " = " << cur_value << ">" << acc_value << "?" << cur << ":" << acc << ";" << std::endl;
os << acc << "= select(" << acc << "," << cur << "," << cur_value << ">" << acc_value << ");" << std::endl;
@@ -247,7 +247,7 @@ protected:
}
static inline void process_all(std::string const & type_key, std::string const & str,
utils::kernel_generation_stream & stream, std::vector<mapping_type> const & mappings)
tools::kernel_generation_stream & stream, std::vector<mapping_type> const & mappings)
{
for (std::vector<mapping_type>::const_iterator mit = mappings.begin(); mit != mappings.end(); ++mit)
for (mapping_type::const_iterator mmit = mit->begin(); mmit != mit->end(); ++mmit)
@@ -257,7 +257,7 @@ protected:
static inline void process_all_at(std::string const & type_key, std::string const & str,
utils::kernel_generation_stream & stream, std::vector<mapping_type> const & mappings,
tools::kernel_generation_stream & stream, std::vector<mapping_type> const & mappings,
size_t root_idx, leaf_t leaf)
{
for (std::vector<mapping_type>::const_iterator mit = mappings.begin(); mit != mappings.end(); ++mit)
@@ -290,8 +290,8 @@ protected:
static std::string generate_arguments(std::vector<mapping_type> const & mappings, std::multimap<std::string, std::string> const & accessors, statements_container const & statements)
{
utils::kernel_generation_stream stream;
tree_parsing::process(stream, PARENT_NODE_TYPE, accessors, statements, mappings);
tools::kernel_generation_stream stream;
tools::process(stream, PARENT_NODE_TYPE, accessors, statements, mappings);
std::string res = stream.str();
res.erase(res.rfind(','));
return res;
@@ -309,7 +309,7 @@ protected:
static std::string generate_arguments(std::string const & data_type, std::vector<mapping_type> const & mappings, statements_container const & statements)
{
return generate_arguments(mappings, utils::create_process_accessors("scalar", "__global #scalartype* #pointer,")
return generate_arguments(mappings, tools::create_process_accessors("scalar", "__global #scalartype* #pointer,")
("host_scalar", "#scalartype #name,")
("matrix", matrix_arguments(data_type))
("vector", vector_arguments(data_type))
@@ -323,7 +323,7 @@ protected:
{
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);
tools::traverse(*itt, itt->root(), set_arguments_functor(*binder,current_arg,kernel), true);
}
class invalid_template_exception : public std::exception
@@ -340,7 +340,7 @@ protected:
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)
static void fetching_loop_info(fetching_policy_type policy, std::string const & bound, tools::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)
{
@@ -396,18 +396,18 @@ protected:
for (statements_container::data_type::const_iterator it = statements.data().begin(); it != statements.data().end(); ++it)
{
//checks for vectors
std::vector<viennacl::scheduler::lhs_rhs_element> vectors = tree_parsing::filter_elements(viennacl::scheduler::DENSE_VECTOR_TYPE, *it);
std::vector<viennacl::scheduler::lhs_rhs_element> vectors = tools::filter_elements(viennacl::scheduler::DENSE_VECTOR_TYPE, *it);
for (std::vector<viennacl::scheduler::lhs_rhs_element>::iterator itt = vectors.begin(); itt != vectors.end(); ++itt)
if (utils::call_on_vector(*itt, utils::stride_fun())>1)
if (tools::call_on_vector(*itt, tools::stride_fun())>1)
return true;
//checks for matrix
std::vector<viennacl::scheduler::lhs_rhs_element> matrices = tree_parsing::filter_elements(viennacl::scheduler::DENSE_MATRIX_TYPE, *it);
std::vector<viennacl::scheduler::lhs_rhs_element> matrices = tools::filter_elements(viennacl::scheduler::DENSE_MATRIX_TYPE, *it);
for (std::vector<viennacl::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)
if (tools::call_on_matrix(*itt, tools::stride1_fun())>1 || tools::call_on_matrix(*itt, tools::stride2_fun())>2)
return true;
if(tree_parsing::filter_nodes(&is_offset_modifier, *it, true).empty()==false)
if(tools::filter_nodes(&is_offset_modifier, *it, true).empty()==false)
return true;
}
return false;
@@ -416,7 +416,7 @@ protected:
static atidlas_int_t vector_size(viennacl::scheduler::statement_node const & node, bool up_to_internal_size)
{
using namespace viennacl::scheduler;
using namespace utils;
using namespace tools;
if (node.op.type==OPERATION_BINARY_MATRIX_DIAG_TYPE)
{
atidlas_int_t size1 = up_to_internal_size?call_on_matrix(node.lhs, internal_size1_fun()):call_on_matrix(node.lhs, size1_fun());
@@ -434,10 +434,10 @@ protected:
//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 void operator()(tools::kernel_generation_stream & stream, unsigned int simd_width) const = 0;
};
static void element_wise_loop_1D(utils::kernel_generation_stream & stream, loop_body_base const & loop_body,
static void element_wise_loop_1D(tools::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);
@@ -468,7 +468,7 @@ protected:
if (simd_width==1)
return "(" + ptr + ")[" + offset + "] = " + value;
else
return utils::append_width("vstore", simd_width) + "(" + value + ", " + offset + ", " + ptr + ")";
return tools::append_width("vstore", simd_width) + "(" + value + ", " + offset + ", " + ptr + ")";
}
static std::string vload(unsigned int simd_width, std::string const & offset, std::string const & ptr)
@@ -476,7 +476,7 @@ protected:
if (simd_width==1)
return "(" + ptr + ")[" + offset + "]";
else
return utils::append_width("vload", simd_width) + "(" + offset + ", " + ptr + ")";
return tools::append_width("vload", simd_width) + "(" + offset + ", " + ptr + ")";
}
private:
@@ -500,7 +500,7 @@ public:
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);
tools::traverse(*sit, sit->root(), map_functor(*binder,*mit), true);
return generate_impl(kernel_prefix, statements, mappings);
}
@@ -530,15 +530,15 @@ protected:
for (statements_container::data_type::const_iterator it = statements.data().begin(); it != statements.data().end(); ++it)
{
//checks for vectors
std::vector<viennacl::scheduler::lhs_rhs_element> vectors = tree_parsing::filter_elements(viennacl::scheduler::DENSE_VECTOR_TYPE, *it);
std::vector<viennacl::scheduler::lhs_rhs_element> vectors = tools::filter_elements(viennacl::scheduler::DENSE_VECTOR_TYPE, *it);
for (std::vector<viennacl::scheduler::lhs_rhs_element>::iterator itt = vectors.begin(); itt != vectors.end(); ++itt)
if (utils::call_on_vector(*itt, utils::stride_fun())>1)
if (tools::call_on_vector(*itt, tools::stride_fun())>1)
return true;
//checks for matrix
std::vector<viennacl::scheduler::lhs_rhs_element> matrices = tree_parsing::filter_elements(viennacl::scheduler::DENSE_MATRIX_TYPE, *it);
std::vector<viennacl::scheduler::lhs_rhs_element> matrices = tools::filter_elements(viennacl::scheduler::DENSE_MATRIX_TYPE, *it);
for (std::vector<viennacl::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)
if (tools::call_on_matrix(*itt, tools::stride1_fun())>1 || tools::call_on_matrix(*itt, tools::stride2_fun())>2)
return true;
}
return false;
@@ -566,7 +566,7 @@ public:
using namespace viennacl::tools;
viennacl::scheduler::statement const & statement = statements.data().front();
unsigned int scalartype_size = utils::size_of(lhs_most(statement.array(), statement.root()).lhs.numeric_type);
unsigned int scalartype_size = tools::size_of(lhs_most(statement.array(), statement.root()).lhs.numeric_type);
//Query device informations
size_t lmem_available = static_cast<size_t>(device.local_mem_size());

20
atidlas/templates/vector_axpy.hpp
View File

@@ -7,10 +7,6 @@
#include "viennacl/scheduler/forwards.h"
#include "viennacl/tools/tools.hpp"
#include "atidlas/mapped_objects.hpp"
#include "atidlas/tree_parsing.hpp"
#include "atidlas/forwards.h"
#include "atidlas/utils.hpp"
#include "atidlas/templates/template_base.hpp"
namespace atidlas
@@ -41,18 +37,18 @@ private:
std::vector<std::string> result;
for (unsigned int i = 0; i < 2; ++i)
{
utils::kernel_generation_stream stream;
tools::kernel_generation_stream stream;
unsigned int simd_width = (i==0)?1:p_.simd_width;
std::string str_simd_width = tools::to_string(simd_width);
std::string data_type = utils::append_width("#scalartype",simd_width);
std::string data_type = tools::append_width("#scalartype",simd_width);
stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << ",1,1)))" << std::endl;
stream << "__kernel void " << kernel_prefix << i << "(unsigned int N," << generate_arguments(data_type, mappings, statements) << ")" << std::endl;
stream << "{" << std::endl;
stream.inc_tab();
tree_parsing::process(stream, PARENT_NODE_TYPE,
utils::create_process_accessors("scalar", "#scalartype #namereg = *#pointer;")
tools::process(stream, PARENT_NODE_TYPE,
tools::create_process_accessors("scalar", "#scalartype #namereg = *#pointer;")
("matrix", "#pointer += #start1 + #start2*#ld;")
("vector", "#pointer += #start;")
("vector", "#start/=" + str_simd_width + ";"), statements, mappings);
@@ -62,20 +58,20 @@ private:
stream << "for(unsigned int i = " << init << "; i < " << upper_bound << "; i += " << inc << ")" << std::endl;
stream << "{" << std::endl;
stream.inc_tab();
tree_parsing::process(stream, PARENT_NODE_TYPE,
utils::create_process_accessors("vector", data_type + " #namereg = #pointer[i*#stride];")
tools::process(stream, PARENT_NODE_TYPE,
tools::create_process_accessors("vector", data_type + " #namereg = #pointer[i*#stride];")
("matrix_row", "#scalartype #namereg = #pointer[$OFFSET{#row*#stride1, i*#stride2}];")
("matrix_column", "#scalartype #namereg = #pointer[$OFFSET{i*#stride1,#column*#stride2}];")
("matrix_diag", "#scalartype #namereg = #pointer[#diag_offset<0?$OFFSET{(i - #diag_offset)*#stride1, i*#stride2}:$OFFSET{i*#stride1, (i + #diag_offset)*#stride2}];")
, statements, mappings);
tree_parsing::evaluate(stream, PARENT_NODE_TYPE, utils::create_evaluate_accessors("vector", "#namereg")
tools::evaluate(stream, PARENT_NODE_TYPE, tools::create_evaluate_accessors("vector", "#namereg")
("matrix_row", "#namereg")
("matrix_column", "#namereg")
("matrix_diag", "#namereg")
("scalar", "#namereg"), statements, mappings);
tree_parsing::process(stream, LHS_NODE_TYPE, utils::create_process_accessors("vector", "#pointer[i*#stride] = #namereg;")
tools::process(stream, LHS_NODE_TYPE, tools::create_process_accessors("vector", "#pointer[i*#stride] = #namereg;")
("matrix_row", "#pointer[$OFFSET{#row, i}] = #namereg;")
("matrix_column", "#pointer[$OFFSET{i, #column}] = #namereg;")
("matrix_diag", "#pointer[#diag_offset<0?$OFFSET{i - #diag_offset, i}:$OFFSET{i, i + #diag_offset}] = #namereg;")

2
atidlas/execution_handler.hpp → atidlas/tools/execution_handler.hpp
View File

@@ -5,7 +5,7 @@
#include "viennacl/tools/shared_ptr.hpp"
#include "atidlas/lazy_program_compiler.hpp"
#include "atidlas/tools/lazy_program_compiler.hpp"
#include "atidlas/templates/template_base.hpp"
namespace atidlas

44
atidlas/tools/lazy_program_compiler.hpp Normal file
View File

@@ -0,0 +1,44 @@
#ifndef ATIDLAS_LAZY_PROGRAM_COMPILER_HPP
#define ATIDLAS_LAZY_PROGRAM_COMPILER_HPP
#include <map>
#include "viennacl/ocl/context.hpp"
namespace atidlas
{
class lazy_program_compiler
{
public:
lazy_program_compiler(viennacl::ocl::context * ctx, std::string const & name, std::string const & src, bool force_recompilation) : ctx_(ctx), program_(NULL), name_(name), src_(src), force_recompilation_(force_recompilation){ }
lazy_program_compiler(viennacl::ocl::context * ctx, std::string const & name, bool force_recompilation) : ctx_(ctx), program_(NULL), name_(name), force_recompilation_(force_recompilation){ }
void add(std::string const & src) { src_+=src; }
std::string const & src() const { return src_; }
viennacl::ocl::program & program()
{
if(program_==NULL)
{
// std::cout << src_ << std::endl;
if (force_recompilation_ && ctx_->has_program(name_))
ctx_->delete_program(name_);
if (!ctx_->has_program(name_))
ctx_->add_program(src_, name_);
program_ = &ctx_->get_program(name_);
}
return *program_;
}
private:
viennacl::ocl::context * ctx_;
viennacl::ocl::program * program_;
std::string name_;
std::string src_;
bool force_recompilation_;
};
}
#endif

4
atidlas/utils.hpp → atidlas/tools/misc.hpp
View File

@@ -19,7 +19,7 @@
namespace atidlas
{
namespace utils
namespace tools
{
//CUDA Conversion
@@ -492,7 +492,7 @@ inline bool elementwise_function(viennacl::scheduler::op_element const & op)
inline viennacl::scheduler::lhs_rhs_element & lhs_rhs_element(viennacl::scheduler::statement const & st, atidlas_int_t idx, leaf_t leaf)
{
using namespace tree_parsing;
using namespace tools;
assert(leaf==LHS_NODE_TYPE || leaf==RHS_NODE_TYPE);
if (leaf==LHS_NODE_TYPE)
return const_cast<viennacl::scheduler::lhs_rhs_element &>(st.array()[idx].lhs);

68
atidlas/tree_parsing.hpp → atidlas/tools/tree_parsing.hpp
View File

@@ -9,13 +9,13 @@
#include "viennacl/scheduler/forwards.h"
#include "atidlas/mapped_objects.hpp"
#include "atidlas/utils.hpp"
#include "atidlas/tools/misc.hpp"
#include "atidlas/forwards.h"
namespace atidlas
{
namespace tree_parsing
namespace tools
{
/** @brief base functor class for traversing a statement */
@@ -31,7 +31,7 @@ template<class Fun>
inline void traverse(viennacl::scheduler::statement const & statement, atidlas_int_t root_idx, Fun const & fun, bool inspect)
{
viennacl::scheduler::statement_node const & root_node = statement.array()[root_idx];
bool recurse = utils::node_leaf(root_node.op)?inspect:true;
bool recurse = tools::node_leaf(root_node.op)?inspect:true;
fun.call_before_expansion(statement, root_idx);
@@ -80,7 +80,7 @@ private:
std::vector<size_t> filter_nodes(bool (*pred)(viennacl::scheduler::statement_node const & node), viennacl::scheduler::statement const & statement, bool inspect)
{
std::vector<size_t> res;
tree_parsing::traverse(statement, statement.root(), filter_fun(pred, res), inspect);
tools::traverse(statement, statement.root(), filter_fun(pred, res), inspect);
return res;
}
@@ -105,7 +105,7 @@ private:
std::vector<viennacl::scheduler::lhs_rhs_element> filter_elements(viennacl::scheduler::statement_node_subtype subtype, viennacl::scheduler::statement const & statement)
{
std::vector<viennacl::scheduler::lhs_rhs_element> res;
tree_parsing::traverse(statement, statement.root(), filter_elements_fun(subtype, res), true);
tools::traverse(statement, statement.root(), filter_elements_fun(subtype, res), true);
return res;
}
@@ -227,7 +227,7 @@ inline const char * operator_string(viennacl::scheduler::operation_node_type typ
}
/** @brief functor for generating the expression string from a statement */
class evaluate_expression_traversal: public tree_parsing::traversal_functor
class evaluate_expression_traversal: public tools::traversal_functor
{
private:
std::map<std::string, std::string> const & accessors_;
@@ -240,9 +240,9 @@ public:
void call_before_expansion(viennacl::scheduler::statement const & statement, atidlas_int_t root_idx) const
{
viennacl::scheduler::statement_node const & root_node = statement.array()[root_idx];
if ((root_node.op.type_family==viennacl::scheduler::OPERATION_UNARY_TYPE_FAMILY || utils::elementwise_function(root_node.op))
&& !utils::node_leaf(root_node.op))
str_+=tree_parsing::evaluate(root_node.op.type);
if ((root_node.op.type_family==viennacl::scheduler::OPERATION_UNARY_TYPE_FAMILY || tools::elementwise_function(root_node.op))
&& !tools::node_leaf(root_node.op))
str_+=tools::evaluate(root_node.op.type);
str_+="(";
}
@@ -258,11 +258,11 @@ public:
mapping_type::key_type key = std::make_pair(root_idx, leaf);
if (leaf==PARENT_NODE_TYPE)
{
if (utils::node_leaf(root_node.op))
if (tools::node_leaf(root_node.op))
str_ += mapping_.at(key)->evaluate(accessors_);
else if (utils::elementwise_operator(root_node.op))
str_ += tree_parsing::evaluate(root_node.op.type);
else if (root_node.op.type_family!=viennacl::scheduler::OPERATION_UNARY_TYPE_FAMILY && utils::elementwise_function(root_node.op))
else if (tools::elementwise_operator(root_node.op))
str_ += tools::evaluate(root_node.op.type);
else if (root_node.op.type_family!=viennacl::scheduler::OPERATION_UNARY_TYPE_FAMILY && tools::elementwise_function(root_node.op))
str_ += ",";
}
else
@@ -292,24 +292,24 @@ inline std::string evaluate(leaf_t leaf, std::map<std::string, std::string> cons
if (leaf==RHS_NODE_TYPE)
{
if (root_node.rhs.type_family==viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
tree_parsing::traverse(statement, root_node.rhs.node_index, traversal_functor, false);
tools::traverse(statement, root_node.rhs.node_index, traversal_functor, false);
else
traversal_functor(statement, root_idx, leaf);
}
else if (leaf==LHS_NODE_TYPE)
{
if (root_node.lhs.type_family==viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
tree_parsing::traverse(statement, root_node.lhs.node_index, traversal_functor, false);
tools::traverse(statement, root_node.lhs.node_index, traversal_functor, false);
else
traversal_functor(statement, root_idx, leaf);
}
else
tree_parsing::traverse(statement, root_idx, traversal_functor, false);
tools::traverse(statement, root_idx, traversal_functor, false);
return res;
}
inline void evaluate(utils::kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors,
inline void evaluate(tools::kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors,
statements_container const & statements, std::vector<mapping_type> const & mappings)
{
statements_container::data_type::const_iterator sit;
@@ -321,10 +321,10 @@ inline void evaluate(utils::kernel_generation_stream & stream, leaf_t leaf, std:
/** @brief functor for fetching or writing-back the elements in a statement */
class process_traversal : public tree_parsing::traversal_functor
class process_traversal : public tools::traversal_functor
{
public:
process_traversal(std::multimap<std::string, std::string> const & accessors, utils::kernel_generation_stream & stream,
process_traversal(std::multimap<std::string, std::string> const & accessors, tools::kernel_generation_stream & stream,
mapping_type const & mapping, std::set<std::string> & already_processed) : accessors_(accessors), stream_(stream), mapping_(mapping), already_processed_(already_processed){ }
void operator()(viennacl::scheduler::statement const & /*statement*/, atidlas_int_t root_idx, leaf_t leaf) const
@@ -347,12 +347,12 @@ public:
private:
std::multimap<std::string, std::string> accessors_;
utils::kernel_generation_stream & stream_;
tools::kernel_generation_stream & stream_;
mapping_type const & mapping_;
std::set<std::string> & already_processed_;
};
inline void process(utils::kernel_generation_stream & stream, leaf_t leaf, std::multimap<std::string, std::string> const & accessors,
inline void process(tools::kernel_generation_stream & stream, leaf_t leaf, std::multimap<std::string, std::string> const & accessors,
viennacl::scheduler::statement const & statement, size_t root_idx, mapping_type const & mapping, std::set<std::string> & already_processed)
{
process_traversal traversal_functor(accessors, stream, mapping, already_processed);
@@ -361,24 +361,24 @@ inline void process(utils::kernel_generation_stream & stream, leaf_t leaf, std::
if (leaf==RHS_NODE_TYPE)
{
if (root_node.rhs.type_family==viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
tree_parsing::traverse(statement, root_node.rhs.node_index, traversal_functor, true);
tools::traverse(statement, root_node.rhs.node_index, traversal_functor, true);
else
traversal_functor(statement, root_idx, leaf);
}
else if (leaf==LHS_NODE_TYPE)
{
if (root_node.lhs.type_family==viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
tree_parsing::traverse(statement, root_node.lhs.node_index, traversal_functor, true);
tools::traverse(statement, root_node.lhs.node_index, traversal_functor, true);
else
traversal_functor(statement, root_idx, leaf);
}
else
{
tree_parsing::traverse(statement, root_idx, traversal_functor, true);
tools::traverse(statement, root_idx, traversal_functor, true);
}
}
inline void process(utils::kernel_generation_stream & stream, leaf_t leaf, std::multimap<std::string, std::string> const & accessors,
inline void process(tools::kernel_generation_stream & stream, leaf_t leaf, std::multimap<std::string, std::string> const & accessors,
statements_container const & statements, std::vector<mapping_type> const & mappings)
{
statements_container::data_type::const_iterator sit;
@@ -414,7 +414,7 @@ public:
{
*ptr_++='h'; //host
*ptr_++='s'; //scalar
*ptr_++=utils::first_letter_of_type<NumericT>::value();
*ptr_++=tools::first_letter_of_type<NumericT>::value();
}
/** @brief Scalar mapping */
@@ -422,7 +422,7 @@ public:
inline result_type operator()(viennacl::scalar<NumericT> const & scal) const
{
*ptr_++='s'; //scalar
*ptr_++=utils::first_letter_of_type<NumericT>::value();
*ptr_++=tools::first_letter_of_type<NumericT>::value();
append_id(ptr_, binder_.get(&viennacl::traits::handle(scal)));
}
@@ -431,7 +431,7 @@ public:
inline result_type operator()(viennacl::vector_base<NumericT> const & vec) const
{
*ptr_++='v'; //vector
*ptr_++=utils::first_letter_of_type<NumericT>::value();
*ptr_++=tools::first_letter_of_type<NumericT>::value();
append_id(ptr_, binder_.get(&viennacl::traits::handle(vec)));
}
@@ -441,7 +441,7 @@ public:
{
*ptr_++='i'; //implicit
*ptr_++='v'; //vector
*ptr_++=utils::first_letter_of_type<NumericT>::value();
*ptr_++=tools::first_letter_of_type<NumericT>::value();
}
/** @brief Matrix mapping */
@@ -449,7 +449,7 @@ public:
inline result_type operator()(viennacl::matrix_base<NumericT> const & mat) const
{
*ptr_++='m'; //Matrix
*ptr_++=utils::first_letter_of_type<NumericT>::value();
*ptr_++=tools::first_letter_of_type<NumericT>::value();
append_id(ptr_, binder_.get(&viennacl::traits::handle(mat)));
}
@@ -459,7 +459,7 @@ public:
{
*ptr_++='i'; //implicit
*ptr_++='m'; //matrix
*ptr_++=utils::first_letter_of_type<NumericT>::value();
*ptr_++=tools::first_letter_of_type<NumericT>::value();
}
static inline void append(char*& p, const char * str)
@@ -473,9 +473,9 @@ public:
{
viennacl::scheduler::statement_node const & root_node = statement.array()[root_idx];
if (leaf_t==LHS_NODE_TYPE && root_node.lhs.type_family != viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
utils::call_on_element(root_node.lhs, *this);
tools::call_on_element(root_node.lhs, *this);
else if (root_node.op.type_family==viennacl::scheduler::OPERATION_BINARY_TYPE_FAMILY && leaf_t==RHS_NODE_TYPE && root_node.rhs.type_family != viennacl::scheduler::COMPOSITE_OPERATION_FAMILY)
utils::call_on_element(root_node.rhs, *this);
tools::call_on_element(root_node.rhs, *this);
else if (leaf_t==PARENT_NODE_TYPE)
append_id(ptr_,root_node.op.type);
}
@@ -495,7 +495,7 @@ inline std::string statements_representation(statements_container const & statem
*program_name++='s';
tools::shared_ptr<symbolic_binder> binder = make_binder(binding_policy);
for (statements_container::data_type::const_iterator it = statements.data().begin(); it != statements.data().end(); ++it)
tree_parsing::traverse(*it, it->root(), tree_parsing::statement_representation_functor(*binder, program_name),true);
tools::traverse(*it, it->root(), tools::statement_representation_functor(*binder, program_name),true);
*program_name='\0';
return std::string(program_name_vector.data());
}

18
tests/matrix_axpy.cpp
View File

@@ -122,15 +122,15 @@ void test_impl(NumericT epsilon)
TEST_OPERATIONS(range, slice, range)
TEST_OPERATIONS(range, slice, slice)
// TEST_OPERATIONS(slice, matrix, matrix)
// TEST_OPERATIONS(slice, matrix, range)
// TEST_OPERATIONS(slice, matrix, slice)
// TEST_OPERATIONS(slice, range, matrix)
// TEST_OPERATIONS(slice, range, range)
// TEST_OPERATIONS(slice, range, slice)
// TEST_OPERATIONS(slice, slice, matrix)
// TEST_OPERATIONS(slice, slice, range)
// TEST_OPERATIONS(slice, slice, slice)
TEST_OPERATIONS(slice, matrix, matrix)
TEST_OPERATIONS(slice, matrix, range)
TEST_OPERATIONS(slice, matrix, slice)
TEST_OPERATIONS(slice, range, matrix)
TEST_OPERATIONS(slice, range, range)
TEST_OPERATIONS(slice, range, slice)
TEST_OPERATIONS(slice, slice, matrix)
TEST_OPERATIONS(slice, slice, range)
TEST_OPERATIONS(slice, slice, slice)
}
int main()