SHOGUN  4.0.0
opencl_util.h
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2014, Shogun Toolbox Foundation
3  * All rights reserved.
4  *
5  * Redistribution and use in source and binary forms, with or without
6  * modification, are permitted provided that the following conditions are met:
7 
8  * 1. Redistributions of source code must retain the above copyright notice,
9  * this list of conditions and the following disclaimer.
10  *
11  * 2. Redistributions in binary form must reproduce the above copyright notice,
12  * this list of conditions and the following disclaimer in the documentation
13  * and/or other materials provided with the distribution.
14  *
15  * 3. Neither the name of the copyright holder nor the names of its
16  * contributors may be used to endorse or promote products derived from this
17  * software without specific prior written permission.
18 
19  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
20  * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
21  * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
22  * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
23  * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
24  * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
25  * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
26  * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
27  * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
28  * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
29  * POSSIBILITY OF SUCH DAMAGE.
30  *
31  * Written (W) 2014 Khaled Nasr
32  */
33 
35 #ifndef __OPENCL_UTIL_H__
36 #define __OPENCL_UTIL_H__
37 
38 #include <shogun/lib/config.h>
39 #ifdef HAVE_VIENNACL
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>
45 
47 
48 #include <string>
49 
50 namespace shogun
51 {
52 
53 namespace linalg
54 {
55 
56 namespace implementation
57 {
58 
59 namespace ocl
60 {
61 
63 template <class T>
64 std::string get_type_string()
65 {
66  return viennacl::ocl::type_to_string<T>::apply();
67 }
68 
75 template <class T>
76 std::string generate_kernel_preamble(std::string kernel_name)
77 {
78  std::string type_string = get_type_string<T>();
79 
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");
86 
87  return source;
88 }
89 
91 inline bool kernel_exists(std::string kernel_name)
92 {
93  return viennacl::ocl::current_context().has_program(kernel_name);
94 }
95 
97 inline viennacl::ocl::kernel& get_kernel(std::string kernel_name)
98 {
99  return viennacl::ocl::current_context().get_program(kernel_name).get_kernel(kernel_name);
100 }
101 
103 inline viennacl::ocl::kernel& compile_kernel(std::string kernel_name, std::string source)
104 {
105  viennacl::ocl::program & prog =
106  viennacl::ocl::current_context().add_program(source, kernel_name);
107 
108  return prog.get_kernel(kernel_name);
109 }
110 
112 inline uint32_t align_to_multiple_1d(uint32_t n)
113 {
114  return viennacl::tools::align_to_multiple<uint32_t>(n, OCL_WORK_GROUP_SIZE_1D);
115 }
116 
118 inline uint32_t align_to_multiple_2d(uint32_t n)
119 {
120  return viennacl::tools::align_to_multiple<uint32_t>(n, OCL_WORK_GROUP_SIZE_2D);
121 }
122 
135 template <class T>
136 viennacl::ocl::kernel& generate_single_arg_elementwise_kernel(
137  std::string kernel_name, std::string operation)
138 {
139  if (ocl::kernel_exists(kernel_name))
140  return ocl::get_kernel(kernel_name);
141 
142  std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
143 
144  source.append("inline DATATYPE operation(DATATYPE element)\n{\n");
145  source.append(operation);
146  source.append("\n}\n");
147 
148  source.append(
149  R"(
150  __kernel void KERNEL_NAME(
151  __global DATATYPE* vec, int size, int vec_offset,
152  __global DATATYPE* result, int result_offset)
153  {
154  int i = get_global_id(0);
155 
156  if (i<size)
157  result[i+result_offset] = operation(vec[i+vec_offset]);
158  }
159  )"
160  );
161 
162  viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
163 
164  kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
165 
166  return kernel;
167 }
168 
182 template <class T>
183 viennacl::ocl::kernel& generate_two_arg_elementwise_kernel(
184  std::string kernel_name, std::string operation)
185 {
186  if (ocl::kernel_exists(kernel_name))
187  return ocl::get_kernel(kernel_name);
188 
189  std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
190 
191  source.append("inline DATATYPE operation(DATATYPE element1, DATATYPE element2)\n{\n");
192  source.append(operation);
193  source.append("\n}\n");
194 
195  source.append(
196  R"(
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)
201  {
202  int i = get_global_id(0);
203 
204  if (i<size)
205  result[i+result_offset] =
206  operation(vec1[i+vec1_offset], vec2[i+vec2_offset]);
207  }
208  )"
209  );
210 
211  viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
212 
213  kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
214 
215  return kernel;
216 }
217 
218 }
219 }
220 }
221 }
222 
223 #endif // HAVE_VIENNACL
224 
225 #endif // __OPENCL_UTIL_H__
all of classes and functions are contained in the shogun namespace
Definition: class_list.h:18

SHOGUN Machine Learning Toolbox - Documentation