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

Clearer array_expression with hopefully lower overhead.

Browse Source 
Also removed pyc's
This commit is contained in:
Philippe Tillet
2015-01-31 22:01:48 -05:00
parent 13ec84fbda
commit d29f1252ad
36 changed files with 566 additions and 572 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
bench/blas.cpp
View File

@@ -40,7 +40,7 @@ void bench(ad::numeric_type dtype)
total_time += times.back();\ total_time += times.back();\
}\ }\
float tres = median(times);\ float tres = median(times);\
std::cout << " " << PERF << std::flush;\ std::cout << " " << tres << std::flush;\
} }
#define CL_BENCHMARK(OP, PERF) BENCHMARK(OP, PERF, ad::cl_ext::synchronize(ad::cl_ext::default_context())) #define CL_BENCHMARK(OP, PERF) BENCHMARK(OP, PERF, ad::cl_ext::synchronize(ad::cl_ext::default_context()))

10
include/atidlas/backend/mapped_object.h
View File

@@ -24,7 +24,7 @@ typedef std::map<mapping_key, tools::shared_ptr<mapped_object> > mapping_type;
/** @brief Mapped Object /** @brief Mapped Object
* *
* This object populates the symbolic mapping associated with a symbolic_expression. (root_id, LHS|RHS|PARENT) => mapped_object * This object populates the symbolic mapping associated with a array_expression. (root_id, LHS|RHS|PARENT) => mapped_object
* The tree can then be reconstructed in its symbolic form * The tree can then be reconstructed in its symbolic form
*/ */
class mapped_object class mapped_object
@@ -46,9 +46,9 @@ protected:
public: public:
struct node_info struct node_info
{ {
node_info(mapping_type const * _mapping, symbolic_expression const * _symbolic_expression, int_t _root_idx); node_info(mapping_type const * _mapping, array_expression const * _array_expression, int_t _root_idx);
mapping_type const * mapping; mapping_type const * mapping;
atidlas::symbolic_expression const * symbolic_expression; atidlas::array_expression const * array_expression;
int_t root_idx; int_t root_idx;
}; };
@@ -100,8 +100,8 @@ public:
mapped_reduction(std::string const & scalartype, unsigned int id, node_info info, std::string const & type_key); mapped_reduction(std::string const & scalartype, unsigned int id, node_info info, std::string const & type_key);
int_t root_idx() const; int_t root_idx() const;
atidlas::symbolic_expression const & symbolic_expression() const; atidlas::array_expression const & array_expression() const;
symbolic_expression_node root_node() const; array_expression::node root_node() const;
bool is_index_reduction() const; bool is_index_reduction() const;
op_element root_op() const; op_element root_op() const;
}; };

74
include/atidlas/backend/parse.h
View File

@@ -13,8 +13,8 @@ namespace detail
{ {
bool is_node_leaf(op_element const & op); bool is_node_leaf(op_element const & op);
bool is_scalar_reduction(symbolic_expression_node const & node); bool is_scalar_reduction(array_expression::node const & node);
bool is_vector_reduction(symbolic_expression_node const & node); bool is_vector_reduction(array_expression::node const & node);
bool is_elementwise_operator(op_element const & op); bool is_elementwise_operator(op_element const & op);
bool is_elementwise_function(op_element const & op); bool is_elementwise_function(op_element const & op);
bool is_cast(op_element const & op); bool is_cast(op_element const & op);
@@ -23,54 +23,54 @@ namespace detail
class scalar; class scalar;
/** @brief base functor class for traversing a symbolic_expression */ /** @brief base functor class for traversing a array_expression */
class traversal_functor class traversal_functor
{ {
public: public:
void call_before_expansion(symbolic_expression const &, int_t) const { } void call_before_expansion(array_expression const &, int_t) const { }
void call_after_expansion(symbolic_expression const &, int_t) const { } void call_after_expansion(array_expression const &, int_t) const { }
}; };
/** @brief Recursively execute a functor on a symbolic_expression */ /** @brief Recursively execute a functor on a array_expression */
template<class Fun> template<class Fun>
inline void traverse(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, Fun const & fun, bool inspect) inline void traverse(atidlas::array_expression const & array_expression, int_t root_idx, Fun const & fun, bool inspect)
{ {
symbolic_expression_node const & root_node = symbolic_expression.tree()[root_idx]; array_expression::node const & root_node = array_expression.tree()[root_idx];
bool recurse = detail::is_node_leaf(root_node.op)?inspect:true; bool recurse = detail::is_node_leaf(root_node.op)?inspect:true;
fun.call_before_expansion(symbolic_expression, root_idx); fun.call_before_expansion(array_expression, root_idx);
//Lhs: //Lhs:
if (recurse) if (recurse)
{ {
if (root_node.lhs.type_family==COMPOSITE_OPERATOR_FAMILY) if (root_node.lhs.type_family==COMPOSITE_OPERATOR_FAMILY)
traverse(symbolic_expression, root_node.lhs.node_index, fun, inspect); traverse(array_expression, root_node.lhs.node_index, fun, inspect);
if (root_node.lhs.type_family != INVALID_TYPE_FAMILY) if (root_node.lhs.type_family != INVALID_TYPE_FAMILY)
fun(symbolic_expression, root_idx, LHS_NODE_TYPE); fun(array_expression, root_idx, LHS_NODE_TYPE);
} }
//Self: //Self:
fun(symbolic_expression, root_idx, PARENT_NODE_TYPE); fun(array_expression, root_idx, PARENT_NODE_TYPE);
//Rhs: //Rhs:
if (recurse && root_node.rhs.type_family!=INVALID_TYPE_FAMILY) if (recurse && root_node.rhs.type_family!=INVALID_TYPE_FAMILY)
{ {
if (root_node.rhs.type_family==COMPOSITE_OPERATOR_FAMILY) if (root_node.rhs.type_family==COMPOSITE_OPERATOR_FAMILY)
traverse(symbolic_expression, root_node.rhs.node_index, fun, inspect); traverse(array_expression, root_node.rhs.node_index, fun, inspect);
if (root_node.rhs.type_family != INVALID_TYPE_FAMILY) if (root_node.rhs.type_family != INVALID_TYPE_FAMILY)
fun(symbolic_expression, root_idx, RHS_NODE_TYPE); fun(array_expression, root_idx, RHS_NODE_TYPE);
} }
fun.call_after_expansion(symbolic_expression, root_idx); fun.call_after_expansion(array_expression, root_idx);
} }
class filter_fun : public traversal_functor class filter_fun : public traversal_functor
{ {
public: public:
typedef bool (*pred_t)(symbolic_expression_node const & node); typedef bool (*pred_t)(array_expression::node const & node);
filter_fun(pred_t pred, std::vector<size_t> & out); filter_fun(pred_t pred, std::vector<size_t> & out);
void operator()(atidlas::symbolic_expression const & symbolic_expression, size_t root_idx, leaf_t) const; void operator()(atidlas::array_expression const & array_expression, size_t root_idx, leaf_t) const;
private: private:
pred_t pred_; pred_t pred_;
std::vector<size_t> & out_; std::vector<size_t> & out_;
@@ -79,22 +79,22 @@ private:
class filter_elements_fun : public traversal_functor class filter_elements_fun : public traversal_functor
{ {
public: public:
filter_elements_fun(symbolic_expression_node_subtype subtype, std::vector<lhs_rhs_element> & out); filter_elements_fun(array_expression_node_subtype subtype, std::vector<lhs_rhs_element> & out);
void operator()(atidlas::symbolic_expression const & symbolic_expression, size_t root_idx, leaf_t) const; void operator()(atidlas::array_expression const & array_expression, size_t root_idx, leaf_t) const;
private: private:
symbolic_expression_node_subtype subtype_; array_expression_node_subtype subtype_;
std::vector<lhs_rhs_element> & out_; std::vector<lhs_rhs_element> & out_;
}; };
std::vector<size_t> filter_nodes(bool (*pred)(symbolic_expression_node const & node), std::vector<size_t> filter_nodes(bool (*pred)(array_expression::node const & node),
atidlas::symbolic_expression const & symbolic_expression, atidlas::array_expression const & array_expression,
bool inspect); bool inspect);
std::vector<lhs_rhs_element> filter_elements(symbolic_expression_node_subtype subtype, std::vector<lhs_rhs_element> filter_elements(array_expression_node_subtype subtype,
atidlas::symbolic_expression const & symbolic_expression); atidlas::array_expression const & array_expression);
const char * evaluate(operation_node_type type); const char * evaluate(operation_node_type type);
/** @brief functor for generating the expression string from a symbolic_expression */ /** @brief functor for generating the expression string from a array_expression */
class evaluate_expression_traversal: public traversal_functor class evaluate_expression_traversal: public traversal_functor
{ {
private: private:
@@ -104,24 +104,24 @@ private:
public: public:
evaluate_expression_traversal(std::map<std::string, std::string> const & accessors, std::string & str, mapping_type const & mapping); evaluate_expression_traversal(std::map<std::string, std::string> const & accessors, std::string & str, mapping_type const & mapping);
void call_before_expansion(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx) const; void call_before_expansion(atidlas::array_expression const & array_expression, int_t root_idx) const;
void call_after_expansion(symbolic_expression const & /*symbolic_expression*/, int_t /*root_idx*/) const; void call_after_expansion(array_expression const & /*array_expression*/, int_t /*root_idx*/) const;
void operator()(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf) const; void operator()(atidlas::array_expression const & array_expression, int_t root_idx, leaf_t leaf) const;
}; };
std::string evaluate(leaf_t leaf, std::map<std::string, std::string> const & accessors, std::string evaluate(leaf_t leaf, std::map<std::string, std::string> const & accessors,
atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, mapping_type const & mapping); atidlas::array_expression const & array_expression, int_t root_idx, mapping_type const & mapping);
void evaluate(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors, void evaluate(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors,
symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings); array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings);
/** @brief functor for fetching or writing-back the elements in a symbolic_expression */ /** @brief functor for fetching or writing-back the elements in a array_expression */
class process_traversal : public traversal_functor class process_traversal : public traversal_functor
{ {
public: public:
process_traversal(std::map<std::string, std::string> const & accessors, kernel_generation_stream & stream, process_traversal(std::map<std::string, std::string> const & accessors, kernel_generation_stream & stream,
mapping_type const & mapping, std::set<std::string> & already_processed); mapping_type const & mapping, std::set<std::string> & already_processed);
void operator()(symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf) const; void operator()(array_expression const & array_expression, int_t root_idx, leaf_t leaf) const;
private: private:
std::map<std::string, std::string> accessors_; std::map<std::string, std::string> accessors_;
kernel_generation_stream & stream_; kernel_generation_stream & stream_;
@@ -130,21 +130,21 @@ private:
}; };
void process(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors, void process(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors,
atidlas::symbolic_expression const & symbolic_expression, size_t root_idx, mapping_type const & mapping, std::set<std::string> & already_processed); atidlas::array_expression const & array_expression, size_t root_idx, mapping_type const & mapping, std::set<std::string> & already_processed);
void process(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors, void process(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors,
symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings); array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings);
class symbolic_expression_representation_functor : public traversal_functor{ class array_expression_representation_functor : public traversal_functor{
private: private:
static void append_id(char * & ptr, unsigned int val); static void append_id(char * & ptr, unsigned int val);
void append(cl_mem h, numeric_type dtype, char prefix) const; void append(cl_mem h, numeric_type dtype, char prefix) const;
void append(lhs_rhs_element const & lhs_rhs) const; void append(lhs_rhs_element const & lhs_rhs) const;
public: public:
symbolic_expression_representation_functor(symbolic_binder & binder, char *& ptr); array_expression_representation_functor(symbolic_binder & binder, char *& ptr);
void append(char*& p, const char * str) const; void append(char*& p, const char * str) const;
void operator()(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf_t) const; void operator()(atidlas::array_expression const & array_expression, int_t root_idx, leaf_t leaf_t) const;
private: private:
symbolic_binder & binder_; symbolic_binder & binder_;
char *& ptr_; char *& ptr_;

46
include/atidlas/backend/templates/base.h
View File

@@ -69,13 +69,13 @@ public:
protected: protected:
/** @brief Functor to map the symbolic_expressions to the types defined in mapped_objects.hpp */ /** @brief Functor to map the array_expressions to the types defined in mapped_objects.hpp */
class map_functor : public traversal_functor class map_functor : public traversal_functor
{ {
/** @brief Accessor for the numeric type */ /** @brief Accessor for the numeric type */
numeric_type get_numeric_type(atidlas::symbolic_expression const * symbolic_expression, int_t root_idx) const; numeric_type get_numeric_type(atidlas::array_expression const * array_expression, int_t root_idx) const;
/** @brief Creates a binary leaf */ /** @brief Creates a binary leaf */
template<class T> tools::shared_ptr<mapped_object> binary_leaf(atidlas::symbolic_expression const * symbolic_expression, int_t root_idx, mapping_type const * mapping) const; template<class T> tools::shared_ptr<mapped_object> binary_leaf(atidlas::array_expression const * array_expression, int_t root_idx, mapping_type const * mapping) const;
/** @brief Creates a value scalar mapping */ /** @brief Creates a value scalar mapping */
tools::shared_ptr<mapped_object> create(numeric_type dtype, values_holder) const; tools::shared_ptr<mapped_object> create(numeric_type dtype, values_holder) const;
/** @brief Creates a vector mapping */ /** @brief Creates a vector mapping */
@@ -87,7 +87,7 @@ protected:
public: public:
map_functor(symbolic_binder & binder, mapping_type & mapping); map_functor(symbolic_binder & binder, mapping_type & mapping);
/** @brief Functor for traversing the tree */ /** @brief Functor for traversing the tree */
void operator()(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf_t) const; void operator()(atidlas::array_expression const & array_expression, int_t root_idx, leaf_t leaf_t) const;
private: private:
symbolic_binder & binder_; symbolic_binder & binder_;
mapping_type & mapping_; mapping_type & mapping_;
@@ -105,7 +105,7 @@ protected:
void set_arguments(repeat_infos const & i) const; void set_arguments(repeat_infos const & i) const;
void set_arguments(lhs_rhs_element const & lhs_rhs) const; void set_arguments(lhs_rhs_element const & lhs_rhs) const;
void operator()(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf_t) const; void operator()(atidlas::array_expression const & array_expression, int_t root_idx, leaf_t leaf_t) const;
private: private:
symbolic_binder & binder_; symbolic_binder & binder_;
unsigned int & current_arg_; unsigned int & current_arg_;
@@ -130,41 +130,41 @@ protected:
size_t root_idx, leaf_t leaf); size_t root_idx, leaf_t leaf);
static std::string neutral_element(op_element const & op); static std::string neutral_element(op_element const & op);
static std::string generate_arguments(std::vector<mapping_type> const & mappings, std::map<std::string, std::string> const & accessors, static std::string generate_arguments(std::vector<mapping_type> const & mappings, std::map<std::string, std::string> const & accessors,
symbolic_expressions_container const & symbolic_expressions); array_expressions_container const & array_expressions);
static std::string generate_arguments(std::string const & data_type, std::vector<mapping_type> const & mappings, static std::string generate_arguments(std::string const & data_type, std::vector<mapping_type> const & mappings,
symbolic_expressions_container const & symbolic_expressions); array_expressions_container const & array_expressions);
static void fill_kernel_name(char * ptr, unsigned int label, const char * suffix); static void fill_kernel_name(char * ptr, unsigned int label, const char * suffix);
static bool is_node_trans(symbolic_expression::container_type const & array, size_t root_idx, leaf_t leaf_type); static bool is_node_trans(array_expression::container_type const & array, size_t root_idx, leaf_t leaf_type);
static std::string append_simd_suffix(std::string const & str, unsigned int i); static std::string append_simd_suffix(std::string const & str, unsigned int i);
static bool is_strided(symbolic_expression_node const & node); static bool is_strided(array_expression::node const & node);
static int_t vector_size(symbolic_expression_node const & node); static int_t vector_size(array_expression::node const & node);
static std::pair<int_t, int_t> matrix_size(symbolic_expression_node const & node); static std::pair<int_t, int_t> matrix_size(array_expression::node const & node);
static unsigned int align(unsigned int to_round, unsigned int base); static unsigned int align(unsigned int to_round, unsigned int base);
static bool is_reduction(symbolic_expression_node const & node); static bool is_reduction(array_expression::node const & node);
static bool is_index_reduction(op_element const & op); static bool is_index_reduction(op_element const & op);
tools::shared_ptr<symbolic_binder> make_binder(); tools::shared_ptr<symbolic_binder> make_binder();
static std::string vstore(unsigned int simd_width, std::string const & value, std::string const & offset, std::string const & ptr); static std::string vstore(unsigned int simd_width, std::string const & value, std::string const & offset, std::string const & ptr);
static std::string vload(unsigned int simd_width, std::string const & offset, std::string const & ptr); static std::string vload(unsigned int simd_width, std::string const & offset, std::string const & ptr);
static std::string append_width(std::string const & str, unsigned int width); static std::string append_width(std::string const & str, unsigned int width);
static bool requires_fallback(symbolic_expressions_container const & symbolic_expressions); static bool requires_fallback(array_expressions_container const & array_expressions);
void set_arguments(symbolic_expressions_container const & symbolic_expressions, cl::Kernel & kernel, unsigned int & current_arg); void set_arguments(array_expressions_container const & array_expressions, cl::Kernel & kernel, unsigned int & current_arg);
private: private:
virtual std::vector<std::string> generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mapping) const = 0; virtual std::vector<std::string> generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mapping) const = 0;
public: public:
base(binding_policy_t binding_policy); base(binding_policy_t binding_policy);
virtual unsigned int lmem_usage(symbolic_expressions_container const &) const; virtual unsigned int lmem_usage(array_expressions_container const &) const;
virtual unsigned int registers_usage(symbolic_expressions_container const &) const; virtual unsigned int registers_usage(array_expressions_container const &) const;
virtual std::vector<int_t> input_sizes(symbolic_expressions_container const & symbolic_expressions) = 0; virtual std::vector<int_t> input_sizes(array_expressions_container const & array_expressions) = 0;
virtual ~base(); virtual ~base();
std::vector<std::string> generate(unsigned int label, symbolic_expressions_container const & symbolic_expressions, cl::Device const & device); std::vector<std::string> generate(unsigned int label, array_expressions_container const & array_expressions, cl::Device const & device);
virtual int check_invalid(symbolic_expressions_container const & symbolic_expressions, cl::Device const & device) const = 0; virtual int check_invalid(array_expressions_container const & array_expressions, cl::Device const & device) const = 0;
virtual void enqueue(cl::CommandQueue & queue, virtual void enqueue(cl::CommandQueue & queue,
std::vector<cl_ext::lazy_compiler> & programs, std::vector<cl_ext::lazy_compiler> & programs,
unsigned int label, symbolic_expressions_container const & symbolic_expressions) = 0; unsigned int label, array_expressions_container const & array_expressions) = 0;
virtual tools::shared_ptr<base> clone() const = 0; virtual tools::shared_ptr<base> clone() const = 0;
private: private:
binding_policy_t binding_policy_; binding_policy_t binding_policy_;
@@ -175,7 +175,7 @@ template<class TemplateType, class ParametersType>
class base_impl : public base class base_impl : public base
{ {
private: private:
virtual int check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const; virtual int check_invalid_impl(cl::Device const &, array_expressions_container const &) const;
public: public:
typedef ParametersType parameters_type; typedef ParametersType parameters_type;
base_impl(parameters_type const & parameters, binding_policy_t binding_policy); base_impl(parameters_type const & parameters, binding_policy_t binding_policy);
@@ -183,7 +183,7 @@ public:
int_t local_size_1() const; int_t local_size_1() const;
tools::shared_ptr<base> clone() const; tools::shared_ptr<base> clone() const;
/** @brief returns whether or not the profile has undefined behavior on particular device */ /** @brief returns whether or not the profile has undefined behavior on particular device */
int check_invalid(symbolic_expressions_container const & symbolic_expressions, cl::Device const & device) const; int check_invalid(array_expressions_container const & array_expressions, cl::Device const & device) const;
protected: protected:
parameters_type p_; parameters_type p_;
binding_policy_t binding_policy_; binding_policy_t binding_policy_;

10
include/atidlas/backend/templates/maxpy.h
View File

@@ -20,14 +20,14 @@ public:
class maxpy : public base_impl<maxpy, maxpy_parameters> class maxpy : public base_impl<maxpy, maxpy_parameters>
{ {
private: private:
int check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const; int check_invalid_impl(cl::Device const &, array_expressions_container const &) const;
std::string generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width) const; std::string generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width) const;
std::vector<std::string> generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const; std::vector<std::string> generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const;
public: public:
maxpy(parameters_type const & parameters, binding_policy_t binding_policy = BIND_ALL_UNIQUE); maxpy(parameters_type const & parameters, binding_policy_t binding_policy = BIND_ALL_UNIQUE);
maxpy(unsigned int simd, unsigned int ls1, unsigned int ls2, unsigned int ng1, unsigned int ng2, fetching_policy_type fetch, binding_policy_t bind = BIND_ALL_UNIQUE); maxpy(unsigned int simd, unsigned int ls1, unsigned int ls2, unsigned int ng1, unsigned int ng2, fetching_policy_type fetch, binding_policy_t bind = BIND_ALL_UNIQUE);
std::vector<int_t> input_sizes(symbolic_expressions_container const & symbolic_expressions); std::vector<int_t> input_sizes(array_expressions_container const & array_expressions);
void enqueue(cl::CommandQueue & queue, std::vector<cl_ext::lazy_compiler> & programs, unsigned int label, symbolic_expressions_container const & symbolic_expressions); void enqueue(cl::CommandQueue & queue, std::vector<cl_ext::lazy_compiler> & programs, unsigned int label, array_expressions_container const & array_expressions);
}; };
} }

16
include/atidlas/backend/templates/mproduct.h
View File

@@ -33,25 +33,25 @@ struct mproduct_parameters : public base::parameters_type
class mproduct : public base_impl<mproduct, mproduct_parameters> class mproduct : public base_impl<mproduct, mproduct_parameters>
{ {
private: private:
unsigned int lmem_usage(symbolic_expressions_container const & symbolic_expressions) const; unsigned int lmem_usage(array_expressions_container const & array_expressions) const;
unsigned int registers_usage(symbolic_expressions_container const & symbolic_expressions) const; unsigned int registers_usage(array_expressions_container const & array_expressions) const;
int check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const; int check_invalid_impl(cl::Device const &, array_expressions_container const &) const;
std::string generate_impl(unsigned int label, const char * id, const symbolic_expressions_container &symbolic_expressions, const std::vector<mapping_type> &, bool fallback) const; std::string generate_impl(unsigned int label, const char * id, const array_expressions_container &array_expressions, const std::vector<mapping_type> &, bool fallback) const;
std::vector<std::string> generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const; std::vector<std::string> generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const;
void enqueue_block(cl::CommandQueue & queue, int_t M, int_t N, int_t K, void enqueue_block(cl::CommandQueue & queue, int_t M, int_t N, int_t K,
array_infos const & A, array_infos const & B, array_infos const & C, array_infos const & A, array_infos const & B, array_infos const & C,
value_scalar const & alpha, value_scalar const & beta, value_scalar const & alpha, value_scalar const & beta,
std::vector<cl_ext::lazy_compiler> & programs, unsigned int label, int id); std::vector<cl_ext::lazy_compiler> & programs, unsigned int label, int id);
array_infos create_slice(array_infos & M, int_t s0_0, int_t s0_1, int_t s1_0, int_t s1_1, bool swap); array_infos create_slice(array_infos & M, int_t s0_0, int_t s0_1, int_t s1_0, int_t s1_1, bool swap);
std::vector<int_t> infos(symbolic_expressions_container const & symbolic_expressions, std::vector<int_t> infos(array_expressions_container const & array_expressions,
lhs_rhs_element & C, lhs_rhs_element & A, lhs_rhs_element & B); lhs_rhs_element & C, lhs_rhs_element & A, lhs_rhs_element & B);
public: public:
mproduct(mproduct::parameters_type const & parameters, char A_trans, char B_trans); mproduct(mproduct::parameters_type const & parameters, char A_trans, char B_trans);
std::vector<int_t> input_sizes(symbolic_expressions_container const & symbolic_expressions); std::vector<int_t> input_sizes(array_expressions_container const & array_expressions);
void enqueue(cl::CommandQueue & queue, void enqueue(cl::CommandQueue & queue,
std::vector<cl_ext::lazy_compiler> & programs, std::vector<cl_ext::lazy_compiler> & programs,
unsigned int label, unsigned int label,
symbolic_expressions_container const & symbolic_expressions); array_expressions_container const & array_expressions);
private: private:
const char A_trans_; const char A_trans_;

10
include/atidlas/backend/templates/mreduction.h
View File

@@ -29,13 +29,13 @@ protected:
}; };
mreduction(mreduction::parameters_type const & , reduction_type, binding_policy_t); mreduction(mreduction::parameters_type const & , reduction_type, binding_policy_t);
private: private:
virtual int check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const; virtual int check_invalid_impl(cl::Device const &, array_expressions_container const &) const;
unsigned int lmem_usage() const; unsigned int lmem_usage() const;
std::string generate_impl(unsigned int, symbolic_expressions_container const &, std::vector<mapping_type> const &, unsigned int, std::vector<mapped_mreduction*> const &) const; std::string generate_impl(unsigned int, array_expressions_container const &, std::vector<mapping_type> const &, unsigned int, std::vector<mapped_mreduction*> const &) const;
std::vector<std::string> generate_impl(unsigned int, symbolic_expressions_container const &, std::vector<mapping_type> const &) const; std::vector<std::string> generate_impl(unsigned int, array_expressions_container const &, std::vector<mapping_type> const &) const;
public: public:
virtual std::vector<int_t> input_sizes(symbolic_expressions_container const & symbolic_expressions); virtual std::vector<int_t> input_sizes(array_expressions_container const & array_expressions);
void enqueue(cl::CommandQueue & queue,std::vector<cl_ext::lazy_compiler> & programs,unsigned int label, symbolic_expressions_container const & symbolic_expressions); void enqueue(cl::CommandQueue & queue,std::vector<cl_ext::lazy_compiler> & programs,unsigned int label, array_expressions_container const & array_expressions);
private: private:
reduction_type reduction_type_; reduction_type reduction_type_;
}; };

