1 #ifndef VIENNACL_GENERATOR_GENERATE_SAXPY_HPP
2 #define VIENNACL_GENERATOR_GENERATE_SAXPY_HPP
47 return "Vec,LSize1,NumGroups1,GlobalDecomposition";
51 std::ostringstream oss;
55 <<
"," << decomposition_;
72 arguments_string += detail::generate_value_kernel_argument(
"unsigned int",
"N");
78 stream <<
"for(unsigned int i = get_global_id(0) ; i < N ; i += get_global_size(0))" << std::endl;
79 stream <<
"{" << std::endl;
83 std::set<std::string> fetched;
84 for(std::vector<detail::mapping_type>::const_iterator it = mapping.begin() ; it != mapping.end() ; ++it)
85 for(detail::mapping_type::const_reverse_iterator iit = it->rbegin() ; iit != it->rend() ; ++iit)
88 p->fetch( std::make_pair(
"i",
"0"),
vector_size_, fetched, stream);
92 for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it){
95 stream << str <<
";" << std::endl;
99 for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it)
101 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()))
102 p->write_back( std::make_pair(
"i",
"0"), fetched, stream);
105 stream <<
"}" << std::endl;
110 unsigned int decomposition_;
123 matrix_saxpy(
unsigned int v,
vcl_size_t gs1,
vcl_size_t gs2,
vcl_size_t ng1,
vcl_size_t ng2,
unsigned int d) :
profile_base(v, gs1, gs2, 1), num_groups_row_(ng1), num_groups_col_(ng2), decomposition_(d){ }
126 return "Vec,LSize1,LSize2,NumGroups1,NumGroups2,GlobalDecomposition";
130 std::ostringstream oss;
134 <<
"," << num_groups_row_
135 <<
"," << num_groups_col_
136 <<
"," << decomposition_;
152 arguments_string += detail::generate_value_kernel_argument(
"unsigned int",
"M");
153 arguments_string += detail::generate_value_kernel_argument(
"unsigned int",
"N");
159 for(std::vector<detail::mapping_type>::const_iterator it = mapping.begin() ; it != mapping.end() ; ++it){
160 for(detail::mapping_type::const_iterator iit = it->begin() ; iit != it->end() ; ++iit){
162 p->bind_sizes(
"M",
"N");
166 stream <<
"for(unsigned int i = get_global_id(0) ; i < M ; i += get_global_size(0))" << std::endl;
167 stream <<
"{" << std::endl;
169 stream <<
"for(unsigned int j = get_global_id(1) ; j < N ; j += get_global_size(1))" << std::endl;
170 stream <<
"{" << std::endl;
174 std::set<std::string> fetched;
175 for(std::vector<detail::mapping_type>::const_iterator it = mapping.begin() ; it != mapping.end() ; ++it)
176 for(detail::mapping_type::const_reverse_iterator it2 = it->rbegin() ; it2 != it->rend() ; ++it2)
177 if(detail::mapped_matrix * p = dynamic_cast<detail::mapped_matrix *>(it2->second.get()))
178 p->fetch(std::make_pair(
"i",
"j"),
vector_size_, fetched, stream);
182 for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it){
184 detail::traverse(it->first, it->second, detail::expression_generation_traversal(std::make_pair(
"i",
"j"), -1, str, mapping[i++]));
185 stream << str <<
";" << std::endl;
189 for(statements_type::const_iterator it = statements.begin() ; it != statements.end() ; ++it){
190 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()))
191 p->write_back(std::make_pair(
"i",
"j"), fetched, stream);
195 stream <<
"}" << std::endl;
197 stream <<
"}" << std::endl;
204 unsigned int decomposition_;
A stream class where the kernel sources are streamed to. Takes care of indentation of the sources...
Definition: utils.hpp:233
void arg(unsigned int pos, cl_char val)
Sets a char argument at the provided position.
Definition: kernel.hpp:124
void configure_range_enqueue_arguments(vcl_size_t kernel_id, statements_type const &statements, viennacl::ocl::kernel &k, unsigned int &n_arg) const
Configures the range and enqueues the arguments associated with the profile.
Definition: saxpy.hpp:61
std::size_t vcl_size_t
Definition: forwards.h:58
OpenCL kernel generation class for matrix expressions of AXPY type, i.e. A = alpha * B + beta * C...
Definition: saxpy.hpp:117
static std::string csv_format()
Definition: saxpy.hpp:125
Internal utils for a dynamic OpenCL kernel generation.
void configure_range_enqueue_arguments(vcl_size_t kernel_id, statements_type const &statements, viennacl::ocl::kernel &k, unsigned int &n_arg) const
Configures the range and enqueues the arguments associated with the profile.
Definition: saxpy.hpp:140
vector_saxpy(unsigned int v, vcl_size_t gs, vcl_size_t ng, unsigned int d)
Definition: saxpy.hpp:59
Represents an OpenCL kernel within ViennaCL.
Definition: kernel.hpp:59
Base class for an operation profile.
Definition: profile_base.hpp:47
lhs_rhs_element lhs
Definition: forwards.h:422
A class representing a compute device (e.g. a GPU)
Definition: device.hpp:49
Base class for datastructures passed by pointer.
Definition: mapped_objects.hpp:133
Mapping of a matrix to a generator class.
Definition: mapped_objects.hpp:236
Functor for obtaining the internal number of columns of a ViennaCL matrix.
Definition: utils.hpp:188
std::list< std::pair< scheduler::statement, scheduler::statement_node > > statements_type
Definition: profile_base.hpp:49
ValueT const & at(std::map< KeyT, ValueT > const &map, KeyT const &key)
Emulation of C++11's .at() member for std::map<>
Definition: forwards.h:97
static std::string csv_format()
Definition: saxpy.hpp:46
Main namespace in ViennaCL. Holds all the basic types such as vector, matrix, etc. and defines operations upon them.
Definition: cpu_ram.hpp:29
several code generation helpers
OpenCL kernel generation class for vector expressions of AXPY type, i.e. x = alpha * y + beta * z...
Definition: saxpy.hpp:44
Base classes for the profiles.
void kernel_arguments(statements_type const &, std::string &arguments_string) const
Definition: saxpy.hpp:71
Map ViennaCL objects to generator wrappers.
void kernel_arguments(statements_type const &, std::string &arguments_string) const
Definition: saxpy.hpp:151
Functor for obtaining the internal number of rows of a ViennaCL matrix.
Definition: utils.hpp:181
Definition: forwards.h:113
std::string csv_representation() const
csv representation of an operation
Definition: saxpy.hpp:129
void configure_local_sizes(viennacl::ocl::kernel &k, vcl_size_t) const
Definition: profile_base.hpp:59
void dec_tab()
Definition: utils.hpp:259
Provides the datastructures for dealing with a single statement such as 'x = y + z;'.
unsigned int vector_size_
Definition: profile_base.hpp:178
matrix_saxpy(unsigned int v, vcl_size_t gs1, vcl_size_t gs2, vcl_size_t ng1, vcl_size_t ng2, unsigned int d)
Definition: saxpy.hpp:123
void inc_tab()
Definition: utils.hpp:257
vcl_size_t local_size_2_
Definition: profile_base.hpp:180
size_type global_work_size(int index=0) const
Returns the global work size at the respective dimension.
Definition: kernel.hpp:759
vcl_size_t local_size_1_
Definition: profile_base.hpp:179
functor for generating the expression string from a statement
Definition: helpers.hpp:224
std::string csv_representation() const
csv representation of an operation
Definition: saxpy.hpp:50
Main datastructure for an node in the statement tree.
Definition: forwards.h:420
Functor for returning the internal size of a vector.
Definition: utils.hpp:167