ViennaCL - The Vienna Computing Library  1.5.2
vector_element.hpp
Go to the documentation of this file.
1 #ifndef VIENNACL_LINALG_OPENCL_KERNELS_VECTOR_ELEMENT_HPP
2 #define VIENNACL_LINALG_OPENCL_KERNELS_VECTOR_ELEMENT_HPP
3 
7 #include "viennacl/ocl/utils.hpp"
8 
11 namespace viennacl
12 {
13  namespace linalg
14  {
15  namespace opencl
16  {
17  namespace kernels
18  {
19 
21 
22 
23  //generate code for C = op1(A) * op2(B), where A, B, C can have different storage layouts and opX(D) = D or trans(D)
24  template <typename StringType>
25  void generate_vector_unary_element_ops(StringType & source, std::string const & numeric_string,
26  std::string const & funcname, std::string const & op, std::string const & op_name)
27  {
28  source.append("__kernel void "); source.append(funcname); source.append("_"); source.append(op_name); source.append("(\n");
29  source.append(" __global "); source.append(numeric_string); source.append(" * vec1, \n");
30  source.append(" uint4 size1, \n");
31  source.append(" __global "); source.append(numeric_string); source.append(" * vec2, \n");
32  source.append(" uint4 size2) { \n");
33  source.append(" for (unsigned int i = get_global_id(0); i < size1.z; i += get_global_size(0)) \n");
34  source.append(" vec1[i*size1.y+size1.x] "); source.append(op); source.append(" "); source.append(funcname); source.append("(vec2[i*size2.y+size2.x]); \n");
35  source.append("} \n");
36  }
37 
38  template <typename StringType>
39  void generate_vector_unary_element_ops(StringType & source, std::string const & numeric_string, std::string const & funcname)
40  {
41  generate_vector_unary_element_ops(source, numeric_string, funcname, "=", "assign");
42  //generate_vector_unary_element_ops(source, numeric_string, funcname, "+=", "plus");
43  //generate_vector_unary_element_ops(source, numeric_string, funcname, "-=", "minus");
44  }
45 
46  template <typename StringType>
47  void generate_vector_binary_element_ops(StringType & source, std::string const & numeric_string)
48  {
49  // generic kernel for the vector operation v1 = alpha * v2 + beta * v3, where v1, v2, v3 are not necessarily distinct vectors
50  source.append("__kernel void element_op( \n");
51  source.append(" __global "); source.append(numeric_string); source.append(" * vec1, \n");
52  source.append(" unsigned int start1, \n");
53  source.append(" unsigned int inc1, \n");
54  source.append(" unsigned int size1, \n");
55 
56  source.append(" __global const "); source.append(numeric_string); source.append(" * vec2, \n");
57  source.append(" unsigned int start2, \n");
58  source.append(" unsigned int inc2, \n");
59 
60  source.append(" __global const "); source.append(numeric_string); source.append(" * vec3, \n");
61  source.append(" unsigned int start3, \n");
62  source.append(" unsigned int inc3, \n");
63 
64  source.append(" unsigned int op_type) \n"); //0: product, 1: division, 2: power
65  source.append("{ \n");
66  if (numeric_string == "float" || numeric_string == "double")
67  {
68  source.append(" if (op_type == 2) \n");
69  source.append(" { \n");
70  source.append(" for (unsigned int i = get_global_id(0); i < size1; i += get_global_size(0)) \n");
71  source.append(" vec1[i*inc1+start1] = pow(vec2[i*inc2+start2], vec3[i*inc3+start3]); \n");
72  source.append(" } else ");
73  }
74  source.append(" if (op_type == 1) \n");
75  source.append(" { \n");
76  source.append(" for (unsigned int i = get_global_id(0); i < size1; i += get_global_size(0)) \n");
77  source.append(" vec1[i*inc1+start1] = vec2[i*inc2+start2] / vec3[i*inc3+start3]; \n");
78  source.append(" } \n");
79  source.append(" else if (op_type == 0)\n");
80  source.append(" { \n");
81  source.append(" for (unsigned int i = get_global_id(0); i < size1; i += get_global_size(0)) \n");
82  source.append(" vec1[i*inc1+start1] = vec2[i*inc2+start2] * vec3[i*inc3+start3]; \n");
83  source.append(" } \n");
84  source.append("} \n");
85  }
86 
88 
89  // main kernel class
91  template <class TYPE>
93  {
94  static std::string program_name()
95  {
96  return viennacl::ocl::type_to_string<TYPE>::apply() + "_vector_element";
97  }
98 
99  static void init(viennacl::ocl::context & ctx)
100  {
102  std::string numeric_string = viennacl::ocl::type_to_string<TYPE>::apply();
103 
104  static std::map<cl_context, bool> init_done;
105  if (!init_done[ctx.handle().get()])
106  {
107  std::string source;
108  source.reserve(8192);
109 
110  viennacl::ocl::append_double_precision_pragma<TYPE>(ctx, source);
111 
112  // unary operations
113  if (numeric_string == "float" || numeric_string == "double")
114  {
115  generate_vector_unary_element_ops(source, numeric_string, "acos");
116  generate_vector_unary_element_ops(source, numeric_string, "asin");
117  generate_vector_unary_element_ops(source, numeric_string, "atan");
118  generate_vector_unary_element_ops(source, numeric_string, "ceil");
119  generate_vector_unary_element_ops(source, numeric_string, "cos");
120  generate_vector_unary_element_ops(source, numeric_string, "cosh");
121  generate_vector_unary_element_ops(source, numeric_string, "exp");
122  generate_vector_unary_element_ops(source, numeric_string, "fabs");
123  generate_vector_unary_element_ops(source, numeric_string, "floor");
124  generate_vector_unary_element_ops(source, numeric_string, "log");
125  generate_vector_unary_element_ops(source, numeric_string, "log10");
126  generate_vector_unary_element_ops(source, numeric_string, "sin");
127  generate_vector_unary_element_ops(source, numeric_string, "sinh");
128  generate_vector_unary_element_ops(source, numeric_string, "sqrt");
129  generate_vector_unary_element_ops(source, numeric_string, "tan");
130  generate_vector_unary_element_ops(source, numeric_string, "tanh");
131  }
132  else
133  {
134  generate_vector_unary_element_ops(source, numeric_string, "abs");
135  }
136 
137  // binary operations
138  generate_vector_binary_element_ops(source, numeric_string);
139 
140  std::string prog_name = program_name();
141  #ifdef VIENNACL_BUILD_INFO
142  std::cout << "Creating program " << prog_name << std::endl;
143  #endif
144  ctx.add_program(source, prog_name);
145  init_done[ctx.handle().get()] = true;
146  } //if
147  } //init
148  };
149 
150  } // namespace kernels
151  } // namespace opencl
152  } // namespace linalg
153 } // namespace viennacl
154 #endif
155 
static std::string program_name()
Definition: vector_element.hpp:94
Implements a OpenCL platform within ViennaCL.
void generate_vector_unary_element_ops(StringType &source, std::string const &numeric_string, std::string const &funcname, std::string const &op, std::string const &op_name)
Definition: vector_element.hpp:25
Various little tools used here and there in ViennaCL.
Manages an OpenCL context and provides the respective convenience functions for creating buffers...
Definition: context.hpp:51
Provides OpenCL-related utilities.
static void init(viennacl::ocl::context &ctx)
Definition: vector_element.hpp:99
void generate_vector_binary_element_ops(StringType &source, std::string const &numeric_string)
Definition: vector_element.hpp:47
const viennacl::ocl::handle< cl_context > & handle() const
Returns the context handle.
Definition: context.hpp:476
const OCL_TYPE & get() const
Definition: handle.hpp:189
static void apply(viennacl::ocl::context const &)
Definition: utils.hpp:40
Main kernel class for generating OpenCL kernels for elementwise operations other than addition and su...
Definition: vector_element.hpp:92
Representation of an OpenCL kernel in ViennaCL.
Helper class for converting a type to its string representation.
Definition: utils.hpp:57