12
include/atidlas/backend/templates/reduction.h
View File

@@ -18,21 +18,21 @@ struct reduction_parameters : public base::parameters_type
class reduction : public base_impl<reduction, reduction_parameters> class reduction : public base_impl<reduction, reduction_parameters>
{ {
private: private:
unsigned int lmem_usage(symbolic_expressions_container const & symbolic_expressions) const; unsigned int lmem_usage(array_expressions_container const & array_expressions) const;
int check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const; int check_invalid_impl(cl::Device const &, array_expressions_container const &) const;
inline void reduce_1d_local_memory(kernel_generation_stream & stream, unsigned int size, std::vector<mapped_scalar_reduction*> exprs, inline void reduce_1d_local_memory(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; std::string const & buf_str, std::string const & buf_value_str) const;
std::string generate_impl(unsigned int label, const char * type, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width) const; std::string generate_impl(unsigned int label, const char * type, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width) const;
std::vector<std::string> generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const; std::vector<std::string> generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const;
public: public:
reduction(reduction::parameters_type const & parameters, binding_policy_t binding_policy = BIND_ALL_UNIQUE); reduction(reduction::parameters_type const & parameters, binding_policy_t binding_policy = BIND_ALL_UNIQUE);
reduction(unsigned int simd, unsigned int ls, unsigned int ng, fetching_policy_type fetch, binding_policy_t bind = BIND_ALL_UNIQUE); reduction(unsigned int simd, unsigned int ls, unsigned int ng, fetching_policy_type fetch, binding_policy_t bind = BIND_ALL_UNIQUE);
std::vector<int_t> input_sizes(symbolic_expressions_container const & symbolic_expressions); std::vector<int_t> input_sizes(array_expressions_container const & array_expressions);
void enqueue(cl::CommandQueue & queue, void enqueue(cl::CommandQueue & queue,
std::vector<cl_ext::lazy_compiler> & programs, std::vector<cl_ext::lazy_compiler> & programs,
unsigned int label, unsigned int label,
symbolic_expressions_container const & symbolic_expressions); array_expressions_container const & array_expressions);
private: private:
std::vector< cl::Buffer > tmp_; std::vector< cl::Buffer > tmp_;
std::vector< cl::Buffer > tmpidx_; std::vector< cl::Buffer > tmpidx_;

8
include/atidlas/backend/templates/vaxpy.h
View File

@@ -17,14 +17,14 @@ public:
class vaxpy : public base_impl<vaxpy, vaxpy_parameters> class vaxpy : public base_impl<vaxpy, vaxpy_parameters>
{ {
private: private:
virtual int check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const; virtual int check_invalid_impl(cl::Device const &, array_expressions_container const &) const;
std::vector<std::string> generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const; std::vector<std::string> generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const;
public: public:
vaxpy(vaxpy::parameters_type const & parameters, binding_policy_t binding_policy = BIND_ALL_UNIQUE); vaxpy(vaxpy::parameters_type const & parameters, binding_policy_t binding_policy = BIND_ALL_UNIQUE);
vaxpy(unsigned int _simd_width, unsigned int _group_size, unsigned int _num_groups, fetching_policy_type _fetching_policy, binding_policy_t binding_policy = BIND_ALL_UNIQUE); vaxpy(unsigned int _simd_width, unsigned int _group_size, unsigned int _num_groups, fetching_policy_type _fetching_policy, binding_policy_t binding_policy = BIND_ALL_UNIQUE);
std::vector<int_t> input_sizes(symbolic_expressions_container const & symbolic_expressions); std::vector<int_t> input_sizes(array_expressions_container const & array_expressions);
void enqueue(cl::CommandQueue & queue, std::vector<cl_ext::lazy_compiler> & programs, void enqueue(cl::CommandQueue & queue, std::vector<cl_ext::lazy_compiler> & programs,
unsigned int label, symbolic_expressions_container const & symbolic_expressions); unsigned int label, array_expressions_container const & array_expressions);
}; };
} }

8
include/atidlas/model/model.h
View File

@@ -20,16 +20,16 @@ namespace atidlas
private: private:
std::string define_extension(std::string const & extensions, std::string const & ext); std::string define_extension(std::string const & extensions, std::string const & ext);
inline void fill_program_name(char* program_name, symbolic_expressions_container const & symbolic_expressions, binding_policy_t binding_policy); inline void fill_program_name(char* program_name, array_expressions_container const & array_expressions, binding_policy_t binding_policy);
std::vector<cl_ext::lazy_compiler>& init(symbolic_expressions_container const & symbolic_expressions, cl::Context const & context, cl::Device const & device, bool force_recompilation); std::vector<cl_ext::lazy_compiler>& init(array_expressions_container const & array_expressions, cl::Context const & context, cl::Device const & device, bool force_recompilation);
public: public:
model(predictors::random_forest const &, std::vector< tools::shared_ptr<base> > const &, cl::CommandQueue &); model(predictors::random_forest const &, std::vector< tools::shared_ptr<base> > const &, cl::CommandQueue &);
model(std::vector< tools::shared_ptr<base> > const &, cl::CommandQueue &); model(std::vector< tools::shared_ptr<base> > const &, cl::CommandQueue &);
model(base const &, cl::CommandQueue &); model(base const &, cl::CommandQueue &);
void execute(symbolic_expressions_container const &, bool bypass_predictor = false, bool force_recompilation = false); void execute(array_expressions_container const &, bool bypass_predictor = false, bool force_recompilation = false);
void tune(symbolic_expressions_container const &); void tune(array_expressions_container const &);
templates_container const & templates() const; templates_container const & templates() const;
private: private:

4
include/atidlas/symbolic/execute.h
View File

@@ -8,8 +8,8 @@
namespace atidlas namespace atidlas
{ {
/** @brief Executes a symbolic_expression on the given queue for the given models map*/ /** @brief Executes a array_expression on the given queue for the given models map*/
void execute(symbolic_expression &, model_map_t &); void execute(array_expression &, model_map_t &);
} }

107
include/atidlas/symbolic/expression.h
View File

@@ -1,5 +1,5 @@
#ifndef _ATIDLAS_SYMBOLIC_EXPRESSION_H #ifndef _ATIDLAS_array_expression_H
#define _ATIDLAS_SYMBOLIC_EXPRESSION_H #define _ATIDLAS_array_expression_H
#include <vector> #include <vector>
#include <list> #include <list>
@@ -117,8 +117,8 @@ enum operation_node_type
OPERATOR_PAIR_TYPE OPERATOR_PAIR_TYPE
}; };
/** @brief Groups the type of a node in the symbolic_expression tree. Used for faster dispatching */ /** @brief Groups the type of a node in the array_expression tree. Used for faster dispatching */
enum symbolic_expression_node_type_family enum array_expression_node_type_family
{ {
INVALID_TYPE_FAMILY = 0, INVALID_TYPE_FAMILY = 0,
COMPOSITE_OPERATOR_FAMILY, COMPOSITE_OPERATOR_FAMILY,
@@ -127,8 +127,8 @@ enum symbolic_expression_node_type_family
INFOS_TYPE_FAMILY INFOS_TYPE_FAMILY
}; };
/** @brief Encodes the type of a node in the symbolic_expression tree. */ /** @brief Encodes the type of a node in the array_expression tree. */
enum symbolic_expression_node_subtype enum array_expression_node_subtype
{ {
INVALID_SUBTYPE = 0, INVALID_SUBTYPE = 0,
VALUE_SCALAR_TYPE, VALUE_SCALAR_TYPE,
@@ -136,6 +136,15 @@ enum symbolic_expression_node_subtype
REPEAT_INFOS_TYPE REPEAT_INFOS_TYPE
}; };
struct op_element
{
op_element();
op_element(operation_node_type_family const & _type_family, operation_node_type const & _type);
operation_node_type_family type_family;
operation_node_type type;
};
struct array_infos struct array_infos
{ {
numeric_type dtype; numeric_type dtype;
@@ -149,17 +158,11 @@ struct array_infos
int_t ld; int_t ld;
}; };
void fill(array const & a, array_infos& i);
struct lhs_rhs_element struct lhs_rhs_element
{ {
lhs_rhs_element(); lhs_rhs_element();
lhs_rhs_element(unsigned int _node_index); array_expression_node_type_family type_family;
lhs_rhs_element(atidlas::array const & x); array_expression_node_subtype subtype;
lhs_rhs_element(value_scalar const & x);
lhs_rhs_element(repeat_infos const & x);
symbolic_expression_node_type_family type_family;
symbolic_expression_node_subtype subtype;
numeric_type dtype; numeric_type dtype;
union union
{ {
@@ -171,73 +174,65 @@ struct lhs_rhs_element
cl::Buffer memory_; cl::Buffer memory_;
}; };
struct invalid_node{};
struct op_element void fill(lhs_rhs_element &x, invalid_node);
void fill(lhs_rhs_element & x, unsigned int node_index);
void fill(lhs_rhs_element & x, array const & a);
void fill(lhs_rhs_element & x, value_scalar const & v);
void fill(lhs_rhs_element & x, repeat_infos const & r);
class array_expression
{ {
op_element(operation_node_type_family const & _type_family, operation_node_type const & _type); public:
operation_node_type_family type_family; struct node
operation_node_type type;
};
struct symbolic_expression_node
{ {
symbolic_expression_node(lhs_rhs_element const & _lhs, op_element const& _op, lhs_rhs_element const & _rhs);
lhs_rhs_element lhs; lhs_rhs_element lhs;
op_element op; op_element op;
lhs_rhs_element rhs; lhs_rhs_element rhs;
}; };
class symbolic_expression typedef std::vector<node> container_type;
{
public: public:
typedef symbolic_expression_node value_type; template<class LT, class RT>
typedef std::vector<value_type> container_type; array_expression(LT const & lhs, RT const & rhs, op_element const & op, cl::Context const & ctx, numeric_type const & dtype, size4 const & shape);
template<class RT>
symbolic_expression(lhs_rhs_element const & lhs, lhs_rhs_element const & rhs, op_element const & op, cl::Context const & context, numeric_type const & dtype); array_expression(array_expression const & lhs, RT const & rhs, op_element const & op, numeric_type const & dtype, size4 const & shape);
symbolic_expression(symbolic_expression const & lhs, lhs_rhs_element const & rhs, op_element const & op, numeric_type const & dtype); template<class LT>
symbolic_expression(lhs_rhs_element const & lhs, symbolic_expression const & rhs, op_element const & op, numeric_type const & dtype); array_expression(LT const & lhs, array_expression const & rhs, op_element const & op, numeric_type const & dtype, size4 const & shape);
symbolic_expression(symbolic_expression const & lhs, symbolic_expression const & rhs, op_element const & op, numeric_type const & dtype); array_expression(array_expression const & lhs, array_expression const & rhs, op_element const & op, numeric_type const & dtype, size4 const & shape);
size4 shape() const;
array_expression& reshape(int_t size1, int_t size2=1);
int_t nshape() const;
container_type & tree(); container_type & tree();
container_type const & tree() const; container_type const & tree() const;
std::size_t root() const; std::size_t root() const;
cl::Context const & context() const; cl::Context const & context() const;
numeric_type const & dtype() const; numeric_type const & dtype() const;
protected:
container_type tree_;
std::size_t root_;
cl::Context context_;
numeric_type dtype_;
};
struct array_expression: public symbolic_expression
{
array_expression(lhs_rhs_element const & lhs, lhs_rhs_element const & rhs, op_element const & op, cl::Context const & ctx, numeric_type const & dtype, size4 shape);
array_expression(symbolic_expression const & lhs, lhs_rhs_element const & rhs, op_element const & op, numeric_type const & dtype, size4 shape);
array_expression(lhs_rhs_element const & lhs, symbolic_expression const & rhs, op_element const & op, numeric_type const & dtype, size4 shape);
array_expression(symbolic_expression const & lhs, symbolic_expression const & rhs, op_element const & op, numeric_type const & dtype, size4 shape);
size4 shape() const;
array_expression& reshape(int_t size1, int_t size2=1);
int_t nshape() const;
array_expression operator-(); array_expression operator-();
array_expression operator!(); array_expression operator!();
private: private:
container_type tree_;
std::size_t root_;
cl::Context context_;
numeric_type dtype_;
size4 shape_; size4 shape_;
}; };
class symbolic_expressions_container class array_expressions_container
{ {
private: private:
tools::shared_ptr<symbolic_expression> create(symbolic_expression const & s); tools::shared_ptr<array_expression> create(array_expression const & s);
public: public:
typedef std::list<tools::shared_ptr<symbolic_expression> > data_type; typedef std::list<tools::shared_ptr<array_expression> > data_type;
enum order_type { SEQUENTIAL, INDEPENDENT }; enum order_type { SEQUENTIAL, INDEPENDENT };
symbolic_expressions_container(symbolic_expression const & s0); array_expressions_container(array_expression const & s0);
symbolic_expressions_container(order_type order, symbolic_expression const & s0, symbolic_expression const & s1); array_expressions_container(order_type order, array_expression const & s0, array_expression const & s1);
symbolic_expressions_container(data_type const & data, order_type order); array_expressions_container(data_type const & data, order_type order);
data_type const & data() const; data_type const & data() const;
cl::Context const & context() const; cl::Context const & context() const;
@@ -247,8 +242,8 @@ private:
order_type order_; order_type order_;
}; };
symbolic_expression_node const & lhs_most(symbolic_expression::container_type const & array, symbolic_expression_node const & init); array_expression::node const & lhs_most(array_expression::container_type const & array, array_expression::node const & init);
symbolic_expression_node const & lhs_most(symbolic_expression::container_type const & array, size_t root); array_expression::node const & lhs_most(array_expression::container_type const & array, size_t root);
} }

6
include/atidlas/symbolic/io.h
View File

@@ -7,10 +7,10 @@
namespace atidlas namespace atidlas
{ {
std::string to_string(symbolic_expression_node_subtype const & f); std::string to_string(array_expression_node_subtype const & f);
std::string to_string(lhs_rhs_element const & e); std::string to_string(lhs_rhs_element const & e);
std::ostream & operator<<(std::ostream & os, symbolic_expression_node const & s_node); std::ostream & operator<<(std::ostream & os, array_expression::node const & s_node);
std::string to_string(atidlas::symbolic_expression const & s); std::string to_string(atidlas::array_expression const & s);
} }

1
lib/CMakeLists.txt
View File

@@ -1,4 +1,5 @@
set(CMAKE_BUILD_TYPE RELEASE) set(CMAKE_BUILD_TYPE RELEASE)
set(CMAKE_CXX_FLAGS_RELEASE "-O2")
file( GLOB_RECURSE LIBATIDLAS_SRC *.cpp) file( GLOB_RECURSE LIBATIDLAS_SRC *.cpp)
include_directories(${CMAKE_CURRENT_SOURCE_DIR}) include_directories(${CMAKE_CURRENT_SOURCE_DIR})

46
lib/array.cpp
View File

