35 #ifndef __OPENCL_UTIL_H__
36 #define __OPENCL_UTIL_H__
40 #include <viennacl/ocl/backend.hpp>
41 #include <viennacl/ocl/kernel.hpp>
42 #include <viennacl/ocl/program.hpp>
43 #include <viennacl/ocl/utils.hpp>
44 #include <viennacl/tools/tools.hpp>
56 namespace implementation
64 std::string get_type_string()
66 return viennacl::ocl::type_to_string<T>::apply();
76 std::string generate_kernel_preamble(std::string kernel_name)
78 std::string type_string = get_type_string<T>();
80 std::string source =
"";
81 viennacl::ocl::append_double_precision_pragma<T>(viennacl::ocl::current_context(), source);
82 source.append(
"#define DATATYPE " + type_string +
"\n");
83 source.append(
"#define KERNEL_NAME " + kernel_name +
"\n");
84 source.append(
"#define WORK_GROUP_SIZE_1D " + std::to_string(OCL_WORK_GROUP_SIZE_1D) +
"\n");
85 source.append(
"#define WORK_GROUP_SIZE_2D " + std::to_string(OCL_WORK_GROUP_SIZE_2D) +
"\n");
91 inline bool kernel_exists(std::string kernel_name)
93 return viennacl::ocl::current_context().has_program(kernel_name);
97 inline viennacl::ocl::kernel& get_kernel(std::string kernel_name)
99 return viennacl::ocl::current_context().get_program(kernel_name).get_kernel(kernel_name);
103 inline viennacl::ocl::kernel& compile_kernel(std::string kernel_name, std::string source)
105 viennacl::ocl::program & prog =
106 viennacl::ocl::current_context().add_program(source, kernel_name);
108 return prog.get_kernel(kernel_name);
112 inline uint32_t align_to_multiple_1d(uint32_t n)
114 return viennacl::tools::align_to_multiple<uint32_t>(n, OCL_WORK_GROUP_SIZE_1D);
118 inline uint32_t align_to_multiple_2d(uint32_t n)
120 return viennacl::tools::align_to_multiple<uint32_t>(n, OCL_WORK_GROUP_SIZE_2D);
136 viennacl::ocl::kernel& generate_single_arg_elementwise_kernel(
137 std::string kernel_name, std::string operation)
139 if (ocl::kernel_exists(kernel_name))
140 return ocl::get_kernel(kernel_name);
142 std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
144 source.append(
"inline DATATYPE operation(DATATYPE element)\n{\n");
145 source.append(operation);
146 source.append(
"\n}\n");
150 __kernel void KERNEL_NAME(
151 __global DATATYPE* vec, int size, int vec_offset,
152 __global DATATYPE* result, int result_offset)
154 int i = get_global_id(0);
157 result[i+result_offset] = operation(vec[i+vec_offset]);
162 viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
164 kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
183 viennacl::ocl::kernel& generate_two_arg_elementwise_kernel(
184 std::string kernel_name, std::string operation)
186 if (ocl::kernel_exists(kernel_name))
187 return ocl::get_kernel(kernel_name);
189 std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
191 source.append(
"inline DATATYPE operation(DATATYPE element1, DATATYPE element2)\n{\n");
192 source.append(operation);
193 source.append(
"\n}\n");
197 __kernel void KERNEL_NAME(
198 __global DATATYPE* vec1, int size, int vec1_offset,
199 __global DATATYPE* vec2, int vec2_offset,
200 __global DATATYPE* result, int result_offset)
202 int i = get_global_id(0);
205 result[i+result_offset] =
206 operation(vec1[i+vec1_offset], vec2[i+vec2_offset]);
211 viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
213 kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
223 #endif // HAVE_VIENNACL
225 #endif // __OPENCL_UTIL_H__
all of classes and functions are contained in the shogun namespace