diff --git a/atidlas/execute.hpp b/atidlas/execute.hpp index 2ad9334d0..390790d48 100644 --- a/atidlas/execute.hpp +++ b/atidlas/execute.hpp @@ -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); diff --git a/atidlas/forwards.h b/atidlas/forwards.h index be3fa26f9..ecf3231ec 100644 --- a/atidlas/forwards.h +++ b/atidlas/forwards.h @@ -137,7 +137,7 @@ private: std::string message_; }; -namespace utils +namespace tools { class kernel_generation_stream; } @@ -157,13 +157,13 @@ typedef std::pair mapping_key; typedef std::map > mapping_type; -namespace tree_parsing +namespace tools { template 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 const & accessors, + inline void process(tools::kernel_generation_stream & stream, leaf_t leaf, std::multimap const & accessors, viennacl::scheduler::statement const & statement, size_t root_idx, mapping_type const & mapping, std::set & already_processed); inline std::string evaluate(leaf_t leaf, std::map const & accessors, viennacl::scheduler::statement const & statement, atidlas_int_t root_idx,mapping_type const & mapping); } diff --git a/atidlas/mapped_objects.hpp b/atidlas/mapped_objects.hpp index af533a414..1abc6b1d0 100644 --- a/atidlas/mapped_objects.hpp +++ b/atidlas/mapped_objects.hpp @@ -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 const & accessors) + void process_recursive(tools::kernel_generation_stream & stream, leaf_t leaf, std::multimap const & accessors) { std::set 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 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 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 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 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 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 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: diff --git a/atidlas/templates/matrix_axpy.hpp b/atidlas/templates/matrix_axpy.hpp index 3d7d1d8a7..7b4f22185 100644 --- a/atidlas/templates/matrix_axpy.hpp +++ b/atidlas/templates/matrix_axpy.hpp @@ -4,9 +4,6 @@ #include -#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 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); diff --git a/atidlas/templates/matrix_product.hpp b/atidlas/templates/matrix_product.hpp index 32aaf2497..30b2a1959 100644 --- a/atidlas/templates/matrix_product.hpp +++ b/atidlas/templates/matrix_product.hpp @@ -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 & 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 & 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(&viennacl::scheduler::lhs_rhs_element::matrix_float, stcopy, A, B, C, beta, beta.host_float, programs, kernel_prefix); diff --git a/atidlas/templates/reduction.hpp b/atidlas/templates/reduction.hpp index ef18b3306..407970791 100644 --- a/atidlas/templates/reduction.hpp +++ b/atidlas/templates/reduction.hpp @@ -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 exprs, + inline void reduce_1d_local_memory(tools::kernel_generation_stream & stream, unsigned int size, std::vector 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 const & mappings, unsigned int simd_width) const { - utils::kernel_generation_stream stream; + tools::kernel_generation_stream stream; std::vector exprs; for (std::vector::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 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::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 reductions_idx = tree_parsing::filter_nodes(&utils::is_reduction, *it, false); + std::vector reductions_idx = tools::filter_nodes(&tools::is_reduction, *it, false); size = static_cast(vector_size(lhs_most(it->array(), reductions_idx[0]), false)); for (std::vector::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::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)); diff --git a/atidlas/templates/row_wise_reduction.hpp b/atidlas/templates/row_wise_reduction.hpp index abca8c5bf..a61c72915 100644 --- a/atidlas/templates/row_wise_reduction.hpp +++ b/atidlas/templates/row_wise_reduction.hpp @@ -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 & 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 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::const_iterator it = exprs.begin(); it != exprs.end(); ++it) { @@ -181,7 +178,7 @@ private: std::map 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()))); } diff --git a/atidlas/templates/template_base.hpp b/atidlas/templates/template_base.hpp index bbcec9813..363cfc2b8 100644 --- a/atidlas/templates/template_base.hpp +++ b/atidlas/templates/template_base.hpp @@ -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 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 result_type operator()(NumericT const & /*scalar*/) const { - return result_type(new mapped_host_scalar(utils::type_to_string::value(), binder_.get(NULL))); + return result_type(new mapped_host_scalar(tools::type_to_string::value(), binder_.get(NULL))); } /** @brief Scalar mapping */ template result_type operator()(viennacl::scalar const & scal) const { - return result_type(new mapped_scalar(utils::type_to_string::value(), binder_.get(&viennacl::traits::handle(scal)))); + return result_type(new mapped_scalar(tools::type_to_string::value(), binder_.get(&viennacl::traits::handle(scal)))); } /** @brief Vector mapping */ template result_type operator()(viennacl::vector_base const & vec) const { - return result_type(new mapped_vector(utils::type_to_string::value(), binder_.get(&viennacl::traits::handle(vec)))); + return result_type(new mapped_vector(tools::type_to_string::value(), binder_.get(&viennacl::traits::handle(vec)))); } /** @brief Implicit vector mapping */ template result_type operator()(viennacl::implicit_vector_base const & /*vec*/) const { - return result_type(new mapped_implicit_vector(utils::type_to_string::value(), binder_.get(NULL))); + return result_type(new mapped_implicit_vector(tools::type_to_string::value(), binder_.get(NULL))); } /** @brief Matrix mapping */ template result_type operator()(viennacl::matrix_base const & mat) const { - return result_type(new mapped_matrix(utils::type_to_string::value(), binder_.get(&viennacl::traits::handle(mat)))); + return result_type(new mapped_matrix(tools::type_to_string::value(), binder_.get(&viennacl::traits::handle(mat)))); } /** @brief Implicit matrix mapping */ template result_type operator()(viennacl::implicit_matrix_base const & /*mat*/) const { - return result_type(new mapped_implicit_matrix(utils::type_to_string::value(), binder_.get(NULL))); + return result_type(new mapped_implicit_matrix(tools::type_to_string::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 const & mappings) + tools::kernel_generation_stream & stream, std::vector const & mappings) { for (std::vector::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 const & mappings, + tools::kernel_generation_stream & stream, std::vector const & mappings, size_t root_idx, leaf_t leaf) { for (std::vector::const_iterator mit = mappings.begin(); mit != mappings.end(); ++mit) @@ -290,8 +290,8 @@ protected: static std::string generate_arguments(std::vector const & mappings, std::multimap 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 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 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 vectors = tree_parsing::filter_elements(viennacl::scheduler::DENSE_VECTOR_TYPE, *it); + std::vector vectors = tools::filter_elements(viennacl::scheduler::DENSE_VECTOR_TYPE, *it); for (std::vector::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 matrices = tree_parsing::filter_elements(viennacl::scheduler::DENSE_MATRIX_TYPE, *it); + std::vector matrices = tools::filter_elements(viennacl::scheduler::DENSE_MATRIX_TYPE, *it); for (std::vector::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 mappings(statements.data().size()); tools::shared_ptr 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 vectors = tree_parsing::filter_elements(viennacl::scheduler::DENSE_VECTOR_TYPE, *it); + std::vector vectors = tools::filter_elements(viennacl::scheduler::DENSE_VECTOR_TYPE, *it); for (std::vector::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 matrices = tree_parsing::filter_elements(viennacl::scheduler::DENSE_MATRIX_TYPE, *it); + std::vector matrices = tools::filter_elements(viennacl::scheduler::DENSE_MATRIX_TYPE, *it); for (std::vector::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(device.local_mem_size()); diff --git a/atidlas/templates/vector_axpy.hpp b/atidlas/templates/vector_axpy.hpp index ef9eb93ef..40b1c0bca 100644 --- a/atidlas/templates/vector_axpy.hpp +++ b/atidlas/templates/vector_axpy.hpp @@ -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 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;") diff --git a/atidlas/execution_handler.hpp b/atidlas/tools/execution_handler.hpp similarity index 97% rename from atidlas/execution_handler.hpp rename to atidlas/tools/execution_handler.hpp index 9557745fd..31e5adee6 100644 --- a/atidlas/execution_handler.hpp +++ b/atidlas/tools/execution_handler.hpp @@ -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 diff --git a/atidlas/tools/lazy_program_compiler.hpp b/atidlas/tools/lazy_program_compiler.hpp new file mode 100644 index 000000000..8cf18436b --- /dev/null +++ b/atidlas/tools/lazy_program_compiler.hpp @@ -0,0 +1,44 @@ +#ifndef ATIDLAS_LAZY_PROGRAM_COMPILER_HPP +#define ATIDLAS_LAZY_PROGRAM_COMPILER_HPP + +#include +#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 diff --git a/atidlas/utils.hpp b/atidlas/tools/misc.hpp similarity index 99% rename from atidlas/utils.hpp rename to atidlas/tools/misc.hpp index b327b90a4..d5c918779 100644 --- a/atidlas/utils.hpp +++ b/atidlas/tools/misc.hpp @@ -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(st.array()[idx].lhs); diff --git a/atidlas/tree_parsing.hpp b/atidlas/tools/tree_parsing.hpp similarity index 88% rename from atidlas/tree_parsing.hpp rename to atidlas/tools/tree_parsing.hpp index 0a46a37a8..b9afadcbc 100644 --- a/atidlas/tree_parsing.hpp +++ b/atidlas/tools/tree_parsing.hpp @@ -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 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 filter_nodes(bool (*pred)(viennacl::scheduler::statement_node const & node), viennacl::scheduler::statement const & statement, bool inspect) { std::vector 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 filter_elements(viennacl::scheduler::statement_node_subtype subtype, viennacl::scheduler::statement const & statement) { std::vector 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 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 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 const & accessors, +inline void evaluate(tools::kernel_generation_stream & stream, leaf_t leaf, std::map const & accessors, statements_container const & statements, std::vector 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 const & accessors, utils::kernel_generation_stream & stream, + process_traversal(std::multimap const & accessors, tools::kernel_generation_stream & stream, mapping_type const & mapping, std::set & 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 accessors_; - utils::kernel_generation_stream & stream_; + tools::kernel_generation_stream & stream_; mapping_type const & mapping_; std::set & already_processed_; }; -inline void process(utils::kernel_generation_stream & stream, leaf_t leaf, std::multimap const & accessors, +inline void process(tools::kernel_generation_stream & stream, leaf_t leaf, std::multimap const & accessors, viennacl::scheduler::statement const & statement, size_t root_idx, mapping_type const & mapping, std::set & 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 const & accessors, +inline void process(tools::kernel_generation_stream & stream, leaf_t leaf, std::multimap const & accessors, statements_container const & statements, std::vector 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::value(); + *ptr_++=tools::first_letter_of_type::value(); } /** @brief Scalar mapping */ @@ -422,7 +422,7 @@ public: inline result_type operator()(viennacl::scalar const & scal) const { *ptr_++='s'; //scalar - *ptr_++=utils::first_letter_of_type::value(); + *ptr_++=tools::first_letter_of_type::value(); append_id(ptr_, binder_.get(&viennacl::traits::handle(scal))); } @@ -431,7 +431,7 @@ public: inline result_type operator()(viennacl::vector_base const & vec) const { *ptr_++='v'; //vector - *ptr_++=utils::first_letter_of_type::value(); + *ptr_++=tools::first_letter_of_type::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::value(); + *ptr_++=tools::first_letter_of_type::value(); } /** @brief Matrix mapping */ @@ -449,7 +449,7 @@ public: inline result_type operator()(viennacl::matrix_base const & mat) const { *ptr_++='m'; //Matrix - *ptr_++=utils::first_letter_of_type::value(); + *ptr_++=tools::first_letter_of_type::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::value(); + *ptr_++=tools::first_letter_of_type::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 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()); } diff --git a/tests/matrix_axpy.cpp b/tests/matrix_axpy.cpp index 1325a9553..386f6baf1 100644 --- a/tests/matrix_axpy.cpp +++ b/tests/matrix_axpy.cpp @@ -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()