@@ -173,10 +173,10 @@ INSTANTIATE(cl_double);
#undef INSTANTIATE #undef INSTANTIATE
array_expression array::operator-() array_expression array::operator-()
{ return array_expression(*this, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_SUB_TYPE), context_, dtype_, shape_); } { return array_expression(*this, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_SUB_TYPE), context_, dtype_, shape_); }
array_expression array::operator!() array_expression array::operator!()
{ return array_expression(*this, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_NEGATE_TYPE), context_, INT_TYPE, shape_); } { return array_expression(*this, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_NEGATE_TYPE), context_, INT_TYPE, shape_); }
// //
array & array::operator+=(value_scalar const & rhs) array & array::operator+=(value_scalar const & rhs)
@@ -461,10 +461,10 @@ array_expression outer(array const & x, array const & y)
//--------------------------------------- //---------------------------------------
#define DEFINE_ELEMENT_UNARY_OPERATOR(OP, OPNAME) \ #define DEFINE_ELEMENT_UNARY_OPERATOR(OP, OPNAME) \
array_expression OPNAME (array const & x) \ array_expression OPNAME (array const & x) \
{ return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OP), x.context(), x.dtype(), x.shape()); }\ { return array_expression(x, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OP), x.context(), x.dtype(), x.shape()); }\
\ \
array_expression OPNAME (array_expression const & x) \ array_expression OPNAME (array_expression const & x) \
{ return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OP), x.dtype(), x.shape()); } { return array_expression(x, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OP), x.dtype(), x.shape()); }
DEFINE_ELEMENT_UNARY_OPERATOR((x.dtype()==FLOAT_TYPE || x.dtype()==DOUBLE_TYPE)?OPERATOR_FABS_TYPE:OPERATOR_ABS_TYPE, abs) DEFINE_ELEMENT_UNARY_OPERATOR((x.dtype()==FLOAT_TYPE || x.dtype()==DOUBLE_TYPE)?OPERATOR_FABS_TYPE:OPERATOR_ABS_TYPE, abs)
DEFINE_ELEMENT_UNARY_OPERATOR(OPERATOR_ACOS_TYPE, acos) DEFINE_ELEMENT_UNARY_OPERATOR(OPERATOR_ACOS_TYPE, acos)
@@ -509,16 +509,16 @@ inline operation_node_type casted(numeric_type dtype)
} }
array_expression cast(array const & x, numeric_type dtype) array_expression cast(array const & x, numeric_type dtype)
{ return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, casted(dtype)), x.context(), dtype, x.shape()); } { return array_expression(x, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, casted(dtype)), x.context(), dtype, x.shape()); }
array_expression cast(array_expression const & x, numeric_type dtype) array_expression cast(array_expression const & x, numeric_type dtype)
{ return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, casted(dtype)), dtype, x.shape()); } { return array_expression(x, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, casted(dtype)), dtype, x.shape()); }
atidlas::array_expression eye(std::size_t M, std::size_t N, atidlas::numeric_type dtype, cl::Context ctx) atidlas::array_expression eye(std::size_t M, std::size_t N, atidlas::numeric_type dtype, cl::Context ctx)
{ return array_expression(value_scalar(1), value_scalar(0), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_VDIAG_TYPE), ctx, dtype, size4(M, N)); } { return array_expression(value_scalar(1), value_scalar(0), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_VDIAG_TYPE), ctx, dtype, size4(M, N)); }
atidlas::array_expression zeros(std::size_t M, std::size_t N, atidlas::numeric_type dtype, cl::Context ctx) atidlas::array_expression zeros(std::size_t M, std::size_t N, atidlas::numeric_type dtype, cl::Context ctx)
{ return array_expression(value_scalar(0), lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_ADD_TYPE), ctx, dtype, size4(M, N)); } { return array_expression(value_scalar(0), invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_ADD_TYPE), ctx, dtype, size4(M, N)); }
inline size4 flip(size4 const & shape) inline size4 flip(size4 const & shape)
{ return size4(shape._2, shape._1);} { return size4(shape._2, shape._1);}
@@ -527,10 +527,10 @@ inline size4 prod(size4 const & shape1, size4 const & shape2)
{ return size4(shape1._1*shape2._1, shape1._2*shape2._2);} { return size4(shape1._1*shape2._1, shape1._2*shape2._2);}
array_expression trans(array const & x) \ array_expression trans(array const & x) \
{ return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_TRANS_TYPE), x.context(), x.dtype(), flip(x.shape())); }\ { return array_expression(x, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_TRANS_TYPE), x.context(), x.dtype(), flip(x.shape())); }\
\ \
array_expression trans(array_expression const & x) \ array_expression trans(array_expression const & x) \
{ return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_TRANS_TYPE), x.dtype(), flip(x.shape())); } { return array_expression(x, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_TRANS_TYPE), x.dtype(), flip(x.shape())); }
array_expression repmat(array const & A, int_t const & rep1, int_t const & rep2) array_expression repmat(array const & A, int_t const & rep1, int_t const & rep2)
{ {
@@ -562,11 +562,11 @@ array_expression OPNAME(array const & x, int_t axis)\
if(axis < -1 || axis > x.nshape())\ if(axis < -1 || axis > x.nshape())\
throw std::out_of_range("The axis entry is out of bounds");\ throw std::out_of_range("The axis entry is out of bounds");\
else if(axis==-1)\ else if(axis==-1)\
return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_VECTOR_REDUCTION_TYPE_FAMILY, OP), x.context(), x.dtype(), size4(1));\ return array_expression(x, invalid_node(), op_element(OPERATOR_VECTOR_REDUCTION_TYPE_FAMILY, OP), x.context(), x.dtype(), size4(1));\
else if(axis==0)\ else if(axis==0)\
return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_ROWS_REDUCTION_TYPE_FAMILY, OP), x.context(), x.dtype(), size4(x.shape()._1));\ return array_expression(x, invalid_node(), op_element(OPERATOR_ROWS_REDUCTION_TYPE_FAMILY, OP), x.context(), x.dtype(), size4(x.shape()._1));\
else\ else\
return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_COLUMNS_REDUCTION_TYPE_FAMILY, OP), x.context(), x.dtype(), size4(x.shape()._2));\ return array_expression(x, invalid_node(), op_element(OPERATOR_COLUMNS_REDUCTION_TYPE_FAMILY, OP), x.context(), x.dtype(), size4(x.shape()._2));\
}\ }\
\ \
array_expression OPNAME(array_expression const & x, int_t axis)\ array_expression OPNAME(array_expression const & x, int_t axis)\
@@ -574,11 +574,11 @@ array_expression OPNAME(array_expression const & x, int_t axis)\
if(axis < -1 || axis > x.nshape())\ if(axis < -1 || axis > x.nshape())\
throw std::out_of_range("The axis entry is out of bounds");\ throw std::out_of_range("The axis entry is out of bounds");\
if(axis==-1)\ if(axis==-1)\
return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_VECTOR_REDUCTION_TYPE_FAMILY, OP), x.dtype(), size4(1));\ return array_expression(x, invalid_node(), op_element(OPERATOR_VECTOR_REDUCTION_TYPE_FAMILY, OP), x.dtype(), size4(1));\
else if(axis==0)\ else if(axis==0)\
return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_ROWS_REDUCTION_TYPE_FAMILY, OP), x.dtype(), size4(x.shape()._1));\ return array_expression(x, invalid_node(), op_element(OPERATOR_ROWS_REDUCTION_TYPE_FAMILY, OP), x.dtype(), size4(x.shape()._1));\
else\ else\
return array_expression(x, lhs_rhs_element(), op_element(OPERATOR_COLUMNS_REDUCTION_TYPE_FAMILY, OP), x.dtype(), size4(x.shape()._2));\ return array_expression(x, invalid_node(), op_element(OPERATOR_COLUMNS_REDUCTION_TYPE_FAMILY, OP), x.dtype(), size4(x.shape()._2));\
} }
DEFINE_REDUCTION(OPERATOR_ADD_TYPE, sum) DEFINE_REDUCTION(OPERATOR_ADD_TYPE, sum)
@@ -603,7 +603,7 @@ namespace detail
operation_node_type type = OPERATOR_MATRIX_PRODUCT_NN_TYPE; operation_node_type type = OPERATOR_MATRIX_PRODUCT_NN_TYPE;
size4 shape(A.shape()._1, B.shape()._2); size4 shape(A.shape()._1, B.shape()._2);
symbolic_expression_node & A_root = const_cast<symbolic_expression_node &>(A.tree()[A.root()]); array_expression::node & A_root = const_cast<array_expression::node &>(A.tree()[A.root()]);
bool A_trans = A_root.op.type==OPERATOR_TRANS_TYPE; bool A_trans = A_root.op.type==OPERATOR_TRANS_TYPE;
if(A_trans){ if(A_trans){
type = OPERATOR_MATRIX_PRODUCT_TN_TYPE; type = OPERATOR_MATRIX_PRODUCT_TN_TYPE;
@@ -611,7 +611,7 @@ namespace detail
} }
array_expression res(A, B, op_element(OPERATOR_MATRIX_PRODUCT_TYPE_FAMILY, type), A.dtype(), shape); array_expression res(A, B, op_element(OPERATOR_MATRIX_PRODUCT_TYPE_FAMILY, type), A.dtype(), shape);
symbolic_expression_node & res_root = const_cast<symbolic_expression_node &>(res.tree()[res.root()]); array_expression::node & res_root = const_cast<array_expression::node &>(res.tree()[res.root()]);
if(A_trans) res_root.lhs = A_root.lhs; if(A_trans) res_root.lhs = A_root.lhs;
return res; return res;
} }
@@ -621,14 +621,14 @@ namespace detail
operation_node_type type = OPERATOR_MATRIX_PRODUCT_NN_TYPE; operation_node_type type = OPERATOR_MATRIX_PRODUCT_NN_TYPE;
size4 shape(A.shape()._1, B.shape()._2); size4 shape(A.shape()._1, B.shape()._2);
symbolic_expression_node & B_root = const_cast<symbolic_expression_node &>(B.tree()[B.root()]); array_expression::node & B_root = const_cast<array_expression::node &>(B.tree()[B.root()]);
bool B_trans = B_root.op.type==OPERATOR_TRANS_TYPE; bool B_trans = B_root.op.type==OPERATOR_TRANS_TYPE;
if(B_trans){ if(B_trans){
type = OPERATOR_MATRIX_PRODUCT_NT_TYPE; type = OPERATOR_MATRIX_PRODUCT_NT_TYPE;
shape._2 = B.shape()._1; shape._2 = B.shape()._1;
} }
array_expression res(A, B, op_element(OPERATOR_MATRIX_PRODUCT_TYPE_FAMILY, type), A.dtype(), shape); array_expression res(A, B, op_element(OPERATOR_MATRIX_PRODUCT_TYPE_FAMILY, type), A.dtype(), shape);
symbolic_expression_node & res_root = const_cast<symbolic_expression_node &>(res.tree()[res.root()]); array_expression::node & res_root = const_cast<array_expression::node &>(res.tree()[res.root()]);
if(B_trans) res_root.rhs = B_root.lhs; if(B_trans) res_root.rhs = B_root.lhs;
return res; return res;
} }
@@ -636,8 +636,8 @@ namespace detail
array_expression matmatprod(array_expression const & A, array_expression const & B) array_expression matmatprod(array_expression const & A, array_expression const & B)
{ {
operation_node_type type = OPERATOR_MATRIX_PRODUCT_NN_TYPE; operation_node_type type = OPERATOR_MATRIX_PRODUCT_NN_TYPE;
symbolic_expression_node & A_root = const_cast<symbolic_expression_node &>(A.tree()[A.root()]); array_expression::node & A_root = const_cast<array_expression::node &>(A.tree()[A.root()]);
symbolic_expression_node & B_root = const_cast<symbolic_expression_node &>(B.tree()[B.root()]); array_expression::node & B_root = const_cast<array_expression::node &>(B.tree()[B.root()]);
size4 shape(A.shape()._1, B.shape()._2); size4 shape(A.shape()._1, B.shape()._2);
bool A_trans = A_root.op.type==OPERATOR_TRANS_TYPE; bool A_trans = A_root.op.type==OPERATOR_TRANS_TYPE;
@@ -650,7 +650,7 @@ namespace detail
else type = OPERATOR_MATRIX_PRODUCT_NN_TYPE; else type = OPERATOR_MATRIX_PRODUCT_NN_TYPE;
array_expression res(A, B, op_element(OPERATOR_MATRIX_PRODUCT_TYPE_FAMILY, type), A.dtype(), shape); array_expression res(A, B, op_element(OPERATOR_MATRIX_PRODUCT_TYPE_FAMILY, type), A.dtype(), shape);
symbolic_expression_node & res_root = const_cast<symbolic_expression_node &>(res.tree()[res.root()]); array_expression::node & res_root = const_cast<array_expression::node &>(res.tree()[res.root()]);
if(A_trans) res_root.lhs = A_root.lhs; if(A_trans) res_root.lhs = A_root.lhs;
if(B_trans) res_root.rhs = B_root.lhs; if(B_trans) res_root.rhs = B_root.lhs;
return res; return res;
@@ -669,7 +669,7 @@ namespace detail
{ {
int_t M = A.shape()._1; int_t M = A.shape()._1;
int_t N = A.shape()._2; int_t N = A.shape()._2;
symbolic_expression_node & A_root = const_cast<symbolic_expression_node &>(A.tree()[A.root()]); array_expression::node & A_root = const_cast<array_expression::node &>(A.tree()[A.root()]);
bool A_trans = A_root.op.type==OPERATOR_TRANS_TYPE; bool A_trans = A_root.op.type==OPERATOR_TRANS_TYPE;
if(A_trans) if(A_trans)
{ {

38
lib/backend/mapped_object.cpp
View File

@@ -49,8 +49,8 @@ void mapped_object::register_attribute(std::string & attribute, std::string cons
keywords_[key] = attribute; keywords_[key] = attribute;
} }
mapped_object::node_info::node_info(mapping_type const * _mapping, atidlas::symbolic_expression const * _symbolic_expression, int_t _root_idx) : mapped_object::node_info::node_info(mapping_type const * _mapping, atidlas::array_expression const * _array_expression, int_t _root_idx) :
mapping(_mapping), symbolic_expression(_symbolic_expression), root_idx(_root_idx) { } mapping(_mapping), array_expression(_array_expression), root_idx(_root_idx) { }
mapped_object::mapped_object(std::string const & scalartype, unsigned int id, std::string const & type_key) : type_key_(type_key) mapped_object::mapped_object(std::string const & scalartype, unsigned int id, std::string const & type_key) : type_key_(type_key)
{ {
@@ -93,12 +93,12 @@ binary_leaf::binary_leaf(mapped_object::node_info info) : info_(info){ }
void binary_leaf::process_recursive(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors) void binary_leaf::process_recursive(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors)
{ {
std::set<std::string> already_fetched; std::set<std::string> already_fetched;
process(stream, leaf, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping, already_fetched); process(stream, leaf, accessors, *info_.array_expression, info_.root_idx, *info_.mapping, already_fetched);
} }
std::string binary_leaf::evaluate_recursive(leaf_t leaf, std::map<std::string, std::string> const & accessors) std::string binary_leaf::evaluate_recursive(leaf_t leaf, std::map<std::string, std::string> const & accessors)
{ {
return evaluate(leaf, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping); return evaluate(leaf, accessors, *info_.array_expression, info_.root_idx, *info_.mapping);
} }
@@ -112,11 +112,11 @@ mapped_reduction::mapped_reduction(std::string const & scalartype, unsigned int
int_t mapped_reduction::root_idx() const int_t mapped_reduction::root_idx() const
{ return info_.root_idx; } { return info_.root_idx; }
atidlas::symbolic_expression const & mapped_reduction::symbolic_expression() const atidlas::array_expression const & mapped_reduction::array_expression() const
{ return *info_.symbolic_expression; } { return *info_.array_expression; }
symbolic_expression_node mapped_reduction::root_node() const array_expression::node mapped_reduction::root_node() const
{ return symbolic_expression().tree()[root_idx()]; } { return array_expression().tree()[root_idx()]; }
bool mapped_reduction::is_index_reduction() const bool mapped_reduction::is_index_reduction() const
{ {
@@ -129,7 +129,7 @@ bool mapped_reduction::is_index_reduction() const
op_element mapped_reduction::root_op() const op_element mapped_reduction::root_op() const
{ {
return info_.symbolic_expression->tree()[info_.root_idx].op; return info_.array_expression->tree()[info_.root_idx].op;
} }
@@ -247,10 +247,10 @@ mapped_array::mapped_array(std::string const & scalartype, unsigned int id, char
void mapped_vdiag::postprocess(std::string &res) const void mapped_vdiag::postprocess(std::string &res) const
{ {
std::map<std::string, std::string> accessors; std::map<std::string, std::string> accessors;
tools::find_and_replace(res, "#diag_offset", atidlas::evaluate(RHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping)); tools::find_and_replace(res, "#diag_offset", atidlas::evaluate(RHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping));
accessors["array1"] = res; accessors["array1"] = res;
accessors["host_scalar"] = res; accessors["host_scalar"] = res;
res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping); res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping);
} }
mapped_vdiag::mapped_vdiag(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "vdiag"), binary_leaf(info){} mapped_vdiag::mapped_vdiag(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "vdiag"), binary_leaf(info){}
@@ -260,9 +260,9 @@ mapped_vdiag::mapped_vdiag(std::string const & scalartype, unsigned int id, node
void mapped_matrix_row::postprocess(std::string &res) const void mapped_matrix_row::postprocess(std::string &res) const
{ {
std::map<std::string, std::string> accessors; std::map<std::string, std::string> accessors;
tools::find_and_replace(res, "#row", atidlas::evaluate(RHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping)); tools::find_and_replace(res, "#row", atidlas::evaluate(RHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping));
accessors["array2"] = res; accessors["array2"] = res;
res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping); res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping);
} }
mapped_matrix_row::mapped_matrix_row(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "matrix_row"), binary_leaf(info) mapped_matrix_row::mapped_matrix_row(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "matrix_row"), binary_leaf(info)
@@ -272,9 +272,9 @@ mapped_matrix_row::mapped_matrix_row(std::string const & scalartype, unsigned in
void mapped_matrix_column::postprocess(std::string &res) const void mapped_matrix_column::postprocess(std::string &res) const
{ {
std::map<std::string, std::string> accessors; std::map<std::string, std::string> accessors;
tools::find_and_replace(res, "#column", atidlas::evaluate(RHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping)); tools::find_and_replace(res, "#column", atidlas::evaluate(RHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping));
accessors["array2"] = res; accessors["array2"] = res;
res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping); res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping);
} }
mapped_matrix_column::mapped_matrix_column(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "matrix_column"), binary_leaf(info) mapped_matrix_column::mapped_matrix_column(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "matrix_column"), binary_leaf(info)
@@ -291,7 +291,7 @@ void mapped_repeat::postprocess(std::string &res) const
tools::find_and_replace(res, "#tuplearg3", args.process("#tuplearg3")); tools::find_and_replace(res, "#tuplearg3", args.process("#tuplearg3"));
accessors["array1"] = res; accessors["array1"] = res;
accessors["array2"] = res; accessors["array2"] = res;
res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping); res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping);
} }
mapped_repeat::mapped_repeat(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "repeat"), binary_leaf(info) mapped_repeat::mapped_repeat(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "repeat"), binary_leaf(info)
@@ -302,9 +302,9 @@ mapped_repeat::mapped_repeat(std::string const & scalartype, unsigned int id, no
void mapped_matrix_diag::postprocess(std::string &res) const void mapped_matrix_diag::postprocess(std::string &res) const
{ {
std::map<std::string, std::string> accessors; std::map<std::string, std::string> accessors;
tools::find_and_replace(res, "#diag_offset", atidlas::evaluate(RHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping)); tools::find_and_replace(res, "#diag_offset", atidlas::evaluate(RHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping));
accessors["array2"] = res; accessors["array2"] = res;
res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.symbolic_expression, info_.root_idx, *info_.mapping); res = atidlas::evaluate(LHS_NODE_TYPE, accessors, *info_.array_expression, info_.root_idx, *info_.mapping);
} }
mapped_matrix_diag::mapped_matrix_diag(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "matrix_diag"), binary_leaf(info) mapped_matrix_diag::mapped_matrix_diag(std::string const & scalartype, unsigned int id, node_info info) : mapped_object(scalartype, id, "matrix_diag"), binary_leaf(info)
@@ -320,7 +320,7 @@ void mapped_outer::postprocess(std::string &res) const
{ {
std::map<std::string, std::string> accessors; std::map<std::string, std::string> accessors;
accessors["array1"] = "$VALUE{"+i+"}"; accessors["array1"] = "$VALUE{"+i+"}";
return atidlas::evaluate(leaf_, accessors, *i_.symbolic_expression, i_.root_idx, *i_.mapping); return atidlas::evaluate(leaf_, accessors, *i_.array_expression, i_.root_idx, *i_.mapping);
} }
std::string operator()(std::string const &, std::string const &) const{return "";} std::string operator()(std::string const &, std::string const &) const{return "";}
private: private:

88
lib/backend/parse.cpp
View File

@@ -10,12 +10,12 @@ namespace detail
bool is_scalar_reduction(symbolic_expression_node const & node) bool is_scalar_reduction(array_expression::node const & node)
{ {
return node.op.type_family==OPERATOR_VECTOR_REDUCTION_TYPE_FAMILY; return node.op.type_family==OPERATOR_VECTOR_REDUCTION_TYPE_FAMILY;
} }
bool is_vector_reduction(symbolic_expression_node const & node) bool is_vector_reduction(array_expression::node const & node)
{ {
return node.op.type_family==OPERATOR_ROWS_REDUCTION_TYPE_FAMILY return node.op.type_family==OPERATOR_ROWS_REDUCTION_TYPE_FAMILY
|| node.op.type_family==OPERATOR_COLUMNS_REDUCTION_TYPE_FAMILY; || node.op.type_family==OPERATOR_COLUMNS_REDUCTION_TYPE_FAMILY;
@@ -109,29 +109,29 @@ namespace detail
filter_fun::filter_fun(pred_t pred, std::vector<size_t> & out) : pred_(pred), out_(out) filter_fun::filter_fun(pred_t pred, std::vector<size_t> & out) : pred_(pred), out_(out)
{ } { }
void filter_fun::operator()(atidlas::symbolic_expression const & symbolic_expression, size_t root_idx, leaf_t) const void filter_fun::operator()(atidlas::array_expression const & array_expression, size_t root_idx, leaf_t) const
{ {
symbolic_expression_node const * root_node = &symbolic_expression.tree()[root_idx]; array_expression::node const * root_node = &array_expression.tree()[root_idx];
if (pred_(*root_node)) if (pred_(*root_node))
out_.push_back(root_idx); out_.push_back(root_idx);
} }
// //
std::vector<size_t> filter_nodes(bool (*pred)(symbolic_expression_node const & node), atidlas::symbolic_expression const & symbolic_expression, bool inspect) std::vector<size_t> filter_nodes(bool (*pred)(array_expression::node const & node), atidlas::array_expression const & array_expression, bool inspect)
{ {
std::vector<size_t> res; std::vector<size_t> res;
traverse(symbolic_expression, symbolic_expression.root(), filter_fun(pred, res), inspect); traverse(array_expression, array_expression.root(), filter_fun(pred, res), inspect);
return res; return res;
} }
// //
filter_elements_fun::filter_elements_fun(symbolic_expression_node_subtype subtype, std::vector<lhs_rhs_element> & out) : filter_elements_fun::filter_elements_fun(array_expression_node_subtype subtype, std::vector<lhs_rhs_element> & out) :
subtype_(subtype), out_(out) subtype_(subtype), out_(out)
{ } { }
void filter_elements_fun::operator()(atidlas::symbolic_expression const & symbolic_expression, size_t root_idx, leaf_t) const void filter_elements_fun::operator()(atidlas::array_expression const & array_expression, size_t root_idx, leaf_t) const
{ {
symbolic_expression_node const * root_node = &symbolic_expression.tree()[root_idx]; array_expression::node const * root_node = &array_expression.tree()[root_idx];
if (root_node->lhs.subtype==subtype_) if (root_node->lhs.subtype==subtype_)
out_.push_back(root_node->lhs); out_.push_back(root_node->lhs);
if (root_node->rhs.subtype==subtype_) if (root_node->rhs.subtype==subtype_)
@@ -139,10 +139,10 @@ void filter_elements_fun::operator()(atidlas::symbolic_expression const & symbol
} }
std::vector<lhs_rhs_element> filter_elements(symbolic_expression_node_subtype subtype, atidlas::symbolic_expression const & symbolic_expression) std::vector<lhs_rhs_element> filter_elements(array_expression_node_subtype subtype, atidlas::array_expression const & array_expression)
{ {
std::vector<lhs_rhs_element> res; std::vector<lhs_rhs_element> res;
traverse(symbolic_expression, symbolic_expression.root(), filter_elements_fun(subtype, res), true); traverse(array_expression, array_expression.root(), filter_elements_fun(subtype, res), true);
return res; return res;
} }
@@ -225,9 +225,9 @@ evaluate_expression_traversal::evaluate_expression_traversal(std::map<std::strin
accessors_(accessors), str_(str), mapping_(mapping) accessors_(accessors), str_(str), mapping_(mapping)
{ } { }
void evaluate_expression_traversal::call_before_expansion(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx) const void evaluate_expression_traversal::call_before_expansion(atidlas::array_expression const & array_expression, int_t root_idx) const
{ {
symbolic_expression_node const & root_node = symbolic_expression.tree()[root_idx]; array_expression::node const & root_node = array_expression.tree()[root_idx];
if(detail::is_cast(root_node.op)) if(detail::is_cast(root_node.op))
str_ += mapping_.at(std::make_pair(root_idx, PARENT_NODE_TYPE))->evaluate(accessors_); str_ += mapping_.at(std::make_pair(root_idx, PARENT_NODE_TYPE))->evaluate(accessors_);
else if ((root_node.op.type_family==OPERATOR_UNARY_TYPE_FAMILY || detail::is_elementwise_function(root_node.op)) else if ((root_node.op.type_family==OPERATOR_UNARY_TYPE_FAMILY || detail::is_elementwise_function(root_node.op))
@@ -237,14 +237,14 @@ void evaluate_expression_traversal::call_before_expansion(atidlas::symbolic_expr
} }
void evaluate_expression_traversal::call_after_expansion(symbolic_expression const & /*symbolic_expression*/, int_t /*root_idx*/) const void evaluate_expression_traversal::call_after_expansion(array_expression const & /*array_expression*/, int_t /*root_idx*/) const
{ {
str_+=")"; str_+=")";
} }
void evaluate_expression_traversal::operator()(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf) const void evaluate_expression_traversal::operator()(atidlas::array_expression const & array_expression, int_t root_idx, leaf_t leaf) const
{ {
symbolic_expression_node const & root_node = symbolic_expression.tree()[root_idx]; array_expression::node const & root_node = array_expression.tree()[root_idx];
mapping_type::key_type key = std::make_pair(root_idx, leaf); mapping_type::key_type key = std::make_pair(root_idx, leaf);
if (leaf==PARENT_NODE_TYPE) if (leaf==PARENT_NODE_TYPE)
{ {
@@ -276,39 +276,39 @@ void evaluate_expression_traversal::operator()(atidlas::symbolic_expression cons
std::string evaluate(leaf_t leaf, std::map<std::string, std::string> const & accessors, std::string evaluate(leaf_t leaf, std::map<std::string, std::string> const & accessors,
atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, mapping_type const & mapping) atidlas::array_expression const & array_expression, int_t root_idx, mapping_type const & mapping)
{ {
std::string res; std::string res;
evaluate_expression_traversal traversal_functor(accessors, res, mapping); evaluate_expression_traversal traversal_functor(accessors, res, mapping);
symbolic_expression_node const & root_node = symbolic_expression.tree()[root_idx]; array_expression::node const & root_node = array_expression.tree()[root_idx];
if (leaf==RHS_NODE_TYPE) if (leaf==RHS_NODE_TYPE)
{ {
if (root_node.rhs.type_family==COMPOSITE_OPERATOR_FAMILY) if (root_node.rhs.type_family==COMPOSITE_OPERATOR_FAMILY)
traverse(symbolic_expression, root_node.rhs.node_index, traversal_functor, false); traverse(array_expression, root_node.rhs.node_index, traversal_functor, false);
else else
traversal_functor(symbolic_expression, root_idx, leaf); traversal_functor(array_expression, root_idx, leaf);
} }
else if (leaf==LHS_NODE_TYPE) else if (leaf==LHS_NODE_TYPE)
{ {
if (root_node.lhs.type_family==COMPOSITE_OPERATOR_FAMILY) if (root_node.lhs.type_family==COMPOSITE_OPERATOR_FAMILY)
traverse(symbolic_expression, root_node.lhs.node_index, traversal_functor, false); traverse(array_expression, root_node.lhs.node_index, traversal_functor, false);
else else
traversal_functor(symbolic_expression, root_idx, leaf); traversal_functor(array_expression, root_idx, leaf);
} }
else else
traverse(symbolic_expression, root_idx, traversal_functor, false); traverse(array_expression, root_idx, traversal_functor, false);
return res; return res;
} }
void evaluate(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors, void evaluate(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors,
symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings)
{ {
symbolic_expressions_container::data_type::const_iterator sit; array_expressions_container::data_type::const_iterator sit;
std::vector<mapping_type>::const_iterator mit; std::vector<mapping_type>::const_iterator mit;
for (mit = mappings.begin(), sit = symbolic_expressions.data().begin(); sit != symbolic_expressions.data().end(); ++mit, ++sit) for (mit = mappings.begin(), sit = array_expressions.data().begin(); sit != array_expressions.data().end(); ++mit, ++sit)
stream << evaluate(leaf, accessors, **sit, (*sit)->root(), *mit) << ";" << std::endl; stream << evaluate(leaf, accessors, **sit, (*sit)->root(), *mit) << ";" << std::endl;
} }
@@ -317,7 +317,7 @@ process_traversal::process_traversal(std::map<std::string, std::string> const &
accessors_(accessors), stream_(stream), mapping_(mapping), already_processed_(already_processed) accessors_(accessors), stream_(stream), mapping_(mapping), already_processed_(already_processed)
{ } { }
void process_traversal::operator()(symbolic_expression const & /*symbolic_expression*/, int_t root_idx, leaf_t leaf) const void process_traversal::operator()(array_expression const & /*array_expression*/, int_t root_idx, leaf_t leaf) const
{ {
mapping_type::const_iterator it = mapping_.find(std::make_pair(root_idx, leaf)); mapping_type::const_iterator it = mapping_.find(std::make_pair(root_idx, leaf));
if (it!=mapping_.end()) if (it!=mapping_.end())
@@ -342,44 +342,44 @@ void process_traversal::operator()(symbolic_expression const & /*symbolic_expres
void process(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors, void process(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors,
atidlas::symbolic_expression const & symbolic_expression, size_t root_idx, mapping_type const & mapping, std::set<std::string> & already_processed) atidlas::array_expression const & array_expression, size_t root_idx, mapping_type const & mapping, std::set<std::string> & already_processed)
{ {
process_traversal traversal_functor(accessors, stream, mapping, already_processed); process_traversal traversal_functor(accessors, stream, mapping, already_processed);
symbolic_expression_node const & root_node = symbolic_expression.tree()[root_idx]; array_expression::node const & root_node = array_expression.tree()[root_idx];
if (leaf==RHS_NODE_TYPE) if (leaf==RHS_NODE_TYPE)
{ {
if (root_node.rhs.type_family==COMPOSITE_OPERATOR_FAMILY) if (root_node.rhs.type_family==COMPOSITE_OPERATOR_FAMILY)
traverse(symbolic_expression, root_node.rhs.node_index, traversal_functor, true); traverse(array_expression, root_node.rhs.node_index, traversal_functor, true);
else else
traversal_functor(symbolic_expression, root_idx, leaf); traversal_functor(array_expression, root_idx, leaf);
} }
else if (leaf==LHS_NODE_TYPE) else if (leaf==LHS_NODE_TYPE)
{ {
if (root_node.lhs.type_family==COMPOSITE_OPERATOR_FAMILY) if (root_node.lhs.type_family==COMPOSITE_OPERATOR_FAMILY)
traverse(symbolic_expression, root_node.lhs.node_index, traversal_functor, true); traverse(array_expression, root_node.lhs.node_index, traversal_functor, true);
else else
traversal_functor(symbolic_expression, root_idx, leaf); traversal_functor(array_expression, root_idx, leaf);
} }
else else
{ {
traverse(symbolic_expression, root_idx, traversal_functor, true); traverse(array_expression, root_idx, traversal_functor, true);
} }
} }
void process(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors, void process(kernel_generation_stream & stream, leaf_t leaf, std::map<std::string, std::string> const & accessors,
symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings)
{ {
symbolic_expressions_container::data_type::const_iterator sit; array_expressions_container::data_type::const_iterator sit;
std::vector<mapping_type>::const_iterator mit; std::vector<mapping_type>::const_iterator mit;
std::set<std::string> already_processed; std::set<std::string> already_processed;
for (mit = mappings.begin(), sit = symbolic_expressions.data().begin(); sit != symbolic_expressions.data().end(); ++mit, ++sit) for (mit = mappings.begin(), sit = array_expressions.data().begin(); sit != array_expressions.data().end(); ++mit, ++sit)
process(stream, leaf, accessors, **sit, (*sit)->root(), *mit, already_processed); process(stream, leaf, accessors, **sit, (*sit)->root(), *mit, already_processed);
} }
void symbolic_expression_representation_functor::append_id(char * & ptr, unsigned int val) void array_expression_representation_functor::append_id(char * & ptr, unsigned int val)
{ {
if (val==0) if (val==0)
*ptr++='0'; *ptr++='0';
@@ -391,31 +391,31 @@ void symbolic_expression_representation_functor::append_id(char * & ptr, unsigne
} }
} }
void symbolic_expression_representation_functor::append(cl_mem h, numeric_type dtype, char prefix) const void array_expression_representation_functor::append(cl_mem h, numeric_type dtype, char prefix) const
{ {
*ptr_++=prefix; *ptr_++=prefix;
*ptr_++=(char)dtype; *ptr_++=(char)dtype;
append_id(ptr_, binder_.get(h)); append_id(ptr_, binder_.get(h));
} }
void symbolic_expression_representation_functor::append(lhs_rhs_element const & lhs_rhs) const void array_expression_representation_functor::append(lhs_rhs_element const & lhs_rhs) const
{ {
if(lhs_rhs.subtype==DENSE_ARRAY_TYPE) if(lhs_rhs.subtype==DENSE_ARRAY_TYPE)
append(lhs_rhs.array.data, lhs_rhs.array.dtype, (char)(((int)'0')+((int)(lhs_rhs.array.shape1>1) + (int)(lhs_rhs.array.shape2>1)))); append(lhs_rhs.array.data, lhs_rhs.array.dtype, (char)(((int)'0')+((int)(lhs_rhs.array.shape1>1) + (int)(lhs_rhs.array.shape2>1))));
} }
symbolic_expression_representation_functor::symbolic_expression_representation_functor(symbolic_binder & binder, char *& ptr) : binder_(binder), ptr_(ptr){ } array_expression_representation_functor::array_expression_representation_functor(symbolic_binder & binder, char *& ptr) : binder_(binder), ptr_(ptr){ }
void symbolic_expression_representation_functor::append(char*& p, const char * str) const void array_expression_representation_functor::append(char*& p, const char * str) const
{ {
std::size_t n = std::strlen(str); std::size_t n = std::strlen(str);
std::memcpy(p, str, n); std::memcpy(p, str, n);
p+=n; p+=n;
} }
void symbolic_expression_representation_functor::operator()(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf_t) const void array_expression_representation_functor::operator()(atidlas::array_expression const & array_expression, int_t root_idx, leaf_t leaf_t) const
{ {
symbolic_expression_node const & root_node = symbolic_expression.tree()[root_idx]; array_expression::node const & root_node = array_expression.tree()[root_idx];
if (leaf_t==LHS_NODE_TYPE && root_node.lhs.type_family != COMPOSITE_OPERATOR_FAMILY) if (leaf_t==LHS_NODE_TYPE && root_node.lhs.type_family != COMPOSITE_OPERATOR_FAMILY)
append(root_node.lhs); append(root_node.lhs);
else if (leaf_t==RHS_NODE_TYPE && root_node.rhs.type_family != COMPOSITE_OPERATOR_FAMILY) else if (leaf_t==RHS_NODE_TYPE && root_node.rhs.type_family != COMPOSITE_OPERATOR_FAMILY)

98
lib/backend/templates/base.cpp
View File

@@ -20,19 +20,19 @@ namespace atidlas
base::parameters_type::parameters_type(unsigned int _simd_width, int_t _local_size_1, int_t _local_size_2, int_t _num_kernels) : simd_width(_simd_width), local_size_0(_local_size_1), local_size_1(_local_size_2), num_kernels(_num_kernels) base::parameters_type::parameters_type(unsigned int _simd_width, int_t _local_size_1, int_t _local_size_2, int_t _num_kernels) : simd_width(_simd_width), local_size_0(_local_size_1), local_size_1(_local_size_2), num_kernels(_num_kernels)
{ } { }
numeric_type base::map_functor::get_numeric_type(atidlas::symbolic_expression const * symbolic_expression, int_t root_idx) const numeric_type base::map_functor::get_numeric_type(atidlas::array_expression const * array_expression, int_t root_idx) const
{ {
symbolic_expression_node const * root_node = &symbolic_expression->tree()[root_idx]; array_expression::node const * root_node = &array_expression->tree()[root_idx];
while (root_node->lhs.dtype==INVALID_NUMERIC_TYPE) while (root_node->lhs.dtype==INVALID_NUMERIC_TYPE)
root_node = &symbolic_expression->tree()[root_node->lhs.node_index]; root_node = &array_expression->tree()[root_node->lhs.node_index];
return root_node->lhs.dtype; return root_node->lhs.dtype;
} }
/** @brief Binary leaf */ /** @brief Binary leaf */
template<class T> template<class T>
tools::shared_ptr<mapped_object> base::map_functor::binary_leaf(atidlas::symbolic_expression const * symbolic_expression, int_t root_idx, mapping_type const * mapping) const tools::shared_ptr<mapped_object> base::map_functor::binary_leaf(atidlas::array_expression const * array_expression, int_t root_idx, mapping_type const * mapping) const
{ {
return tools::shared_ptr<mapped_object>(new T(numeric_type_to_string(symbolic_expression->dtype()), binder_.get(NULL), mapped_object::node_info(mapping, symbolic_expression, root_idx))); return tools::shared_ptr<mapped_object>(new T(numeric_type_to_string(array_expression->dtype()), binder_.get(NULL), mapped_object::node_info(mapping, array_expression, root_idx)));
} }
/** @brief Scalar mapping */ /** @brief Scalar mapping */
@@ -82,9 +82,9 @@ tools::shared_ptr<mapped_object> base::map_functor::create(lhs_rhs_element const
base::map_functor::map_functor(symbolic_binder & binder, mapping_type & mapping) : binder_(binder), mapping_(mapping){ } base::map_functor::map_functor(symbolic_binder & binder, mapping_type & mapping) : binder_(binder), mapping_(mapping){ }
/** @brief Traversal functor */ /** @brief Traversal functor */
void base::map_functor::operator()(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf_t) const { void base::map_functor::operator()(atidlas::array_expression const & array_expression, int_t root_idx, leaf_t leaf_t) const {
mapping_type::key_type key(root_idx, leaf_t); mapping_type::key_type key(root_idx, leaf_t);
symbolic_expression_node const & root_node = symbolic_expression.tree()[root_idx]; array_expression::node const & root_node = array_expression.tree()[root_idx];
if (leaf_t == LHS_NODE_TYPE && root_node.lhs.type_family != COMPOSITE_OPERATOR_FAMILY) if (leaf_t == LHS_NODE_TYPE && root_node.lhs.type_family != COMPOSITE_OPERATOR_FAMILY)
mapping_.insert(mapping_type::value_type(key, create(root_node.lhs))); mapping_.insert(mapping_type::value_type(key, create(root_node.lhs)));
@@ -93,23 +93,23 @@ void base::map_functor::operator()(atidlas::symbolic_expression const & symbolic
else if ( leaf_t== PARENT_NODE_TYPE) else if ( leaf_t== PARENT_NODE_TYPE)
{ {
if (root_node.op.type==OPERATOR_VDIAG_TYPE) if (root_node.op.type==OPERATOR_VDIAG_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_vdiag>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_vdiag>(&array_expression, root_idx, &mapping_)));
else if (root_node.op.type==OPERATOR_MATRIX_DIAG_TYPE) else if (root_node.op.type==OPERATOR_MATRIX_DIAG_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_diag>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_diag>(&array_expression, root_idx, &mapping_)));
else if (root_node.op.type==OPERATOR_MATRIX_ROW_TYPE) else if (root_node.op.type==OPERATOR_MATRIX_ROW_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_row>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_row>(&array_expression, root_idx, &mapping_)));
else if (root_node.op.type==OPERATOR_MATRIX_COLUMN_TYPE) else if (root_node.op.type==OPERATOR_MATRIX_COLUMN_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_column>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_matrix_column>(&array_expression, root_idx, &mapping_)));
else if (detail::is_scalar_reduction(root_node)) else if (detail::is_scalar_reduction(root_node))
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_scalar_reduction>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_scalar_reduction>(&array_expression, root_idx, &mapping_)));
else if (detail::is_vector_reduction(root_node)) else if (detail::is_vector_reduction(root_node))
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_mreduction>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_mreduction>(&array_expression, root_idx, &mapping_)));
else if (root_node.op.type_family == OPERATOR_MATRIX_PRODUCT_TYPE_FAMILY) else if (root_node.op.type_family == OPERATOR_MATRIX_PRODUCT_TYPE_FAMILY)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_mproduct>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_mproduct>(&array_expression, root_idx, &mapping_)));
else if (root_node.op.type == OPERATOR_REPEAT_TYPE) else if (root_node.op.type == OPERATOR_REPEAT_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_repeat>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_repeat>(&array_expression, root_idx, &mapping_)));
else if (root_node.op.type == OPERATOR_OUTER_PROD_TYPE) else if (root_node.op.type == OPERATOR_OUTER_PROD_TYPE)
mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_outer>(&symbolic_expression, root_idx, &mapping_))); mapping_.insert(mapping_type::value_type(key, binary_leaf<mapped_outer>(&array_expression, root_idx, &mapping_)));
else if (detail::is_cast(root_node.op)) else if (detail::is_cast(root_node.op))
mapping_.insert(mapping_type::value_type(key, tools::shared_ptr<mapped_object>(new mapped_cast(root_node.op.type, binder_.get(NULL))))); mapping_.insert(mapping_type::value_type(key, tools::shared_ptr<mapped_object>(new mapped_cast(root_node.op.type, binder_.get(NULL)))));
} }
@@ -187,9 +187,9 @@ void base::set_arguments_functor::set_arguments(lhs_rhs_element const & lhs_rhs)
} }
/** @brief Traversal functor: */ /** @brief Traversal functor: */
void base::set_arguments_functor::operator()(atidlas::symbolic_expression const & symbolic_expression, int_t root_idx, leaf_t leaf_t) const void base::set_arguments_functor::operator()(atidlas::array_expression const & array_expression, int_t root_idx, leaf_t leaf_t) const
{ {
symbolic_expression_node const & root_node = symbolic_expression.tree()[root_idx]; array_expression::node const & root_node = array_expression.tree()[root_idx];
if (leaf_t==LHS_NODE_TYPE && root_node.lhs.type_family != COMPOSITE_OPERATOR_FAMILY) if (leaf_t==LHS_NODE_TYPE && root_node.lhs.type_family != COMPOSITE_OPERATOR_FAMILY)
set_arguments(root_node.lhs); set_arguments(root_node.lhs);
else if (leaf_t==RHS_NODE_TYPE && root_node.rhs.type_family != COMPOSITE_OPERATOR_FAMILY) else if (leaf_t==RHS_NODE_TYPE && root_node.rhs.type_family != COMPOSITE_OPERATOR_FAMILY)
@@ -258,30 +258,30 @@ std::string base::neutral_element(op_element const & op)
} }
} }
std::string base::generate_arguments(std::vector<mapping_type> const & mappings, std::map<std::string, std::string> const & accessors, symbolic_expressions_container const & symbolic_expressions) std::string base::generate_arguments(std::vector<mapping_type> const & mappings, std::map<std::string, std::string> const & accessors, array_expressions_container const & array_expressions)
{ {
kernel_generation_stream stream; kernel_generation_stream stream;
process(stream, PARENT_NODE_TYPE, accessors, symbolic_expressions, mappings); process(stream, PARENT_NODE_TYPE, accessors, array_expressions, mappings);
std::string res = stream.str(); std::string res = stream.str();
res.erase(res.rfind(',')); res.erase(res.rfind(','));
return res; return res;
} }
std::string base::generate_arguments(std::string const & data_type, std::vector<mapping_type> const & mappings, symbolic_expressions_container const & symbolic_expressions) std::string base::generate_arguments(std::string const & data_type, std::vector<mapping_type> const & mappings, array_expressions_container const & array_expressions)
{ {
return generate_arguments(mappings, tools::make_map<std::map<std::string, std::string> >("array0", "__global #scalartype* #pointer, uint #start,") return generate_arguments(mappings, tools::make_map<std::map<std::string, std::string> >("array0", "__global #scalartype* #pointer, uint #start,")
("host_scalar", "#scalartype #name,") ("host_scalar", "#scalartype #name,")
("array1", "__global " + data_type + "* #pointer, uint #start, uint #stride,") ("array1", "__global " + data_type + "* #pointer, uint #start, uint #stride,")
("array2", "__global " + data_type + "* #pointer, uint #ld, uint #start1, uint #start2, uint #stride1, uint #stride2,") ("array2", "__global " + data_type + "* #pointer, uint #ld, uint #start1, uint #start2, uint #stride1, uint #stride2,")
("tuple4", "#scalartype #name0, #scalartype #name1, #scalartype #name2, #scalartype #name3,"), symbolic_expressions); ("tuple4", "#scalartype #name0, #scalartype #name1, #scalartype #name2, #scalartype #name3,"), array_expressions);
} }
void base::set_arguments(symbolic_expressions_container const & symbolic_expressions, cl::Kernel & kernel, unsigned int & current_arg) void base::set_arguments(array_expressions_container const & array_expressions, cl::Kernel & kernel, unsigned int & current_arg)
{ {
tools::shared_ptr<symbolic_binder> binder = make_binder(); tools::shared_ptr<symbolic_binder> binder = make_binder();
for (symbolic_expressions_container::data_type::const_iterator itt = symbolic_expressions.data().begin(); itt != symbolic_expressions.data().end(); ++itt) for (array_expressions_container::data_type::const_iterator itt = array_expressions.data().begin(); itt != array_expressions.data().end(); ++itt)
traverse(**itt, (*itt)->root(), set_arguments_functor(*binder, current_arg, kernel), true); traverse(**itt, (*itt)->root(), set_arguments_functor(*binder, current_arg, kernel), true);
} }
@@ -304,7 +304,7 @@ void base::fill_kernel_name(char * ptr, unsigned int label, const char * suffix)
base::invalid_exception::invalid_exception() : message_() {} base::invalid_exception::invalid_exception() : message_() {}
base::invalid_exception::invalid_exception(std::string message) : base::invalid_exception::invalid_exception(std::string message) :
message_("ViennaCL: Internal error: The generator cannot apply the given template to the given symbolic_expression: " + message + "\n" message_("ViennaCL: Internal error: The generator cannot apply the given template to the given array_expression: " + message + "\n"
"If you are using a builtin template, please report on viennacl-support@lists.sourceforge.net! We will provide a fix as soon as possible\n" "If you are using a builtin template, please report on viennacl-support@lists.sourceforge.net! We will provide a fix as soon as possible\n"
"If you are using your own template, please try using other parameters") {} "If you are using your own template, please try using other parameters") {}
@@ -335,15 +335,15 @@ void base::fetching_loop_info(fetching_policy_type policy, std::string const & b
} }
} }
bool base::is_node_trans(symbolic_expression::container_type const & array, size_t root_idx, leaf_t leaf_type) bool base::is_node_trans(array_expression::container_type const & array, size_t root_idx, leaf_t leaf_type)
{ {
bool res = false; bool res = false;
lhs_rhs_element symbolic_expression_node::*ptr; lhs_rhs_element array_expression::node::*ptr;
if (leaf_type==LHS_NODE_TYPE) if (leaf_type==LHS_NODE_TYPE)
ptr = &symbolic_expression_node::lhs; ptr = &array_expression::node::lhs;
else else
ptr = &symbolic_expression_node::rhs; ptr = &array_expression::node::rhs;
symbolic_expression_node const * node = &array[root_idx]; array_expression::node const * node = &array[root_idx];
while ((node->*ptr).type_family==COMPOSITE_OPERATOR_FAMILY) while ((node->*ptr).type_family==COMPOSITE_OPERATOR_FAMILY)
{ {
if (array[(node->*ptr).node_index].op.type==OPERATOR_TRANS_TYPE) if (array[(node->*ptr).node_index].op.type==OPERATOR_TRANS_TYPE)
@@ -361,7 +361,7 @@ std::string base::append_simd_suffix(std::string const & str, unsigned int i)
return str + tools::to_string(suffixes[i]); return str + tools::to_string(suffixes[i]);
} }
bool base::is_strided(symbolic_expression_node const & node) bool base::is_strided(array_expression::node const & node)
{ {
return node.op.type==OPERATOR_VDIAG_TYPE return node.op.type==OPERATOR_VDIAG_TYPE
|| node.op.type==OPERATOR_MATRIX_DIAG_TYPE || node.op.type==OPERATOR_MATRIX_DIAG_TYPE
@@ -370,17 +370,17 @@ bool base::is_strided(symbolic_expression_node const & node)
|| node.op.type==OPERATOR_OUTER_PROD_TYPE; || node.op.type==OPERATOR_OUTER_PROD_TYPE;
} }
bool base::requires_fallback(symbolic_expressions_container const & symbolic_expressions) bool base::requires_fallback(array_expressions_container const & array_expressions)
{ {
for (symbolic_expressions_container::data_type::const_iterator it = symbolic_expressions.data().begin(); it != symbolic_expressions.data().end(); ++it) for (array_expressions_container::data_type::const_iterator it = array_expressions.data().begin(); it != array_expressions.data().end(); ++it)
for(symbolic_expression::container_type::const_iterator itt = (*it)->tree().begin(); itt != (*it)->tree().end() ; ++itt) for(array_expression::container_type::const_iterator itt = (*it)->tree().begin(); itt != (*it)->tree().end() ; ++itt)
if( (itt->lhs.subtype==DENSE_ARRAY_TYPE && (std::max(itt->lhs.array.stride1, itt->lhs.array.stride2)>1 || std::max(itt->lhs.array.start1,itt->lhs.array.start2)>0)) if( (itt->lhs.subtype==DENSE_ARRAY_TYPE && (std::max(itt->lhs.array.stride1, itt->lhs.array.stride2)>1 || std::max(itt->lhs.array.start1,itt->lhs.array.start2)>0))
|| (itt->rhs.subtype==DENSE_ARRAY_TYPE && (std::max(itt->rhs.array.stride1, itt->rhs.array.stride2)>1 || std::max(itt->rhs.array.start1,itt->rhs.array.start2)>0))) || (itt->rhs.subtype==DENSE_ARRAY_TYPE && (std::max(itt->rhs.array.stride1, itt->rhs.array.stride2)>1 || std::max(itt->rhs.array.start1,itt->rhs.array.start2)>0)))
return true; return true;
return false; return false;
} }
int_t base::vector_size(symbolic_expression_node const & node) int_t base::vector_size(array_expression::node const & node)
{ {
using namespace tools; using namespace tools;
if (node.op.type==OPERATOR_MATRIX_DIAG_TYPE) if (node.op.type==OPERATOR_MATRIX_DIAG_TYPE)
@@ -394,7 +394,7 @@ int_t base::vector_size(symbolic_expression_node const & node)
} }
std::pair<int_t, int_t> base::matrix_size(symbolic_expression_node const & node) std::pair<int_t, int_t> base::matrix_size(array_expression::node const & node)
{ {
if (node.op.type==OPERATOR_VDIAG_TYPE) if (node.op.type==OPERATOR_VDIAG_TYPE)
{ {
@@ -433,7 +433,7 @@ void base::element_wise_loop_1D(kernel_generation_stream & stream, loop_body_bas
} }
} }
bool base::is_reduction(symbolic_expression_node const & node) bool base::is_reduction(array_expression::node const & node)
{ {
return node.op.type_family==OPERATOR_VECTOR_REDUCTION_TYPE_FAMILY return node.op.type_family==OPERATOR_VECTOR_REDUCTION_TYPE_FAMILY
|| node.op.type_family==OPERATOR_COLUMNS_REDUCTION_TYPE_FAMILY || node.op.type_family==OPERATOR_COLUMNS_REDUCTION_TYPE_FAMILY
@@ -490,34 +490,34 @@ tools::shared_ptr<symbolic_binder> base::make_binder()
base::base(binding_policy_t binding_policy) : binding_policy_(binding_policy) base::base(binding_policy_t binding_policy) : binding_policy_(binding_policy)
{} {}
unsigned int base::lmem_usage(symbolic_expressions_container const &) const unsigned int base::lmem_usage(array_expressions_container const &) const
{ return 0; } { return 0; }
unsigned int base::registers_usage(symbolic_expressions_container const &) const unsigned int base::registers_usage(array_expressions_container const &) const
{ return 0; } { return 0; }
base::~base() base::~base()
{ } { }
std::vector<std::string> base::generate(unsigned int label, symbolic_expressions_container const & symbolic_expressions, cl::Device const & device) std::vector<std::string> base::generate(unsigned int label, array_expressions_container const & array_expressions, cl::Device const & device)
{ {
symbolic_expressions_container::data_type::const_iterator sit; array_expressions_container::data_type::const_iterator sit;
std::vector<mapping_type>::iterator mit; std::vector<mapping_type>::iterator mit;
if(int err = check_invalid(symbolic_expressions, device)) if(int err = check_invalid(array_expressions, device))
throw operation_not_supported_exception("The supplied parameters for this template are invalid : err " + tools::to_string(err)); throw operation_not_supported_exception("The supplied parameters for this template are invalid : err " + tools::to_string(err));
//Create mapping //Create mapping
std::vector<mapping_type> mappings(symbolic_expressions.data().size()); std::vector<mapping_type> mappings(array_expressions.data().size());
tools::shared_ptr<symbolic_binder> binder = make_binder(); tools::shared_ptr<symbolic_binder> binder = make_binder();
for (mit = mappings.begin(), sit = symbolic_expressions.data().begin(); sit != symbolic_expressions.data().end(); ++sit, ++mit) for (mit = mappings.begin(), sit = array_expressions.data().begin(); sit != array_expressions.data().end(); ++sit, ++mit)
traverse(**sit, (*sit)->root(), map_functor(*binder,*mit), true); traverse(**sit, (*sit)->root(), map_functor(*binder,*mit), true);
return generate_impl(label, symbolic_expressions, mappings); return generate_impl(label, array_expressions, mappings);
} }
template<class TType, class PType> template<class TType, class PType>
int base_impl<TType, PType>::check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const int base_impl<TType, PType>::check_invalid_impl(cl::Device const &, array_expressions_container const &) const
{ return TEMPLATE_VALID; } { return TEMPLATE_VALID; }
template<class TType, class PType> template<class TType, class PType>
@@ -537,11 +537,11 @@ tools::shared_ptr<base> base_impl<TType, PType>::clone() const
{ return tools::shared_ptr<base>(new TType(*dynamic_cast<TType const *>(this))); } { return tools::shared_ptr<base>(new TType(*dynamic_cast<TType const *>(this))); }
template<class TType, class PType> template<class TType, class PType>
int base_impl<TType, PType>::check_invalid(symbolic_expressions_container const & symbolic_expressions, cl::Device const & device) const int base_impl<TType, PType>::check_invalid(array_expressions_container const & array_expressions, cl::Device const & device) const
{ {
//Query device informations //Query device informations
size_t lmem_available = device.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>(); size_t lmem_available = device.getInfo<CL_DEVICE_LOCAL_MEM_SIZE>();
size_t lmem_used = lmem_usage(symbolic_expressions); size_t lmem_used = lmem_usage(array_expressions);
if (lmem_used>lmem_available) if (lmem_used>lmem_available)
return TEMPLATE_LOCAL_MEMORY_OVERFLOW; return TEMPLATE_LOCAL_MEMORY_OVERFLOW;
@@ -575,7 +575,7 @@ int base_impl<TType, PType>::check_invalid(symbolic_expressions_container const
p_.simd_width!=16) p_.simd_width!=16)
return TEMPLATE_INVALID_SIMD_WIDTH; return TEMPLATE_INVALID_SIMD_WIDTH;
return check_invalid_impl(device, symbolic_expressions); return check_invalid_impl(device, array_expressions);
} }
template class base_impl<vaxpy, vaxpy_parameters>; template class base_impl<vaxpy, vaxpy_parameters>;

30
lib/backend/templates/maxpy.cpp
View File

@@ -14,7 +14,7 @@ maxpy_parameters::maxpy_parameters(unsigned int _simd_width,
int maxpy::check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const int maxpy::check_invalid_impl(cl::Device const &, array_expressions_container const &) const
{ {
if (p_.simd_width>1) if (p_.simd_width>1)
return TEMPLATE_INVALID_SIMD_WIDTH; return TEMPLATE_INVALID_SIMD_WIDTH;
@@ -23,7 +23,7 @@ int maxpy::check_invalid_impl(cl::Device const &, symbolic_expressions_container
return TEMPLATE_VALID; return TEMPLATE_VALID;
} }
std::string maxpy::generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width) const std::string maxpy::generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width) const
{ {
kernel_generation_stream stream; kernel_generation_stream stream;
@@ -33,13 +33,13 @@ std::string maxpy::generate_impl(unsigned int label, symbolic_expressions_contai
fill_kernel_name(kprefix, label, "d"); fill_kernel_name(kprefix, label, "d");
stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl; stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl;
stream << "__kernel void " << kprefix << "(unsigned int M, unsigned int N, " << generate_arguments("#scalartype", mappings, symbolic_expressions) << ")" << std::endl; stream << "__kernel void " << kprefix << "(unsigned int M, unsigned int N, " << generate_arguments("#scalartype", mappings, array_expressions) << ")" << std::endl;
stream << "{" << std::endl; stream << "{" << std::endl;
stream.inc_tab(); stream.inc_tab();
process(stream, PARENT_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array0", "#scalartype #namereg = #pointer[#start];") process(stream, PARENT_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array0", "#scalartype #namereg = #pointer[#start];")
("array1", "#pointer += #start;") ("array1", "#pointer += #start;")
("array2", "#pointer = &$VALUE{#start1, #start2};"), symbolic_expressions, mappings); ("array2", "#pointer = &$VALUE{#start1, #start2};"), array_expressions, mappings);
fetching_loop_info(p_.fetching_policy, "M", stream, init0, upper_bound0, inc0, "get_global_id(0)", "get_global_size(0)"); fetching_loop_info(p_.fetching_policy, "M", stream, init0, upper_bound0, inc0, "get_global_id(0)", "get_global_size(0)");
stream << "for(unsigned int i = " << init0 << "; i < " << upper_bound0 << "; i += " << inc0 << ")" << std::endl; stream << "for(unsigned int i = " << init0 << "; i < " << upper_bound0 << "; i += " << inc0 << ")" << std::endl;
@@ -55,7 +55,7 @@ std::string maxpy::generate_impl(unsigned int label, symbolic_expressions_contai
("vdiag", "#scalartype #namereg = ((i + ((#diag_offset<0)?#diag_offset:0))!=(j-((#diag_offset>0)?#diag_offset:0)))?0:$VALUE{min(i*#stride1, j*#stride1)};") ("vdiag", "#scalartype #namereg = ((i + ((#diag_offset<0)?#diag_offset:0))!=(j-((#diag_offset>0)?#diag_offset:0)))?0:$VALUE{min(i*#stride1, j*#stride1)};")
("repeat", "#scalartype #namereg = $VALUE{(i%#tuplearg0)*#stride1, (j%#tuplearg1)*#stride2};") ("repeat", "#scalartype #namereg = $VALUE{(i%#tuplearg0)*#stride1, (j%#tuplearg1)*#stride2};")
("outer", "#scalartype #namereg = ($LVALUE{i*#stride})*($RVALUE{j*#stride});") ("outer", "#scalartype #namereg = ($LVALUE{i*#stride})*($RVALUE{j*#stride});")
, symbolic_expressions, mappings); , array_expressions, mappings);
evaluate(stream, PARENT_NODE_TYPE, tools::make_map<std::map<std::string, std::string> > evaluate(stream, PARENT_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >
("array2", "#namereg") ("array2", "#namereg")
@@ -64,10 +64,10 @@ std::string maxpy::generate_impl(unsigned int label, symbolic_expressions_contai
("array0", "#namereg") ("array0", "#namereg")
("outer", "#namereg") ("outer", "#namereg")
("cast", "convert_"+data_type) ("cast", "convert_"+data_type)
, symbolic_expressions, mappings); , array_expressions, mappings);
process(stream, LHS_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array2", "$VALUE{i*#stride1,j*#stride2} = #namereg;") process(stream, LHS_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array2", "$VALUE{i*#stride1,j*#stride2} = #namereg;")
, symbolic_expressions, mappings); , array_expressions, mappings);
stream.dec_tab(); stream.dec_tab();
stream << "}" << std::endl; stream << "}" << std::endl;
@@ -81,10 +81,10 @@ std::string maxpy::generate_impl(unsigned int label, symbolic_expressions_contai
return stream.str(); return stream.str();
} }
std::vector<std::string> maxpy::generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const std::vector<std::string> maxpy::generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const
{ {
std::vector<std::string> res; std::vector<std::string> res;
res.push_back(generate_impl(label, symbolic_expressions, mappings, 1)); res.push_back(generate_impl(label, array_expressions, mappings, 1));
return res; return res;
} }
@@ -97,17 +97,17 @@ maxpy::maxpy(unsigned int simd, unsigned int ls1, unsigned int ls2,
base_impl<maxpy, maxpy_parameters>(maxpy_parameters(simd, ls1, ls2, ng1, ng2, fetch), bind) base_impl<maxpy, maxpy_parameters>(maxpy_parameters(simd, ls1, ls2, ng1, ng2, fetch), bind)
{} {}
std::vector<int_t> maxpy::input_sizes(symbolic_expressions_container const & symbolic_expressions) std::vector<int_t> maxpy::input_sizes(array_expressions_container const & array_expressions)
{ {
atidlas::symbolic_expression const & symbolic_expression = *(symbolic_expressions.data().front()); atidlas::array_expression const & array_expression = *(array_expressions.data().front());
std::pair<int_t, int_t> size = matrix_size(lhs_most(symbolic_expression.tree(), symbolic_expression.root())); std::pair<int_t, int_t> size = matrix_size(lhs_most(array_expression.tree(), array_expression.root()));
return tools::make_vector<int_t>() << size.first << size.second; return tools::make_vector<int_t>() << size.first << size.second;
} }
void maxpy::enqueue(cl::CommandQueue & queue, void maxpy::enqueue(cl::CommandQueue & queue,
std::vector<cl_ext::lazy_compiler> & programs, std::vector<cl_ext::lazy_compiler> & programs,
unsigned int label, unsigned int label,
symbolic_expressions_container const & symbolic_expressions) array_expressions_container const & array_expressions)
{ {
char kname[10]; char kname[10];
fill_kernel_name(kname, label, "d"); fill_kernel_name(kname, label, "d");
@@ -116,10 +116,10 @@ void maxpy::enqueue(cl::CommandQueue & queue,
cl::NDRange grange(p_.local_size_0*p_.num_groups_0, p_.local_size_1*p_.num_groups_1); cl::NDRange grange(p_.local_size_0*p_.num_groups_0, p_.local_size_1*p_.num_groups_1);
cl::NDRange lrange(p_.local_size_0, p_.local_size_1); cl::NDRange lrange(p_.local_size_0, p_.local_size_1);
unsigned int current_arg = 0; unsigned int current_arg = 0;
std::vector<int_t> MN = input_sizes(symbolic_expressions); std::vector<int_t> MN = input_sizes(array_expressions);
kernel.setArg(current_arg++, cl_uint(MN[0])); kernel.setArg(current_arg++, cl_uint(MN[0]));
kernel.setArg(current_arg++, cl_uint(MN[1])); kernel.setArg(current_arg++, cl_uint(MN[1]));
set_arguments(symbolic_expressions, kernel, current_arg); set_arguments(array_expressions, kernel, current_arg);
queue.enqueueNDRangeKernel(kernel, cl::NullRange, grange, lrange); queue.enqueueNDRangeKernel(kernel, cl::NullRange, grange, lrange);
} }

40
lib/backend/templates/mproduct.cpp
View File

@@ -17,10 +17,10 @@ mproduct_parameters::mproduct_parameters(unsigned int simd_width
mL(ms*local_size_0), nL(ns*local_size_1){} mL(ms*local_size_0), nL(ns*local_size_1){}
unsigned int mproduct::lmem_usage(symbolic_expressions_container const & symbolic_expressions) const unsigned int mproduct::lmem_usage(array_expressions_container const & array_expressions) const
{ {
atidlas::symbolic_expression const & symbolic_expression = (*symbolic_expressions.data().front()); atidlas::array_expression const & array_expression = (*array_expressions.data().front());
numeric_type numeric_t = lhs_most(symbolic_expression.tree(), symbolic_expression.root()).lhs.dtype; numeric_type numeric_t = lhs_most(array_expression.tree(), array_expression.root()).lhs.dtype;
unsigned int N = 0; unsigned int N = 0;
if (p_.A_fetching_policy==FETCH_FROM_LOCAL) if (p_.A_fetching_policy==FETCH_FROM_LOCAL)
@@ -30,16 +30,16 @@ mproduct_parameters::mproduct_parameters(unsigned int simd_width
return N*size_of(numeric_t); return N*size_of(numeric_t);
} }
unsigned int mproduct::registers_usage(symbolic_expressions_container const & symbolic_expressions) const unsigned int mproduct::registers_usage(array_expressions_container const & array_expressions) const
{ {
atidlas::symbolic_expression const & symbolic_expression = (*symbolic_expressions.data().front()); atidlas::array_expression const & array_expression = (*array_expressions.data().front());
numeric_type numeric_t = lhs_most(symbolic_expression.tree(), symbolic_expression.root()).lhs.dtype; numeric_type numeric_t = lhs_most(array_expression.tree(), array_expression.root()).lhs.dtype;
unsigned int N = p_.mS * p_.nS + p_.mS * p_.kS + p_.kS * p_.nS; unsigned int N = p_.mS * p_.nS + p_.mS * p_.kS + p_.kS * p_.nS;
return N*size_of(numeric_t); return N*size_of(numeric_t);
} }
int mproduct::check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const int mproduct::check_invalid_impl(cl::Device const &, array_expressions_container const &) const
{ {
if (p_.A_fetching_policy!=FETCH_FROM_LOCAL && p_.B_fetching_policy!=FETCH_FROM_LOCAL&& (p_.local_fetch_0!=0 || p_.local_fetch_1!=0)) if (p_.A_fetching_policy!=FETCH_FROM_LOCAL && p_.B_fetching_policy!=FETCH_FROM_LOCAL&& (p_.local_fetch_0!=0 || p_.local_fetch_1!=0))
return TEMPLATE_GLOBAL_MEMORY_REQUIRES_ZERO_LOCAL_FETCH; return TEMPLATE_GLOBAL_MEMORY_REQUIRES_ZERO_LOCAL_FETCH;
@@ -87,7 +87,7 @@ mproduct_parameters::mproduct_parameters(unsigned int simd_width
return TEMPLATE_VALID; return TEMPLATE_VALID;
} }
std::string mproduct::generate_impl(unsigned int label, const char * id, const symbolic_expressions_container &symbolic_expressions, const std::vector<mapping_type> &, bool fallback) const std::string mproduct::generate_impl(unsigned int label, const char * id, const array_expressions_container &array_expressions, const std::vector<mapping_type> &, bool fallback) const
{ {
using std::string; using std::string;
using tools::to_string; using tools::to_string;
@@ -106,7 +106,7 @@ mproduct_parameters::mproduct_parameters(unsigned int simd_width
/// INIT /// INIT
/// ////////////// /// //////////////
kernel_generation_stream stream; kernel_generation_stream stream;
symbolic_expression const & st = (*symbolic_expressions.data().front()); array_expression const & st = (*array_expressions.data().front());
numeric_type dtype = lhs_most(st.tree(), st.root()).lhs.dtype; numeric_type dtype = lhs_most(st.tree(), st.root()).lhs.dtype;
std::string dtypestr = numeric_type_to_string(dtype); std::string dtypestr = numeric_type_to_string(dtype);
@@ -557,11 +557,11 @@ mproduct_parameters::mproduct_parameters(unsigned int simd_width
#undef VST0RE #undef VST0RE
} }
std::vector<std::string> mproduct::generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const std::vector<std::string> mproduct::generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const
{ {
std::vector<std::string> res; std::vector<std::string> res;
res.push_back(generate_impl(label, "o", symbolic_expressions, mappings, false)); res.push_back(generate_impl(label, "o", array_expressions, mappings, false));
res.push_back(generate_impl(label, "f", symbolic_expressions, mappings, true)); res.push_back(generate_impl(label, "f", array_expressions, mappings, true));
return res; return res;
} }
@@ -615,12 +615,12 @@ mproduct_parameters::mproduct_parameters(unsigned int simd_width
return result; return result;
} }
std::vector<int_t> mproduct::infos(symbolic_expressions_container const & symbolic_expressions, std::vector<int_t> mproduct::infos(array_expressions_container const & array_expressions,
lhs_rhs_element & C, lhs_rhs_element & A, lhs_rhs_element & B) lhs_rhs_element & C, lhs_rhs_element & A, lhs_rhs_element & B)
{ {
atidlas::symbolic_expression const & symbolic_expression = (*symbolic_expressions.data().front()); atidlas::array_expression const & array_expression = (*array_expressions.data().front());
symbolic_expression::container_type const & array = symbolic_expression.tree(); array_expression::container_type const & array = array_expression.tree();
std::size_t root = symbolic_expression.root(); std::size_t root = array_expression.root();
C = array[root].lhs; C = array[root].lhs;
@@ -640,18 +640,18 @@ mproduct_parameters::mproduct_parameters(unsigned int simd_width
mproduct::mproduct(mproduct_parameters const & parameters, char A_trans, char B_trans) : base_impl<mproduct, mproduct_parameters>(parameters, BIND_ALL_UNIQUE), A_trans_(A_trans), B_trans_(B_trans) mproduct::mproduct(mproduct_parameters const & parameters, char A_trans, char B_trans) : base_impl<mproduct, mproduct_parameters>(parameters, BIND_ALL_UNIQUE), A_trans_(A_trans), B_trans_(B_trans)
{ } { }
std::vector<int_t> mproduct::input_sizes(symbolic_expressions_container const & symbolic_expressions) std::vector<int_t> mproduct::input_sizes(array_expressions_container const & array_expressions)
{ {
lhs_rhs_element d0, d1, d2; lhs_rhs_element d0, d1, d2;
return infos(symbolic_expressions, d0, d1, d2); return infos(array_expressions, d0, d1, d2);
} }
void mproduct::enqueue(cl::CommandQueue & queue, std::vector<cl_ext::lazy_compiler> & programs, unsigned int label, symbolic_expressions_container const & symbolic_expressions) void mproduct::enqueue(cl::CommandQueue & queue, std::vector<cl_ext::lazy_compiler> & programs, unsigned int label, array_expressions_container const & array_expressions)
{ {
using namespace tools; using namespace tools;
lhs_rhs_element C, A, B; lhs_rhs_element C, A, B;
std::vector<int_t> MNK = infos(symbolic_expressions, C, A, B); std::vector<int_t> MNK = infos(array_expressions, C, A, B);
int_t M = MNK[0]; int_t M = MNK[0];
int_t N = MNK[1]; int_t N = MNK[1];

36
lib/backend/templates/mreduction.cpp
View File

@@ -14,7 +14,7 @@ mreduction_parameters::mreduction_parameters(unsigned int _simd_width,
num_groups_0(_num_groups_0), fetch_policy(_fetch_policy) { } num_groups_0(_num_groups_0), fetch_policy(_fetch_policy) { }
int mreduction::check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const int mreduction::check_invalid_impl(cl::Device const &, array_expressions_container const &) const
{ {
if (p_.fetch_policy==FETCH_FROM_LOCAL) if (p_.fetch_policy==FETCH_FROM_LOCAL)
return TEMPLATE_INVALID_FETCHING_POLICY_TYPE; return TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
@@ -26,7 +26,7 @@ unsigned int mreduction::lmem_usage() const
return p_.local_size_0*(p_.local_size_1+1); return p_.local_size_0*(p_.local_size_1+1);
} }
std::string mreduction::generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width, std::vector<mapped_mreduction*> const & exprs) const std::string mreduction::generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width, std::vector<mapped_mreduction*> const & exprs) const
{ {
using tools::to_string; using tools::to_string;
@@ -40,7 +40,7 @@ std::string mreduction::generate_impl(unsigned int label, symbolic_expressions_c
fill_kernel_name(kprefix, label, "d"); fill_kernel_name(kprefix, label, "d");
stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl; stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << "," << p_.local_size_1 << ",1)))" << std::endl;
stream << "__kernel void " << kprefix << "(unsigned int M, unsigned int N, " << generate_arguments("#scalartype", mappings, symbolic_expressions) << ")" << std::endl; stream << "__kernel void " << kprefix << "(unsigned int M, unsigned int N, " << generate_arguments("#scalartype", mappings, array_expressions) << ")" << std::endl;
stream << "{" << std::endl; stream << "{" << std::endl;
stream.inc_tab(); stream.inc_tab();
@@ -48,7 +48,7 @@ std::string mreduction::generate_impl(unsigned int label, symbolic_expressions_c
tools::make_map<std::map<std::string, std::string> >("array0", "#scalartype #namereg = #pointer[#start];") tools::make_map<std::map<std::string, std::string> >("array0", "#scalartype #namereg = #pointer[#start];")
("array1", "#pointer += #start;") ("array1", "#pointer += #start;")
("array2", "#pointer += #start1 + #start2*#ld; " ("array2", "#pointer += #start1 + #start2*#ld; "
"#ld *= #nldstride; "), symbolic_expressions, mappings); "#ld *= #nldstride; "), array_expressions, mappings);
for (std::vector<mapped_mreduction*>::const_iterator it = exprs.begin(); it != exprs.end(); ++it) for (std::vector<mapped_mreduction*>::const_iterator it = exprs.begin(); it != exprs.end(); ++it)
stream << (*it)->process("__local #scalartype #name_buf[" + to_string(lsize0*lsize1) + "];") << std::endl; stream << (*it)->process("__local #scalartype #name_buf[" + to_string(lsize0*lsize1) + "];") << std::endl;
@@ -160,7 +160,7 @@ std::string mreduction::generate_impl(unsigned int label, symbolic_expressions_c
std::map<std::string, std::string> accessors; std::map<std::string, std::string> accessors;
accessors["mreduction"] = "#name_buf[lid0*" + lsize1str + "]"; accessors["mreduction"] = "#name_buf[lid0*" + lsize1str + "]";
accessors["array1"] = "#pointer[r*#stride]"; accessors["array1"] = "#pointer[r*#stride]";
evaluate(stream, PARENT_NODE_TYPE, accessors, symbolic_expressions, mappings); evaluate(stream, PARENT_NODE_TYPE, accessors, array_expressions, mappings);
stream.dec_tab(); stream.dec_tab();
stream << "}" << std::endl; stream << "}" << std::endl;
@@ -174,14 +174,14 @@ std::string mreduction::generate_impl(unsigned int label, symbolic_expressions_c
return stream.str(); return stream.str();
} }
std::vector<std::string> mreduction::generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const std::vector<std::string> mreduction::generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const
{ {
std::vector<mapped_mreduction*> exprs; std::vector<mapped_mreduction*> exprs;
symbolic_expressions_container::data_type::const_iterator sit; array_expressions_container::data_type::const_iterator sit;
std::vector<mapping_type>::const_iterator mit; std::vector<mapping_type>::const_iterator mit;
for (mit = mappings.begin(), sit = symbolic_expressions.data().begin(); mit != mappings.end(); ++mit, ++sit) for (mit = mappings.begin(), sit = array_expressions.data().begin(); mit != mappings.end(); ++mit, ++sit)
{ {
symbolic_expression const & first_expression = *symbolic_expressions.data().front(); array_expression const & first_expression = *array_expressions.data().front();
std::vector<size_t> idx = filter_nodes(&is_reduction, first_expression, false); std::vector<size_t> idx = filter_nodes(&is_reduction, first_expression, false);
for (unsigned int j = 0; j < idx.size(); ++j) for (unsigned int j = 0; j < idx.size(); ++j)
exprs.push_back((mapped_mreduction*)(mit->at(mapping_key(idx[j], PARENT_NODE_TYPE)).get())); exprs.push_back((mapped_mreduction*)(mit->at(mapping_key(idx[j], PARENT_NODE_TYPE)).get()));
@@ -190,11 +190,11 @@ std::vector<std::string> mreduction::generate_impl(unsigned int label, symbolic_
std::vector<std::string> res; std::vector<std::string> res;
if (reduction_type_ && p_.simd_width>1) if (reduction_type_ && p_.simd_width>1)
{ {
res.push_back(generate_impl(label, symbolic_expressions, mappings, p_.simd_width, exprs)); res.push_back(generate_impl(label, array_expressions, mappings, p_.simd_width, exprs));
res.push_back(generate_impl(label, symbolic_expressions, mappings, 1, exprs)); res.push_back(generate_impl(label, array_expressions, mappings, 1, exprs));
} }
else else
res.push_back(generate_impl(label, symbolic_expressions, mappings, 1, exprs)); res.push_back(generate_impl(label, array_expressions, mappings, 1, exprs));
return res; return res;
} }
@@ -204,9 +204,9 @@ mreduction::mreduction(mreduction::parameters_type const & parameters,
base_impl<mreduction, mreduction_parameters>(parameters, binding_policy), base_impl<mreduction, mreduction_parameters>(parameters, binding_policy),
reduction_type_(rtype){ } reduction_type_(rtype){ }
std::vector<int_t> mreduction::input_sizes(symbolic_expressions_container const & symbolic_expressions) std::vector<int_t> mreduction::input_sizes(array_expressions_container const & array_expressions)
{ {
symbolic_expression const & first_expression = *symbolic_expressions.data().front(); array_expression const & first_expression = *array_expressions.data().front();
std::vector<std::size_t> idx = filter_nodes(&is_reduction, first_expression, false); std::vector<std::size_t> idx = filter_nodes(&is_reduction, first_expression, false);
std::pair<int_t, int_t> MN = matrix_size(lhs_most(first_expression.tree(), idx[0])); std::pair<int_t, int_t> MN = matrix_size(lhs_most(first_expression.tree(), idx[0]));
if(reduction_type_==REDUCE_COLUMNS) if(reduction_type_==REDUCE_COLUMNS)
@@ -217,15 +217,15 @@ std::vector<int_t> mreduction::input_sizes(symbolic_expressions_container const
void mreduction::enqueue(cl::CommandQueue & queue, void mreduction::enqueue(cl::CommandQueue & queue,
std::vector<cl_ext::lazy_compiler> & programs, std::vector<cl_ext::lazy_compiler> & programs,
unsigned int label, unsigned int label,
symbolic_expressions_container const & symbolic_expressions) array_expressions_container const & array_expressions)
{ {
char kname[10]; char kname[10];
fill_kernel_name(kname, label, "d"); fill_kernel_name(kname, label, "d");
std::vector<int_t> MN = input_sizes(symbolic_expressions); std::vector<int_t> MN = input_sizes(array_expressions);
//Kernel //Kernel
int idx = 0; int idx = 0;
if(reduction_type_==REDUCE_COLUMNS && p_.simd_width>1 && requires_fallback(symbolic_expressions)) if(reduction_type_==REDUCE_COLUMNS && p_.simd_width>1 && requires_fallback(array_expressions))
idx = 1; idx = 1;
cl::Program & program = programs[idx].program(); cl::Program & program = programs[idx].program();
cl::Kernel kernel(program, kname); cl::Kernel kernel(program, kname);
@@ -237,7 +237,7 @@ void mreduction::enqueue(cl::CommandQueue & queue,
unsigned int current_arg = 0; unsigned int current_arg = 0;
kernel.setArg(current_arg++, cl_uint(MN[0])); kernel.setArg(current_arg++, cl_uint(MN[0]));
kernel.setArg(current_arg++, cl_uint(MN[1])); kernel.setArg(current_arg++, cl_uint(MN[1]));
set_arguments(symbolic_expressions, kernel, current_arg); set_arguments(array_expressions, kernel, current_arg);
queue.enqueueNDRangeKernel(kernel, cl::NullRange, grange, lrange); queue.enqueueNDRangeKernel(kernel, cl::NullRange, grange, lrange);
} }

50
lib/backend/templates/reduction.cpp
View File

@@ -13,10 +13,10 @@ reduction_parameters::reduction_parameters(unsigned int _simd_width,
fetching_policy_type _fetching_policy) : base::parameters_type(_simd_width, _group_size, 1, 2), num_groups(_num_groups), fetching_policy(_fetching_policy) fetching_policy_type _fetching_policy) : base::parameters_type(_simd_width, _group_size, 1, 2), num_groups(_num_groups), fetching_policy(_fetching_policy)
{ } { }
unsigned int reduction::lmem_usage(symbolic_expressions_container const & symbolic_expressions) const unsigned int reduction::lmem_usage(array_expressions_container const & array_expressions) const
{ {
unsigned int res = 0; unsigned int res = 0;
for(symbolic_expressions_container::data_type::const_iterator it = symbolic_expressions.data().begin() ; it != symbolic_expressions.data().end() ; ++it) for(array_expressions_container::data_type::const_iterator it = array_expressions.data().begin() ; it != array_expressions.data().end() ; ++it)
{ {
numeric_type numeric_t= lhs_most((*it)->tree(), (*it)->root()).lhs.dtype; numeric_type numeric_t= lhs_most((*it)->tree(), (*it)->root()).lhs.dtype;
res += p_.local_size_0*size_of(numeric_t); res += p_.local_size_0*size_of(numeric_t);
@@ -24,7 +24,7 @@ unsigned int reduction::lmem_usage(symbolic_expressions_container const & symbol
return res; return res;
} }
int reduction::check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const int reduction::check_invalid_impl(cl::Device const &, array_expressions_container const &) const
{ {
if (p_.fetching_policy==FETCH_FROM_LOCAL) if (p_.fetching_policy==FETCH_FROM_LOCAL)
return TEMPLATE_INVALID_FETCHING_POLICY_TYPE; return TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
@@ -56,7 +56,7 @@ inline void reduction::reduce_1d_local_memory(kernel_generation_stream & stream,
stream << "}" << std::endl; stream << "}" << std::endl;
} }
std::string reduction::generate_impl(unsigned int label, const char * type, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width) const std::string reduction::generate_impl(unsigned int label, const char * type, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings, unsigned int simd_width) const
{ {
kernel_generation_stream stream; kernel_generation_stream stream;
@@ -70,8 +70,8 @@ std::string reduction::generate_impl(unsigned int label, const char * type, symb
std::string arguments = "unsigned int N, "; std::string arguments = "unsigned int N, ";
for (unsigned int k = 0; k < N; ++k) for (unsigned int k = 0; k < N; ++k)
{ {
std::string numeric_type = numeric_type_to_string(lhs_most(exprs[k]->symbolic_expression().tree(), std::string numeric_type = numeric_type_to_string(lhs_most(exprs[k]->array_expression().tree(),
exprs[k]->symbolic_expression().root()).lhs.dtype); exprs[k]->array_expression().root()).lhs.dtype);
if (exprs[k]->is_index_reduction()) if (exprs[k]->is_index_reduction())
{ {
arguments += exprs[k]->process("__global unsigned int* #name_temp, "); arguments += exprs[k]->process("__global unsigned int* #name_temp, ");
@@ -89,13 +89,13 @@ std::string reduction::generate_impl(unsigned int label, const char * type, symb
fill_kernel_name(kprefix, label, type); fill_kernel_name(kprefix, label, type);
stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << ",1,1)))" << std::endl; stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << ",1,1)))" << std::endl;
stream << "__kernel void " << kprefix << "0" << "(" << arguments << generate_arguments("#scalartype", mappings, symbolic_expressions) << ")" << std::endl; stream << "__kernel void " << kprefix << "0" << "(" << arguments << generate_arguments("#scalartype", mappings, array_expressions) << ")" << std::endl;
stream << "{" << std::endl; stream << "{" << std::endl;
stream.inc_tab(); stream.inc_tab();
stream << "unsigned int lid = get_local_id(0);" << std::endl; stream << "unsigned int lid = get_local_id(0);" << std::endl;
process(stream, PARENT_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array0", "#scalartype #namereg = #pointer[#start];") process(stream, PARENT_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array0", "#scalartype #namereg = #pointer[#start];")
("array1", "#pointer += #start;"), symbolic_expressions, mappings); ("array1", "#pointer += #start;"), array_expressions, mappings);
for (unsigned int k = 0; k < N; ++k) for (unsigned int k = 0; k < N; ++k)
{ {
@@ -194,7 +194,7 @@ std::string reduction::generate_impl(unsigned int label, const char * type, symb
* Second kernel * Second kernel
* -----------------------*/ * -----------------------*/
stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << ",1,1)))" << std::endl; stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << ",1,1)))" << std::endl;
stream << "__kernel void " << kprefix << "1" << "(" << arguments << generate_arguments("#scalartype", mappings, symbolic_expressions) << ")" << std::endl; stream << "__kernel void " << kprefix << "1" << "(" << arguments << generate_arguments("#scalartype", mappings, array_expressions) << ")" << std::endl;
stream << "{" << std::endl; stream << "{" << std::endl;
stream.inc_tab(); stream.inc_tab();
@@ -246,7 +246,7 @@ std::string reduction::generate_impl(unsigned int label, const char * type, symb
std::map<std::string, std::string> accessors; std::map<std::string, std::string> accessors;
accessors["scalar_reduction"] = "#name_buf[0]"; accessors["scalar_reduction"] = "#name_buf[0]";
accessors["array0"] = "#pointer[#start]"; accessors["array0"] = "#pointer[#start]";
evaluate(stream, PARENT_NODE_TYPE, accessors, symbolic_expressions, mappings); evaluate(stream, PARENT_NODE_TYPE, accessors, array_expressions, mappings);
stream.dec_tab(); stream.dec_tab();
stream << "}" << std::endl; stream << "}" << std::endl;
@@ -256,11 +256,11 @@ std::string reduction::generate_impl(unsigned int label, const char * type, symb
return stream.str(); return stream.str();
} }
std::vector<std::string> reduction::generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const std::vector<std::string> reduction::generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const
{ {
std::vector<std::string> result; std::vector<std::string> result;
result.push_back(generate_impl(label, "f", symbolic_expressions, mappings, 1)); result.push_back(generate_impl(label, "f", array_expressions, mappings, 1));
result.push_back(generate_impl(label, "o", symbolic_expressions, mappings, p_.simd_width)); result.push_back(generate_impl(label, "o", array_expressions, mappings, p_.simd_width));
return result; return result;
} }
@@ -273,22 +273,22 @@ reduction::reduction(unsigned int simd, unsigned int ls, unsigned int ng,
base_impl<reduction, reduction_parameters>(reduction_parameters(simd,ls,ng,fetch), bind) base_impl<reduction, reduction_parameters>(reduction_parameters(simd,ls,ng,fetch), bind)
{} {}
std::vector<int_t> reduction::input_sizes(symbolic_expressions_container const & symbolic_expressions) std::vector<int_t> reduction::input_sizes(array_expressions_container const & array_expressions)
{ {
std::vector<size_t> reductions_idx = filter_nodes(&is_reduction, *(symbolic_expressions.data().front()), false); std::vector<size_t> reductions_idx = filter_nodes(&is_reduction, *(array_expressions.data().front()), false);
int_t N = vector_size(lhs_most(symbolic_expressions.data().front()->tree(), reductions_idx[0])); int_t N = vector_size(lhs_most(array_expressions.data().front()->tree(), reductions_idx[0]));
return tools::make_vector<int_t>() << N; return tools::make_vector<int_t>() << N;
} }
void reduction::enqueue(cl::CommandQueue & queue, void reduction::enqueue(cl::CommandQueue & queue,
std::vector<cl_ext::lazy_compiler> & programs, std::vector<cl_ext::lazy_compiler> & programs,
unsigned int label, unsigned int label,
symbolic_expressions_container const & symbolic_expressions) array_expressions_container const & array_expressions)
{ {
//Preprocessing //Preprocessing
int_t size = input_sizes(symbolic_expressions)[0]; int_t size = input_sizes(array_expressions)[0];
std::vector<symbolic_expression_node const *> reductions; std::vector<array_expression::node const *> reductions;
for (symbolic_expressions_container::data_type::const_iterator it = symbolic_expressions.data().begin(); it != symbolic_expressions.data().end(); ++it) for (array_expressions_container::data_type::const_iterator it = array_expressions.data().begin(); it != array_expressions.data().end(); ++it)
{ {
std::vector<size_t> reductions_idx = filter_nodes(&is_reduction, **it, false); std::vector<size_t> reductions_idx = filter_nodes(&is_reduction, **it, false);
for (std::vector<size_t>::iterator itt = reductions_idx.begin(); itt != reductions_idx.end(); ++itt) for (std::vector<size_t>::iterator itt = reductions_idx.begin(); itt != reductions_idx.end(); ++itt)
@@ -303,7 +303,7 @@ void reduction::enqueue(cl::CommandQueue & queue,
fill_kernel_name(kopt[0], label, "o0"); fill_kernel_name(kopt[0], label, "o0");
fill_kernel_name(kopt[1], label, "o1"); fill_kernel_name(kopt[1], label, "o1");
bool fallback = p_.simd_width > 1 && (requires_fallback(symbolic_expressions) || (size%p_.simd_width>0)); bool fallback = p_.simd_width > 1 && (requires_fallback(array_expressions) || (size%p_.simd_width>0));
cl::Program & program = programs[fallback?0:1].program(); cl::Program & program = programs[fallback?0:1].program();
cl::Kernel kernels[2] = { cl::Kernel(program, fallback?kfallback[0]:kopt[0]), cl::Kernel kernels[2] = { cl::Kernel(program, fallback?kfallback[0]:kopt[0]),
cl::Kernel(program, fallback?kfallback[1]:kopt[1]) }; cl::Kernel(program, fallback?kfallback[1]:kopt[1]) };
@@ -313,8 +313,8 @@ void reduction::enqueue(cl::CommandQueue & queue,
cl::NDRange lrange[2] = { cl::NDRange(p_.local_size_0), cl::NDRange(p_.local_size_0) }; cl::NDRange lrange[2] = { cl::NDRange(p_.local_size_0), cl::NDRange(p_.local_size_0) };
//Arguments //Arguments
cl::Context context = symbolic_expressions.context(); cl::Context context = array_expressions.context();
symbolic_expression const & s = *(symbolic_expressions.data().front()); array_expression const & s = *(array_expressions.data().front());
unsigned int dtype_size = size_of(lhs_most(s.tree(), s.root()).lhs.dtype); unsigned int dtype_size = size_of(lhs_most(s.tree(), s.root()).lhs.dtype);
for (unsigned int k = 0; k < 2; k++) for (unsigned int k = 0; k < 2; k++)
{ {
@@ -324,7 +324,7 @@ void reduction::enqueue(cl::CommandQueue & queue,
//Temporary buffers //Temporary buffers
unsigned int i = 0; unsigned int i = 0;
unsigned int j = 0; unsigned int j = 0;
for (std::vector<symbolic_expression_node const *>::const_iterator it = reductions.begin(); it != reductions.end(); ++it) for (std::vector<array_expression::node const *>::const_iterator it = reductions.begin(); it != reductions.end(); ++it)
{ {
if (is_index_reduction((*it)->op)) if (is_index_reduction((*it)->op))
{ {
@@ -338,7 +338,7 @@ void reduction::enqueue(cl::CommandQueue & queue,
kernels[k].setArg(n_arg++, tmp_[i]); kernels[k].setArg(n_arg++, tmp_[i]);
i++; i++;
} }
set_arguments(symbolic_expressions, kernels[k], n_arg); set_arguments(array_expressions, kernels[k], n_arg);
} }
for (unsigned int k = 0; k < 2; k++) for (unsigned int k = 0; k < 2; k++)

29
lib/backend/templates/vaxpy.cpp
View File

@@ -16,14 +16,14 @@ vaxpy_parameters::vaxpy_parameters(unsigned int _simd_width,
{ } { }
int vaxpy::check_invalid_impl(cl::Device const &, symbolic_expressions_container const &) const int vaxpy::check_invalid_impl(cl::Device const &, array_expressions_container const &) const
{ {
if (p_.fetching_policy==FETCH_FROM_LOCAL) if (p_.fetching_policy==FETCH_FROM_LOCAL)
return TEMPLATE_INVALID_FETCHING_POLICY_TYPE; return TEMPLATE_INVALID_FETCHING_POLICY_TYPE;
return TEMPLATE_VALID; return TEMPLATE_VALID;
} }
std::vector<std::string> vaxpy::generate_impl(unsigned int label, symbolic_expressions_container const & symbolic_expressions, std::vector<mapping_type> const & mappings) const std::vector<std::string> vaxpy::generate_impl(unsigned int label, array_expressions_container const & array_expressions, std::vector<mapping_type> const & mappings) const
{ {
std::vector<std::string> result; std::vector<std::string> result;
for (unsigned int i = 0; i < 2; ++i) for (unsigned int i = 0; i < 2; ++i)
@@ -36,14 +36,14 @@ std::vector<std::string> vaxpy::generate_impl(unsigned int label, symbolic_expre
stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << ",1,1)))" << std::endl; stream << " __attribute__((reqd_work_group_size(" << p_.local_size_0 << ",1,1)))" << std::endl;
char kprefix[10]; char kprefix[10];
fill_kernel_name(kprefix, label, (i==0?"f":"o")); fill_kernel_name(kprefix, label, (i==0?"f":"o"));
stream << "__kernel void " << kprefix << "(unsigned int N," << generate_arguments(data_type, mappings, symbolic_expressions) << ")" << std::endl; stream << "__kernel void " << kprefix << "(unsigned int N," << generate_arguments(data_type, mappings, array_expressions) << ")" << std::endl;
stream << "{" << std::endl; stream << "{" << std::endl;
stream.inc_tab(); stream.inc_tab();
process(stream, PARENT_NODE_TYPE, process(stream, PARENT_NODE_TYPE,
tools::make_map<std::map<std::string, std::string> >("array0", "#scalartype #namereg = #pointer[#start];") tools::make_map<std::map<std::string, std::string> >("array0", "#scalartype #namereg = #pointer[#start];")
("array1", "#pointer += #start;") ("array1", "#pointer += #start;")
("array1", "#start1/=" + str_simd_width + ";"), symbolic_expressions, mappings); ("array1", "#start1/=" + str_simd_width + ";"), array_expressions, mappings);
std::string init, upper_bound, inc; std::string init, upper_bound, inc;
fetching_loop_info(p_.fetching_policy, "N/"+str_simd_width, stream, init, upper_bound, inc, "get_global_id(0)", "get_global_size(0)"); fetching_loop_info(p_.fetching_policy, "N/"+str_simd_width, stream, init, upper_bound, inc, "get_global_id(0)", "get_global_size(0)");
@@ -55,7 +55,7 @@ std::vector<std::string> vaxpy::generate_impl(unsigned int label, symbolic_expre
("matrix_row", "#scalartype #namereg = $VALUE{#row*#stride1, i*#stride2};") ("matrix_row", "#scalartype #namereg = $VALUE{#row*#stride1, i*#stride2};")
("matrix_column", "#scalartype #namereg = $VALUE{i*#stride1,#column*#stride2};") ("matrix_column", "#scalartype #namereg = $VALUE{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}];") ("matrix_diag", "#scalartype #namereg = #pointer[#diag_offset<0?$OFFSET{(i - #diag_offset)*#stride1, i*#stride2}:$OFFSET{i*#stride1, (i + #diag_offset)*#stride2}];")
, symbolic_expressions, mappings); , array_expressions, mappings);
evaluate(stream, PARENT_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array1", "#namereg") evaluate(stream, PARENT_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array1", "#namereg")
("matrix_row", "#namereg") ("matrix_row", "#namereg")
@@ -63,13 +63,13 @@ std::vector<std::string> vaxpy::generate_impl(unsigned int label, symbolic_expre
("matrix_diag", "#namereg") ("matrix_diag", "#namereg")
("array0", "#namereg") ("array0", "#namereg")
("cast", "convert_"+data_type) ("cast", "convert_"+data_type)
, symbolic_expressions, mappings); , array_expressions, mappings);
process(stream, LHS_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array1", "#pointer[i*#stride] = #namereg;") process(stream, LHS_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array1", "#pointer[i*#stride] = #namereg;")
("matrix_row", "$VALUE{#row, i} = #namereg;") ("matrix_row", "$VALUE{#row, i} = #namereg;")
("matrix_column", "$VALUE{i, #column} = #namereg;") ("matrix_column", "$VALUE{i, #column} = #namereg;")
("matrix_diag", "#diag_offset<0?$VALUE{(i - #diag_offset)*#stride1, i*#stride2}:$VALUE{i*#stride1, (i + #diag_offset)*#stride2} = #namereg;") ("matrix_diag", "#diag_offset<0?$VALUE{(i - #diag_offset)*#stride1, i*#stride2}:$VALUE{i*#stride1, (i + #diag_offset)*#stride2} = #namereg;")
,symbolic_expressions, mappings); ,array_expressions, mappings);
stream.dec_tab(); stream.dec_tab();
stream << "}" << std::endl; stream << "}" << std::endl;
@@ -77,7 +77,7 @@ std::vector<std::string> vaxpy::generate_impl(unsigned int label, symbolic_expre
stream << "if(get_global_id(0)==0)" << std::endl; stream << "if(get_global_id(0)==0)" << std::endl;
stream << "{" << std::endl; stream << "{" << std::endl;
stream.inc_tab(); stream.inc_tab();
process(stream, LHS_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array0", "#pointer[#start] = #namereg;"), symbolic_expressions, mappings); process(stream, LHS_NODE_TYPE, tools::make_map<std::map<std::string, std::string> >("array0", "#pointer[#start] = #namereg;"), array_expressions, mappings);
stream.dec_tab(); stream.dec_tab();
stream << "}" << std::endl; stream << "}" << std::endl;
@@ -102,25 +102,25 @@ vaxpy::vaxpy(unsigned int simd, unsigned int ls, unsigned int ng,
{} {}
std::vector<int_t> vaxpy::input_sizes(symbolic_expressions_container const & symbolic_expressions) std::vector<int_t> vaxpy::input_sizes(array_expressions_container const & array_expressions)
{ {
int_t size = static_cast<array_expression const *>(symbolic_expressions.data().front().get())->shape()._1; int_t size = static_cast<array_expression const *>(array_expressions.data().front().get())->shape()._1;
return tools::make_vector<int_t>() << size; return tools::make_vector<int_t>() << size;
} }
void vaxpy::enqueue(cl::CommandQueue & queue, void vaxpy::enqueue(cl::CommandQueue & queue,
std::vector<cl_ext::lazy_compiler> & programs, std::vector<cl_ext::lazy_compiler> & programs,
unsigned int label, unsigned int label,
symbolic_expressions_container const & symbolic_expressions) array_expressions_container const & array_expressions)
{ {
//Size //Size
int_t size = input_sizes(symbolic_expressions)[0]; int_t size = input_sizes(array_expressions)[0];
//Kernel //Kernel
char kfb[10]; char kfb[10];
char kopt[10]; char kopt[10];
fill_kernel_name(kfb, label, "f"); fill_kernel_name(kfb, label, "f");
fill_kernel_name(kopt, label, "o"); fill_kernel_name(kopt, label, "o");
bool fallback = p_.simd_width > 1 && (requires_fallback(symbolic_expressions) || (size%p_.simd_width>0)); bool fallback = p_.simd_width > 1 && (requires_fallback(array_expressions) || (size%p_.simd_width>0));
cl::Program const & program = programs[fallback?0:1].program(); cl::Program const & program = programs[fallback?0:1].program();
cl_ext::kernels_t::key_type key(program(), label); cl_ext::kernels_t::key_type key(program(), label);
@@ -135,8 +135,9 @@ void vaxpy::enqueue(cl::CommandQueue & queue,
//Arguments //Arguments
unsigned int current_arg = 0; unsigned int current_arg = 0;
kernel.setArg(current_arg++, cl_uint(size)); kernel.setArg(current_arg++, cl_uint(size));
set_arguments(symbolic_expressions, kernel, current_arg); set_arguments(array_expressions, kernel, current_arg);
queue.enqueueNDRangeKernel(kernel, cl::NullRange, grange, lrange); queue.enqueueNDRangeKernel(kernel, cl::NullRange, grange, lrange);
queue.flush();
} }

46
lib/model/model.cpp
View File

@@ -28,9 +28,9 @@ std::string model::define_extension(std::string const & extensions, std::string
return std::string(""); return std::string("");
} }
void model::fill_program_name(char* program_name, symbolic_expressions_container const & symbolic_expressions, binding_policy_t binding_policy) void model::fill_program_name(char* program_name, array_expressions_container const & array_expressions, binding_policy_t binding_policy)
{ {
if (symbolic_expressions.order()==symbolic_expressions_container::INDEPENDENT) if (array_expressions.order()==array_expressions_container::INDEPENDENT)
*program_name++='i'; *program_name++='i';
else else
*program_name++='s'; *program_name++='s';
@@ -39,16 +39,16 @@ void model::fill_program_name(char* program_name, symbolic_expressions_container
binder = new bind_to_handle(); binder = new bind_to_handle();
else else
binder = new bind_all_unique(); binder = new bind_all_unique();
for (symbolic_expressions_container::data_type::const_iterator it = symbolic_expressions.data().begin(); it != symbolic_expressions.data().end(); ++it) for (array_expressions_container::data_type::const_iterator it = array_expressions.data().begin(); it != array_expressions.data().end(); ++it)
traverse(**it, (*it)->root(), symbolic_expression_representation_functor(*binder, program_name),true); traverse(**it, (*it)->root(), array_expression_representation_functor(*binder, program_name),true);
*program_name='\0'; *program_name='\0';
delete binder; delete binder;
} }
std::vector<cl_ext::lazy_compiler>& model::init(symbolic_expressions_container const & symbolic_expressions, cl::Context const & context, cl::Device const & device, bool force_recompilation) std::vector<cl_ext::lazy_compiler>& model::init(array_expressions_container const & array_expressions, cl::Context const & context, cl::Device const & device, bool force_recompilation)
{ {
char program_name[256]; char program_name[256];
fill_program_name(program_name, symbolic_expressions, BIND_TO_HANDLE); fill_program_name(program_name, array_expressions, BIND_TO_HANDLE);
std::string pname(program_name); std::string pname(program_name);
std::vector<cl_ext::lazy_compiler> & to_init = lazy_programs_[context()][pname]; std::vector<cl_ext::lazy_compiler> & to_init = lazy_programs_[context()][pname];
if(to_init.empty()) if(to_init.empty())
@@ -63,7 +63,7 @@ std::vector<cl_ext::lazy_compiler>& model::init(symbolic_expressions_container c
for(size_t i = 0 ; i < templates_.size() ; ++i) for(size_t i = 0 ; i < templates_.size() ; ++i)
{ {
std::vector<std::string> cur = templates_[i]->generate(i, symbolic_expressions, device); std::vector<std::string> cur = templates_[i]->generate(i, array_expressions, device);
for(size_t j = 0 ; j < cur.size() ; ++j){ for(size_t j = 0 ; j < cur.size() ; ++j){
to_init[j].add(cur[j]); to_init[j].add(cur[j]);
} }
@@ -82,17 +82,17 @@ model::model(std::vector< tools::shared_ptr<base> > const & templates, cl::Comma
model::model(base const & tp, cl::CommandQueue & queue) : templates_(1,tp.clone()), queue_(queue) model::model(base const & tp, cl::CommandQueue & queue) : templates_(1,tp.clone()), queue_(queue)
{} {}
void model::execute(symbolic_expressions_container const & symbolic_expressions, bool bypass_predictor, bool force_recompilation) void model::execute(array_expressions_container const & array_expressions, bool bypass_predictor, bool force_recompilation)
{ {
bypass_predictor = bypass_predictor || predictor_.get()==NULL; bypass_predictor = bypass_predictor || predictor_.get()==NULL;
cl::Context const & context = symbolic_expressions.context(); cl::Context const & context = array_expressions.context();
assert(context() == queue_.getInfo<CL_QUEUE_CONTEXT>()()); assert(context() == queue_.getInfo<CL_QUEUE_CONTEXT>()());
cl::Device const & device = queue_.getInfo<CL_QUEUE_DEVICE>(); cl::Device const & device = queue_.getInfo<CL_QUEUE_DEVICE>();
std::vector<cl_ext::lazy_compiler> & compilers = init(symbolic_expressions, context, device, force_recompilation); std::vector<cl_ext::lazy_compiler> & compilers = init(array_expressions, context, device, force_recompilation);
//Prediction //Prediction
std::vector<int_t> x = templates_[0]->input_sizes(symbolic_expressions); std::vector<int_t> x = templates_[0]->input_sizes(array_expressions);
int label; int label;
//The user tuned the model specifically for this input size //The user tuned the model specifically for this input size
if(hardcoded_.find(x)!=hardcoded_.end()) if(hardcoded_.find(x)!=hardcoded_.end())
@@ -108,16 +108,16 @@ void model::execute(symbolic_expressions_container const & symbolic_expressions,
} }
//Execution //Execution
templates_[label]->enqueue(queue_, compilers, label, symbolic_expressions); templates_[label]->enqueue(queue_, compilers, label, array_expressions);
} }
void model::tune(symbolic_expressions_container const & symbolic_expressions) void model::tune(array_expressions_container const & array_expressions)
{ {
cl::Context const & context = symbolic_expressions.context(); cl::Context const & context = array_expressions.context();
assert(context() == queue_.getInfo<CL_QUEUE_CONTEXT>()()); assert(context() == queue_.getInfo<CL_QUEUE_CONTEXT>()());
cl::Device device = queue_.getInfo<CL_QUEUE_DEVICE>(); cl::Device device = queue_.getInfo<CL_QUEUE_DEVICE>();
std::vector<cl_ext::lazy_compiler> & compilers = init(symbolic_expressions, context, device, false); std::vector<cl_ext::lazy_compiler> & compilers = init(array_expressions, context, device, false);
//Collect the timings //Collect the timings
std::vector<float> timings(templates_.size()); std::vector<float> timings(templates_.size());
@@ -125,13 +125,13 @@ void model::tune(symbolic_expressions_container const & symbolic_expressions)
for(size_t i = 0 ; i < templates_.size() ; ++i) for(size_t i = 0 ; i < templates_.size() ; ++i)
{ {
timer.start(); timer.start();
templates_[i]->enqueue(queue_, compilers, i, symbolic_expressions); templates_[i]->enqueue(queue_, compilers, i, array_expressions);
queue_.finish(); queue_.finish();
timings[i] = timer.get(); timings[i] = timer.get();
} }
//Fill the override //Fill the override
std::vector<int_t> x = templates_[0]->input_sizes(symbolic_expressions); std::vector<int_t> x = templates_[0]->input_sizes(array_expressions);
hardcoded_[x] = std::distance(timings.begin(),std::min_element(timings.begin(), timings.end())); hardcoded_[x] = std::distance(timings.begin(),std::min_element(timings.begin(), timings.end()));
} }
@@ -203,9 +203,7 @@ void import(std::string const & fname, cl::CommandQueue & queue, model_map_t& re
str.assign((std::istreambuf_iterator<char>(t)), std::istreambuf_iterator<char>()); str.assign((std::istreambuf_iterator<char>(t)), std::istreambuf_iterator<char>());
document.Parse<0>(str.c_str()); document.Parse<0>(str.c_str());
//Deserialize //Deserialize
std::vector<std::string> operations = tools::make_vector<std::string>() << "vaxpy" << "dot" std::vector<std::string> operations = tools::make_vector<std::string>() << "vaxpy" << "dot" << "maxpy" << "gemvN" << "gemvT" << "gemmNN" << "gemmTN" << "gemmTT";
<< "maxpy" << "gemvN" << "gemvT"
<< "gemmNN" << "gemmTN" << "gemmTT";
std::vector<std::string> dtype = tools::make_vector<std::string>() << "float32" << "float64"; std::vector<std::string> dtype = tools::make_vector<std::string>() << "float32" << "float64";
for(std::vector<std::string>::iterator op = operations.begin() ; op != operations.end() ; ++op) for(std::vector<std::string>::iterator op = operations.begin() ; op != operations.end() ; ++op)
{ {
@@ -258,13 +256,9 @@ model_map_t init_models(cl::CommandQueue & queue)
res[std::make_pair(MATRIX_PRODUCT_NT_TYPE, DTYPE)] = ptr_t(new model(mproduct_nt(1, 8, 8, 8, 4, 1, 4, FETCH_FROM_LOCAL, FETCH_FROM_LOCAL, 8, 8), queue)); res[std::make_pair(MATRIX_PRODUCT_NT_TYPE, DTYPE)] = ptr_t(new model(mproduct_nt(1, 8, 8, 8, 4, 1, 4, FETCH_FROM_LOCAL, FETCH_FROM_LOCAL, 8, 8), queue));
res[std::make_pair(MATRIX_PRODUCT_TT_TYPE, DTYPE)] = ptr_t(new model(mproduct_tt(1, 8, 8, 8, 4, 1, 4, FETCH_FROM_LOCAL, FETCH_FROM_LOCAL, 8, 8), queue)); res[std::make_pair(MATRIX_PRODUCT_TT_TYPE, DTYPE)] = ptr_t(new model(mproduct_tt(1, 8, 8, 8, 4, 1, 4, FETCH_FROM_LOCAL, FETCH_FROM_LOCAL, 8, 8), queue));
} }
if(const char * cmodel_file = std::getenv("ATIDLAS_MODEL_DEVICE_0")) // if(const char * cmodel_file = std::getenv("ATIDLAS_MODEL_DEVICE_0"))
import(std::string(cmodel_file), queue, res); // import(std::string(cmodel_file), queue, res);
return res; return res;
// else
// throw "Please specify a model file";
} }
model_map_t& get_model_map(cl::CommandQueue & queue) model_map_t& get_model_map(cl::CommandQueue & queue)

33
lib/symbolic/execute.cpp
View File

@@ -116,12 +116,12 @@ namespace atidlas
/** @brief Parses the breakpoints for a given expression tree */ /** @brief Parses the breakpoints for a given expression tree */
static void parse(symbolic_expression::container_type&array, size_t idx, static void parse(array_expression::container_type&array, size_t idx,
expression_type current_type, expression_type current_type,
breakpoints_t & breakpoints, breakpoints_t & breakpoints,
expression_type & final_type) expression_type & final_type)
{ {
symbolic_expression_node & node = array[idx]; array_expression::node & node = array[idx];
if (node.lhs.type_family == COMPOSITE_OPERATOR_FAMILY) if (node.lhs.type_family == COMPOSITE_OPERATOR_FAMILY)
{ {
std::pair<bool, expression_type> breakpoint = is_breakpoint(current_type, array[node.lhs.node_index].op); std::pair<bool, expression_type> breakpoint = is_breakpoint(current_type, array[node.lhs.node_index].op);
@@ -146,13 +146,13 @@ namespace atidlas
} }
/** @brief Executes a symbolic_expression on the given models map*/ /** @brief Executes a array_expression on the given models map*/
void execute(atidlas::symbolic_expression & symbolic_expression, model_map_t & models) void execute(atidlas::array_expression & array_expression, model_map_t & models)
{ {
cl::Context const & context = symbolic_expression.context(); cl::Context const & context = array_expression.context();
size_t rootidx = symbolic_expression.root(); size_t rootidx = array_expression.root();
symbolic_expression::container_type & tree = const_cast<symbolic_expression::container_type &>(symbolic_expression.tree()); array_expression::container_type & tree = const_cast<array_expression::container_type &>(array_expression.tree());
symbolic_expression_node root_save = tree[rootidx]; array_expression::node root_save = tree[rootidx];
//Todo: technically the datatype should be per temporary //Todo: technically the datatype should be per temporary
numeric_type dtype = root_save.lhs.dtype; numeric_type dtype = root_save.lhs.dtype;
@@ -178,8 +178,8 @@ namespace atidlas
for(detail::breakpoints_t::reverse_iterator rit = breakpoints.rbegin() ; rit != breakpoints.rend() ; ++rit) for(detail::breakpoints_t::reverse_iterator rit = breakpoints.rbegin() ; rit != breakpoints.rend() ; ++rit)
{ {
tools::shared_ptr<model> const & pmodel = models[std::make_pair(rit->first, dtype)]; tools::shared_ptr<model> const & pmodel = models[std::make_pair(rit->first, dtype)];
symbolic_expression_node const & node = tree[rit->second->node_index]; array_expression::node const & node = tree[rit->second->node_index];
symbolic_expression_node const & lmost = lhs_most(tree, node); array_expression::node const & lmost = lhs_most(tree, node);
//Creates temporary //Creates temporary
tools::shared_ptr<obj_base> tmp; tools::shared_ptr<obj_base> tmp;
@@ -202,23 +202,20 @@ namespace atidlas
temporaries_.push_back(tmp); temporaries_.push_back(tmp);
tree[rootidx].op.type = OPERATOR_ASSIGN_TYPE; tree[rootidx].op.type = OPERATOR_ASSIGN_TYPE;
tree[rootidx].lhs = lhs_rhs_element((array const &)*tmp); fill(tree[rootidx].lhs, (array&)*tmp);
tree[rootidx].rhs = *rit->second; tree[rootidx].rhs = *rit->second;
tree[rootidx].rhs.type_family = rit->second->type_family; tree[rootidx].rhs.type_family = rit->second->type_family;
//Execute //Execute
pmodel->execute(symbolic_expression); pmodel->execute(array_expression);
tree[rootidx] = root_save; tree[rootidx] = root_save;
//Incorporates the temporary within the symbolic_expression //Incorporates the temporary within the array_expression
rit->second->dtype = dtype; fill(*rit->second, (array&)*tmp);
rit->second->type_family = ARRAY_TYPE_FAMILY;
rit->second->subtype = DENSE_ARRAY_TYPE;
fill((array&)*tmp, rit->second->array);
} }
/*-----Compute final expression-----*/ /*-----Compute final expression-----*/
models[std::make_pair(final_type, dtype)]->execute(symbolic_expression); models[std::make_pair(final_type, dtype)]->execute(array_expression);
} }
} }

228
lib/symbolic/expression.cpp
View File

@@ -8,107 +8,101 @@
namespace atidlas namespace atidlas
{ {
void fill(array const & a, array_infos& i) void fill(lhs_rhs_element &x, invalid_node)
{ {
i.dtype = a.dtype(); x.type_family = INVALID_TYPE_FAMILY;
i.data = a.data()(); x.subtype = INVALID_SUBTYPE;
i.shape1 = a.shape()._1; x.dtype = INVALID_NUMERIC_TYPE;
i.shape2 = a.shape()._2;
i.start1 = a.start()._1;
i.start2 = a.start()._2;
i.stride1 = a.stride()._1;
i.stride2 = a.stride()._2;
i.ld = a.ld();
} }
array_expression array_expression::operator-() void fill(lhs_rhs_element & x, unsigned int node_index)
{ return array_expression(*this, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_SUB_TYPE), dtype_, shape_); }
array_expression array_expression::operator!()
{ return array_expression(*this, lhs_rhs_element(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_NEGATE_TYPE), INT_TYPE, shape_); }
lhs_rhs_element::lhs_rhs_element()
{ {
type_family = INVALID_TYPE_FAMILY; x.type_family = COMPOSITE_OPERATOR_FAMILY;
subtype = INVALID_SUBTYPE; x.subtype = INVALID_SUBTYPE;
dtype = INVALID_NUMERIC_TYPE; x.dtype = INVALID_NUMERIC_TYPE;
x.node_index = node_index;
} }
lhs_rhs_element::lhs_rhs_element(unsigned int _node_index) void fill(lhs_rhs_element & x, array const & a)
{ {
type_family = COMPOSITE_OPERATOR_FAMILY; x.type_family = ARRAY_TYPE_FAMILY;
subtype = INVALID_SUBTYPE; x.subtype = DENSE_ARRAY_TYPE;
dtype = INVALID_NUMERIC_TYPE; x.dtype = a.dtype();
node_index = _node_index;
x.array.dtype = a.dtype();
x.array.data = a.data()();
x.array.shape1 = a.shape()._1;
x.array.shape2 = a.shape()._2;
x.array.start1 = a.start()._1;
x.array.start2 = a.start()._2;
x.array.stride1 = a.stride()._1;
x.array.stride2 = a.stride()._2;
x.array.ld = a.ld();
x.memory_ = a.data();
} }
lhs_rhs_element::lhs_rhs_element(atidlas::array const & x) void fill(lhs_rhs_element & x, value_scalar const & v)
{ {
type_family = ARRAY_TYPE_FAMILY; x.type_family = VALUE_TYPE_FAMILY;
subtype = DENSE_ARRAY_TYPE; x.dtype = v.dtype();
dtype = x.dtype(); x.subtype = VALUE_SCALAR_TYPE;
fill(x, array); x.vscalar = v.values();
memory_ = x.data();
} }
lhs_rhs_element::lhs_rhs_element(atidlas::value_scalar const & x) void fill(lhs_rhs_element & x, repeat_infos const & r)
{ {
type_family = VALUE_TYPE_FAMILY; x.type_family = INFOS_TYPE_FAMILY;
dtype = x.dtype(); x.subtype = REPEAT_INFOS_TYPE;
subtype = VALUE_SCALAR_TYPE; x.dtype = INVALID_NUMERIC_TYPE;
vscalar = x.values(); x.tuple = r;
} }
lhs_rhs_element::lhs_rhs_element(atidlas::repeat_infos const & x) lhs_rhs_element::lhs_rhs_element(){}
{
type_family = INFOS_TYPE_FAMILY;
subtype = REPEAT_INFOS_TYPE;
dtype = INVALID_NUMERIC_TYPE;
tuple = x;
}
// //
op_element::op_element(operation_node_type_family const & _type_family, operation_node_type const & _type) : op_element::op_element() {}
type_family(_type_family), type(_type) op_element::op_element(operation_node_type_family const & _type_family, operation_node_type const & _type) : type_family(_type_family), type(_type){}
{ }
// //
symbolic_expression_node::symbolic_expression_node(lhs_rhs_element const & _lhs, op_element const& _op, lhs_rhs_element const & _rhs) : template<class LT, class RT>
lhs(_lhs), op(_op), rhs(_rhs) array_expression::array_expression(LT const & lhs, RT const & rhs, op_element const & op, cl::Context const & context, numeric_type const & dtype, size4 const & shape) :
{ } tree_(1), root_(0), context_(context), dtype_(dtype), shape_(shape)
//
symbolic_expression::symbolic_expression(lhs_rhs_element const & lhs, lhs_rhs_element const & rhs, op_element const & op, cl::Context const & context, numeric_type const & dtype) :
tree_(1, symbolic_expression_node(lhs, op, rhs)), root_(0), context_(context), dtype_(dtype)
{ }
symbolic_expression::symbolic_expression(symbolic_expression const & lhs, lhs_rhs_element const & rhs, op_element const & op, numeric_type const & dtype) :
context_(lhs.context_), dtype_(dtype)
{ {
tree_.reserve(lhs.tree_.size() + 1); fill(tree_[0].lhs, lhs);
tree_.insert(tree_.end(), lhs.tree_.begin(), lhs.tree_.end()); tree_[0].op = op;
tree_.push_back(value_type(lhs_rhs_element(lhs.root_), op, rhs)); fill(tree_[0].rhs, rhs);
root_ = tree_.size() - 1;
} }
symbolic_expression::symbolic_expression(lhs_rhs_element const & lhs, symbolic_expression const & rhs, op_element const & op, numeric_type const & dtype) : template<class RT>
context_(rhs.context_), dtype_(dtype) array_expression::array_expression(array_expression const & lhs, RT const & rhs, op_element const & op, numeric_type const & dtype, size4 const & shape) :
tree_(lhs.tree_.size() + 1), root_(tree_.size()-1), context_(lhs.context_), dtype_(dtype), shape_(shape)
{ {
tree_.reserve(rhs.tree_.size() + 1); std::copy(lhs.tree_.begin(), lhs.tree_.end(), tree_.begin());
tree_.insert(tree_.end(), rhs.tree_.begin(), rhs.tree_.end()); fill(tree_[root_].lhs, lhs.root_);
tree_.push_back(value_type(lhs, op, lhs_rhs_element(rhs.root_))); tree_[root_].op = op;
root_ = tree_.size() - 1; fill(tree_[root_].rhs, rhs);
} }
symbolic_expression::symbolic_expression(symbolic_expression const & lhs, symbolic_expression const & rhs, op_element const & op, numeric_type const & dtype): template<class LT>
context_(lhs.context_), dtype_(dtype) array_expression::array_expression(LT const & lhs, array_expression const & rhs, op_element const & op, numeric_type const & dtype, size4 const & shape) :
tree_(rhs.tree_.size() + 1), root_(tree_.size() - 1), context_(rhs.context_), dtype_(dtype), shape_(shape)
{
std::copy(rhs.tree_.begin(), rhs.tree_.end(), tree_.begin());
fill(tree_[root_].lhs, lhs);
tree_[root_].op = op;
fill(tree_[root_].rhs, rhs.root_);
}
array_expression::array_expression(array_expression const & lhs, array_expression const & rhs, op_element const & op, numeric_type const & dtype, size4 const & shape):
tree_(lhs.tree_.size() + rhs.tree_.size() + 1), root_(tree_.size()-1), context_(lhs.context_), dtype_(dtype), shape_(shape)
{ {
std::size_t lsize = lhs.tree_.size(); std::size_t lsize = lhs.tree_.size();
std::size_t rsize = rhs.tree_.size(); std::copy(lhs.tree_.begin(), lhs.tree_.end(), tree_.begin());
tree_.reserve(lsize + rsize + 1); std::copy(rhs.tree_.begin(), rhs.tree_.end(), tree_.begin() + lsize);
tree_.insert(tree_.end(), lhs.tree_.begin(), lhs.tree_.end()); fill(tree_[root_].lhs, lhs.root_);
tree_.insert(tree_.end(), rhs.tree_.begin(), rhs.tree_.end()); tree_[root_].op = op;
tree_.push_back(value_type(lhs_rhs_element(lhs.root_), op, lhs_rhs_element(lsize+rhs.root_))); fill(tree_[root_].rhs, lsize + rhs.root_);
for(container_type::iterator it = tree_.begin() + lsize ; it != tree_.end() - 1 ; ++it){ for(container_type::iterator it = tree_.begin() + lsize ; it != tree_.end() - 1 ; ++it){
if(it->lhs.type_family==COMPOSITE_OPERATOR_FAMILY) it->lhs.node_index+=lsize; if(it->lhs.type_family==COMPOSITE_OPERATOR_FAMILY) it->lhs.node_index+=lsize;
if(it->rhs.type_family==COMPOSITE_OPERATOR_FAMILY) it->rhs.node_index+=lsize; if(it->rhs.type_family==COMPOSITE_OPERATOR_FAMILY) it->rhs.node_index+=lsize;
@@ -116,39 +110,51 @@ symbolic_expression::symbolic_expression(symbolic_expression const & lhs, symbol
root_ = tree_.size() - 1; root_ = tree_.size() - 1;
} }
symbolic_expression::container_type & symbolic_expression::tree() template array_expression::array_expression(array_expression const &, value_scalar const &, op_element const &, numeric_type const &, size4 const &);
template array_expression::array_expression(array_expression const &, invalid_node const &, op_element const &, numeric_type const &, size4 const &);
template array_expression::array_expression(array_expression const &, array const &, op_element const &, numeric_type const &, size4 const &);
template array_expression::array_expression(array_expression const &, repeat_infos const &, op_element const &, numeric_type const &, size4 const &);
template array_expression::array_expression(value_scalar const &, array_expression const &, op_element const &, numeric_type const &, size4 const &);
template array_expression::array_expression(invalid_node const &, array_expression const &, op_element const &, numeric_type const &, size4 const &);
template array_expression::array_expression(array const &, array_expression const &, op_element const &, numeric_type const &, size4 const &);
template array_expression::array_expression(repeat_infos const &, array_expression const &, op_element const &, numeric_type const &, size4 const &);
template array_expression::array_expression(value_scalar const &, value_scalar const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(invalid_node const &, value_scalar const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(array const &, value_scalar const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(repeat_infos const &, value_scalar const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(value_scalar const &, invalid_node const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(invalid_node const &, invalid_node const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(array const &, invalid_node const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(repeat_infos const &, invalid_node const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(value_scalar const &, array const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(invalid_node const &, array const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(array const &, array const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(repeat_infos const &, array const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(value_scalar const &, repeat_infos const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(invalid_node const &, repeat_infos const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(array const &, repeat_infos const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
template array_expression::array_expression(repeat_infos const &, repeat_infos const &, op_element const &, cl::Context const &, numeric_type const &, size4 const &);
array_expression::container_type & array_expression::tree()
{ return tree_; } { return tree_; }
symbolic_expression::container_type const & symbolic_expression::tree() const array_expression::container_type const & array_expression::tree() const
{ return tree_; } { return tree_; }
std::size_t symbolic_expression::root() const std::size_t array_expression::root() const
{ return root_; } { return root_; }
cl::Context const & symbolic_expression::context() const cl::Context const & array_expression::context() const
{ return context_; } { return context_; }
numeric_type const & symbolic_expression::dtype() const numeric_type const & array_expression::dtype() const
{ return dtype_; } { return dtype_; }
//
array_expression::array_expression(lhs_rhs_element const & lhs, lhs_rhs_element const & rhs, op_element const & op, cl::Context const & ctx, numeric_type const & dtype, size4 shape):
symbolic_expression(lhs, rhs, op, ctx, dtype), shape_(shape)
{ }
array_expression::array_expression(symbolic_expression const & lhs, lhs_rhs_element const & rhs, op_element const & op, numeric_type const & dtype, size4 shape):
symbolic_expression(lhs, rhs, op, dtype), shape_(shape)
{ }
array_expression::array_expression(lhs_rhs_element const & lhs, symbolic_expression const & rhs, op_element const & op, numeric_type const & dtype, size4 shape):
symbolic_expression(lhs, rhs, op, dtype), shape_(shape)
{ }
array_expression::array_expression(symbolic_expression const & lhs, symbolic_expression const & rhs, op_element const & op, numeric_type const & dtype, size4 shape):
symbolic_expression(lhs, rhs, op, dtype), shape_(shape)
{ }
size4 array_expression::shape() const size4 array_expression::shape() const
{ return shape_; } { return shape_; }
@@ -162,45 +168,51 @@ array_expression& array_expression::reshape(int_t size1, int_t size2)
return *this; return *this;
} }
array_expression array_expression::operator-()
{ return array_expression(*this, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_SUB_TYPE), dtype_, shape_); }
array_expression array_expression::operator!()
{ return array_expression(*this, invalid_node(), op_element(OPERATOR_UNARY_TYPE_FAMILY, OPERATOR_NEGATE_TYPE), INT_TYPE, shape_); }
// //
tools::shared_ptr<symbolic_expression> symbolic_expressions_container::create(symbolic_expression const & s) tools::shared_ptr<array_expression> array_expressions_container::create(array_expression const & s)
{ {
return tools::shared_ptr<symbolic_expression>(new array_expression(static_cast<array_expression const &>(s))); return tools::shared_ptr<array_expression>(new array_expression(static_cast<array_expression const &>(s)));
} }
symbolic_expressions_container::symbolic_expressions_container(data_type const & data, order_type order) : data_(data), order_(order) array_expressions_container::array_expressions_container(data_type const & data, order_type order) : data_(data), order_(order)
{ } { }
symbolic_expressions_container::symbolic_expressions_container(symbolic_expression const & s0) : order_(INDEPENDENT) array_expressions_container::array_expressions_container(array_expression const & s0) : order_(INDEPENDENT)
{ {
data_.push_back(create(s0)); data_.push_back(create(s0));
} }
symbolic_expressions_container::symbolic_expressions_container(order_type order, symbolic_expression const & s0, symbolic_expression const & s1) : order_(order) array_expressions_container::array_expressions_container(order_type order, array_expression const & s0, array_expression const & s1) : order_(order)
{ {
data_.push_back(create(s0)); data_.push_back(create(s0));
data_.push_back(create(s1)); data_.push_back(create(s1));
} }
symbolic_expressions_container::data_type const & symbolic_expressions_container::data() const array_expressions_container::data_type const & array_expressions_container::data() const
{ return data_; } { return data_; }
cl::Context const & symbolic_expressions_container::context() const cl::Context const & array_expressions_container::context() const
{ return data_.front()->context(); } { return data_.front()->context(); }
symbolic_expressions_container::order_type symbolic_expressions_container::order() const array_expressions_container::order_type array_expressions_container::order() const
{ return order_; } { return order_; }
symbolic_expression_node const & lhs_most(symbolic_expression::container_type const & array, symbolic_expression_node const & init) array_expression::node const & lhs_most(array_expression::container_type const & array, array_expression::node const & init)
{ {
symbolic_expression_node const * current = &init; array_expression::node const * current = &init;
while (current->lhs.type_family==COMPOSITE_OPERATOR_FAMILY) while (current->lhs.type_family==COMPOSITE_OPERATOR_FAMILY)
current = &array[current->lhs.node_index]; current = &array[current->lhs.node_index];
return *current; return *current;
} }
symbolic_expression_node const & lhs_most(symbolic_expression::container_type const & array, size_t root) array_expression::node const & lhs_most(array_expression::container_type const & array, size_t root)
{ return lhs_most(array, array[root]); } { return lhs_most(array, array[root]); }

14
lib/symbolic/io.cpp
View File

@@ -9,7 +9,7 @@ namespace atidlas
#define ATIDLAS_MAP_TO_STRING(NAME) case NAME: return #NAME #define ATIDLAS_MAP_TO_STRING(NAME) case NAME: return #NAME
inline std::string to_string(symbolic_expression_node_subtype const & f) inline std::string to_string(array_expression_node_subtype const & f)
{ {
switch(f) switch(f)
{ {
@@ -29,7 +29,7 @@ inline std::string to_string(lhs_rhs_element const & e)
return tools::to_string(e.subtype); return tools::to_string(e.subtype);
} }
inline std::ostream & operator<<(std::ostream & os, symbolic_expression_node const & s_node) inline std::ostream & operator<<(std::ostream & os, array_expression::node const & s_node)
{ {
os << "LHS: " << to_string(s_node.lhs) << "|" << s_node.lhs.dtype << ", " os << "LHS: " << to_string(s_node.lhs) << "|" << s_node.lhs.dtype << ", "
<< "OP: " << s_node.op.type_family << " | " << s_node.op.type << ", " << "OP: " << s_node.op.type_family << " | " << s_node.op.type << ", "
@@ -41,11 +41,11 @@ inline std::ostream & operator<<(std::ostream & os, symbolic_expression_node con
namespace detail namespace detail
{ {
/** @brief Recursive worker routine for printing a whole symbolic_expression */ /** @brief Recursive worker routine for printing a whole array_expression */
inline void print_node(std::ostream & os, atidlas::symbolic_expression const & s, size_t node_index, size_t indent = 0) inline void print_node(std::ostream & os, atidlas::array_expression const & s, size_t node_index, size_t indent = 0)
{ {
symbolic_expression::container_type const & nodes = s.tree(); array_expression::container_type const & nodes = s.tree();
symbolic_expression_node const & current_node = nodes[node_index]; array_expression::node const & current_node = nodes[node_index];
for (size_t i=0; i<indent; ++i) for (size_t i=0; i<indent; ++i)
os << " "; os << " ";
@@ -60,7 +60,7 @@ namespace detail
} }
} }
std::string to_string(atidlas::symbolic_expression const & s) std::string to_string(atidlas::array_expression const & s)
{ {
std::ostringstream os; std::ostringstream os;
detail::print_node(os, s, s.root()); detail::print_node(os, s, s.root());

BIN
python/autotune/pysrc/dataset.pyc
View File

Binary file not shown.

6
python/autotune/pysrc/genetic.py
View File

@@ -99,9 +99,9 @@ class GeneticOperators(object):
bincode = bincode + [str(random.randint(0,1)) for i in range(x)] bincode = bincode + [str(random.randint(0,1)) for i in range(x)]
parameters = self.decode(bincode) parameters = self.decode(bincode)
template = self.Template(*parameters) template = self.Template(*parameters)
symbolic_expressions = atd.symbolic_expression_container(self.symbolic) array_expressions = atd.array_expression_container(self.symbolic)
registers_usage = template.registers_usage(symbolic_expressions)/4 registers_usage = template.registers_usage(array_expressions)/4
lmem_usage = template.lmem_usage(symbolic_expressions) lmem_usage = template.lmem_usage(array_expressions)
local_size = parameters[1]*parameters[3] local_size = parameters[1]*parameters[3]
occupancy_record = misc_tools.OccupancyRecord(self.device, local_size, lmem_usage, registers_usage) occupancy_record = misc_tools.OccupancyRecord(self.device, local_size, lmem_usage, registers_usage)
if not misc_tools.skip(template, self.symbolic): if not misc_tools.skip(template, self.symbolic):

BIN
python/autotune/pysrc/genetic.pyc
View File

Binary file not shown.

14
python/autotune/pysrc/misc_tools.py
View File

@@ -197,21 +197,21 @@ class OccupancyRecord:
def skip(template, symbolic): def skip(template, symbolic):
device = symbolic.context.queues[0].device device = symbolic.context.queues[0].device
symbolic_expressions = atd.symbolic_expression_container(symbolic) array_expressions = atd.array_expression_container(symbolic)
registers_usage = template.registers_usage(symbolic_expressions)/4 registers_usage = template.registers_usage(array_expressions)/4
lmem_usage = template.lmem_usage(symbolic_expressions) lmem_usage = template.lmem_usage(array_expressions)
local_size = template.local_size_0*template.local_size_1 local_size = template.local_size_0*template.local_size_1
occupancy_record = OccupancyRecord(device, local_size, lmem_usage, registers_usage) occupancy_record = OccupancyRecord(device, local_size, lmem_usage, registers_usage)
if template.check_invalid(symbolic_expressions, device) or occupancy_record.occupancy < 15: if template.check_invalid(array_expressions, device) or occupancy_record.occupancy < 15:
return True return True
return False return False
def benchmark(template, symbolic): def benchmark(template, symbolic):
queue = symbolic.context.queues[0] queue = symbolic.context.queues[0]
device = queue.device device = queue.device
symbolic_expressions = atd.symbolic_expression_container(symbolic) array_expressions = atd.array_expression_container(symbolic)
registers_usage = template.registers_usage(symbolic_expressions)/4 registers_usage = template.registers_usage(array_expressions)/4
lmem_usage = template.lmem_usage(symbolic_expressions) lmem_usage = template.lmem_usage(array_expressions)
local_size = template.local_size_0*template.local_size_1 local_size = template.local_size_0*template.local_size_1
occupancy_record = OccupancyRecord(device, local_size, lmem_usage, registers_usage) occupancy_record = OccupancyRecord(device, local_size, lmem_usage, registers_usage)
if occupancy_record.occupancy < 15 : if occupancy_record.occupancy < 15 :

BIN
python/autotune/pysrc/misc_tools.pyc
View File

Binary file not shown.

BIN
python/autotune/pysrc/model.pyc
View File

Binary file not shown.

BIN
python/autotune/pysrc/optimize.pyc
View File

Binary file not shown.

2
python/postinstall.py
View File

@@ -31,6 +31,6 @@ if retune=='y':
print '----------------' print '----------------'
opts += ['--blas3-size'] + [add_input('BLAS3 sizes (M,N,K)', '1024,1024,1024').split(',')] opts += ['--blas3-size'] + [add_input('BLAS3 sizes (M,N,K)', '1024,1024,1024').split(',')]
print '----------------' print '----------------'
subprocess.call(["${CMAKE_BINARY_DIR}/python/autotune/dist/autotune", "tune"] + opts) subprocess.call(["${CMAKE_BINARY_DIR}/python/autotune/dist/autotune", "tune"] + opts +['--json', 'tmp.json'])

18
python/pyatidlas/src/_atidlas.cpp
View File

@@ -153,16 +153,6 @@ void export_core()
; ;
} }
void export_symbolic()
{
bp::class_<atd::symbolic_expressions_container>
("symbolic_expression_container", bp::init<atd::symbolic_expression const &>())
;
bp::class_<atd::symbolic_expression>("symbolic_expression", bp::no_init)
.add_property("context", bp::make_function(&atd::symbolic_expression::context, bp::return_internal_reference<>()))
;
}
namespace detail namespace detail
{ {
@@ -487,7 +477,11 @@ void export_array()
.def(bp::self OP bp::self)\ .def(bp::self OP bp::self)\
ADD_SCALAR_HANDLING(OP) ADD_SCALAR_HANDLING(OP)
bp::class_<atd::array_expression, bp::bases<atd::symbolic_expression> >("array_expression", bp::no_init) bp::class_<atd::array_expressions_container>
("array_expression_container", bp::init<atd::array_expression const &>())
;
bp::class_<atd::array_expression >("array_expression", bp::no_init)
ADD_ARRAY_OPERATOR(+) ADD_ARRAY_OPERATOR(+)
ADD_ARRAY_OPERATOR(-) ADD_ARRAY_OPERATOR(-)
ADD_ARRAY_OPERATOR(*) ADD_ARRAY_OPERATOR(*)
@@ -498,6 +492,7 @@ void export_array()
ADD_ARRAY_OPERATOR(<=) ADD_ARRAY_OPERATOR(<=)
ADD_ARRAY_OPERATOR(==) ADD_ARRAY_OPERATOR(==)
ADD_ARRAY_OPERATOR(!=) ADD_ARRAY_OPERATOR(!=)
.add_property("context", bp::make_function(&atd::array_expression::context, bp::return_internal_reference<>()))
.def(bp::self_ns::abs(bp::self)) .def(bp::self_ns::abs(bp::self))
// .def(bp::self_ns::pow(bp::self)) // .def(bp::self_ns::pow(bp::self))
; ;
@@ -657,6 +652,5 @@ BOOST_PYTHON_MODULE(_atidlas)
export_core(); export_core();
export_cl(); export_cl();
export_model(); export_model();
export_symbolic();
export_array(); export_array();
} }