1 #ifndef VIENNACL_LINALG_OPENCL_KERNELS_VECTOR_ELEMENT_HPP
2 #define VIENNACL_LINALG_OPENCL_KERNELS_VECTOR_ELEMENT_HPP
24 template <
typename StringType>
26 std::string
const & funcname, std::string
const & op, std::string
const & op_name)
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");
38 template <
typename StringType>
46 template <
typename StringType>
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");
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");
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");
64 source.append(
" unsigned int op_type) \n");
65 source.append(
"{ \n");
66 if (numeric_string ==
"float" || numeric_string ==
"double")
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 ");
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");
104 static std::map<cl_context, bool> init_done;
108 source.reserve(8192);
110 viennacl::ocl::append_double_precision_pragma<TYPE>(ctx, source);
113 if (numeric_string ==
"float" || numeric_string ==
"double")
141 #ifdef VIENNACL_BUILD_INFO
142 std::cout <<
"Creating program " << prog_name << std::endl;
144 ctx.add_program(source, prog_name);
145 init_done[ctx.handle().get()] =
true;
static std::string program_name()
Definition: vector_element.hpp:94
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
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