void core(vcl_size_t /*kernel_id*/, utils::kernel_generation_stream& stream, statements_type const & statements, std::vector<detail::mapping_type> const & mapping) const {
stream << "for(unsigned int i = get_global_id(0) ; i < N ; i += get_global_size(0))" << std::endl;
stream << "{" << std::endl;
stream.inc_tab();
//Fetches entries to registers
std::set<std::string> fetched;
for(std::vector<detail::mapping_type>::const_iterator it = mapping.begin() ; it != mapping.end() ; ++it)
for(detail::mapping_type::const_reverse_iterator iit = it->rbegin() ; iit != it->rend() ; ++iit)
//Useless to fetch cpu scalars into registers
if(detail::mapped_handle * p = dynamic_cast<detail::mapped_handle *>(iit->second.get()))
p->fetch( std::make_pair("i","0"), vector_size_, fetched, stream);
//Generates all the expression, in order
vcl_size_t i = 0;
for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it){
std::string str;
detail::traverse(it->first, it->second, detail::expression_generation_traversal(std::make_pair("i","0"), -1, str, mapping[i++]));
stream << str << ";" << std::endl;
}
//Writes back
for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it)
//Gets the mapped object at the LHS of each expression
if(detail::mapped_handle * p = dynamic_cast<detail::mapped_handle *>(at(mapping.at(std::distance(statements.begin(),it)), std::make_pair(&it->second, detail::LHS_NODE_TYPE)).get()))
p->write_back( std::make_pair("i", "0"), fetched, stream);
stream.dec_tab();
stream << "}" << std::endl;
}
ref_ptr<statement> fetch(std::string const &query)
{
ref_ptr<statement> st;
statements_type::iterator p = statements.find(query);
if(p==statements.end())
return st;
st=p->second.stat;
lru.erase(p->second.lru_ptr);
statements.erase(p);
size --;
return st;
}
void configure_range_enqueue_arguments(vcl_size_t kernel_id, statements_type const & statements, viennacl::ocl::kernel & k, unsigned int & n_arg) const{
configure_local_sizes(k, kernel_id);
k.global_work_size(0,local_size_1_*num_groups_);
k.global_work_size(1,1);
scheduler::statement_node const & first_node = statements.front().second;
viennacl::vcl_size_t N = utils::call_on_vector(first_node.lhs, utils::internal_size_fun());
k.arg(n_arg++, cl_uint(N/vector_size_));
}
void configure_range_enqueue_arguments(std::size_t kernel_id, statements_type const & statements, viennacl::ocl::kernel & k, unsigned int & n_arg) const{
configure_local_sizes(k, kernel_id);
k.global_work_size(0,group_size_row_*num_groups_row_);
k.global_work_size(1,group_size_col_*num_groups_col_);
scheduler::statement_node const & first_node = statements.front().second;
k.arg(n_arg++, cl_uint(utils::call_on_matrix(first_node.lhs, utils::internal_size1_fun())));
k.arg(n_arg++, cl_uint(utils::call_on_matrix(first_node.lhs, utils::internal_size2_fun())));
}
void core(vcl_size_t /*kernel_id*/, utils::kernel_generation_stream& stream, statements_type const & statements, std::vector<detail::mapping_type> const & mapping) const {
for(std::vector<detail::mapping_type>::const_iterator it = mapping.begin() ; it != mapping.end() ; ++it){
for(detail::mapping_type::const_iterator iit = it->begin() ; iit != it->end() ; ++iit){
if(detail::mapped_matrix * p = dynamic_cast<detail::mapped_matrix*>(iit->second.get()))
p->bind_sizes("M","N");
}
}
stream << "for(unsigned int i = get_global_id(0) ; i < M ; i += get_global_size(0))" << std::endl;
stream << "{" << std::endl;
stream.inc_tab();
stream << "for(unsigned int j = get_global_id(1) ; j < N ; j += get_global_size(1))" << std::endl;
stream << "{" << std::endl;
stream.inc_tab();
//Fetches entries to registers
std::set<std::string> fetched;
for(std::vector<detail::mapping_type>::const_iterator it = mapping.begin() ; it != mapping.end() ; ++it)
for(detail::mapping_type::const_reverse_iterator it2 = it->rbegin() ; it2 != it->rend() ; ++it2)
if(detail::mapped_matrix * p = dynamic_cast<detail::mapped_matrix *>(it2->second.get()))
p->fetch(std::make_pair("i", "j"), vector_size_, fetched, stream);
vcl_size_t i = 0;
for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it){
std::string str;
detail::traverse(it->first, it->second, detail::expression_generation_traversal(std::make_pair("i", "j"), -1, str, mapping[i++]));
stream << str << ";" << std::endl;
}
//Writes back
for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it){
if(detail::mapped_handle * p = dynamic_cast<detail::mapped_handle *>(at(mapping.at(std::distance(statements.begin(),it)), std::make_pair(&it->second,detail::LHS_NODE_TYPE)).get()))
p->write_back(std::make_pair("i", "j"), fetched, stream);
}
stream.dec_tab();
stream << "}" << std::endl;
stream.dec_tab();
stream << "}" << std::endl;
}
virtual void operator()(utils::kernel_generation_stream & stream, std::size_t device_offset, statements_type const & statements) const {
std::vector<detail::mapping_type> mapping(statements.size());
///Get Prototype, initialize mapping
std::string prototype;
std::set<std::string> already_generated;
kernel_arguments(statements, prototype);
{
std::map<void *, std::size_t> memory;
unsigned int current_arg = 0;
std::size_t i = 0;
for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it)
detail::traverse(it->first, it->second, detail::map_functor(memory,current_arg,mapping[i++]));
}
for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it){
detail::traverse(it->first, it->second, detail::prototype_generation_traversal(already_generated, prototype, vectorization(), mapping[std::distance(statements.begin(), it)]));
}
prototype.erase(prototype.size()-1); //Last comma pruned
//Generate
for(std::size_t n = 0 ; n < num_kernels() ; ++n){
//stream << "__attribute__((vec_type_hint()))" << std::endl;
stream << " __attribute__((reqd_work_group_size(" << local_size_1_ << "," << local_size_2_ << "," << 1 << ")))" << std::endl;
stream << "__kernel " << "void " << "kernel_" << device_offset << "_" << n << "(" << std::endl;
stream << prototype << std::endl;
stream << ")" << std::endl;
//core:
stream << "{" << std::endl;
stream.inc_tab();
core(n, stream, statements, mapping);
stream.dec_tab();
stream << "}" << std::endl;
}
}
void insert(ref_ptr<statement> st)
{
statements_type::iterator p;
if((p=statements.find(st->sql_query()))!=statements.end()) {
p->second.stat = st;
lru.erase(p->second.lru_ptr);
lru.push_front(p);
p->second.lru_ptr = lru.begin();
}
else {
if(size > 0 && size >= max_size) {
statements.erase(lru.back());
lru.pop_back();
size--;
}
std::pair<statements_type::iterator, bool> ins =
statements.insert(std::make_pair(st->sql_query(), entry()));
p = ins.first;
p->second.stat = st;
lru.push_front(p);
p->second.lru_ptr = lru.begin();
size ++;
}
}
void clear()
{
lru.clear();
statements.clear();
size=0;
}
