SHOGUN  4.1.0
 全部  命名空间 文件 函数 变量 类型定义 枚举 枚举值 友元 宏定义  
opencl_util.h
浏览该文件的文档.
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 #if defined(HAVE_CXX0X) || defined(HAVE_CXX11)
51 #include <initializer_list>
53 #endif // defined(HAVE_CXX0X) || defined(HAVE_CXX11)
54 
55 namespace shogun
56 {
57 
58 namespace linalg
59 {
60 
61 namespace implementation
62 {
63 
64 namespace ocl
65 {
66 
68 template <class T>
69 std::string get_type_string()
70 {
71  return viennacl::ocl::type_to_string<T>::apply();
72 }
73 
80 template <class T>
81 std::string generate_kernel_preamble(std::string kernel_name)
82 {
83  std::string type_string = get_type_string<T>();
84 
85  std::string source = "";
86  viennacl::ocl::append_double_precision_pragma<T>(viennacl::ocl::current_context(), source);
87  source.append("#define DATATYPE " + type_string + "\n");
88  source.append("#define KERNEL_NAME " + kernel_name + "\n");
89  source.append("#define WORK_GROUP_SIZE_1D " + std::to_string(OCL_WORK_GROUP_SIZE_1D) + "\n");
90  source.append("#define WORK_GROUP_SIZE_2D " + std::to_string(OCL_WORK_GROUP_SIZE_2D) + "\n");
91 
92  return source;
93 }
94 
96 inline bool kernel_exists(std::string kernel_name)
97 {
98  return viennacl::ocl::current_context().has_program(kernel_name);
99 }
100 
102 inline viennacl::ocl::kernel& get_kernel(std::string kernel_name)
103 {
104  return viennacl::ocl::current_context().get_program(kernel_name).get_kernel(kernel_name);
105 }
106 
108 inline viennacl::ocl::kernel& compile_kernel(std::string kernel_name, std::string source)
109 {
110  viennacl::ocl::program & prog =
111  viennacl::ocl::current_context().add_program(source, kernel_name);
112 
113  return prog.get_kernel(kernel_name);
114 }
115 
117 inline uint32_t align_to_multiple_1d(uint32_t n)
118 {
119  return viennacl::tools::align_to_multiple<uint32_t>(n, OCL_WORK_GROUP_SIZE_1D);
120 }
121 
123 inline uint32_t align_to_multiple_2d(uint32_t n)
124 {
125  return viennacl::tools::align_to_multiple<uint32_t>(n, OCL_WORK_GROUP_SIZE_2D);
126 }
127 
140 template <class T>
141 viennacl::ocl::kernel& generate_single_arg_elementwise_kernel(
142  std::string kernel_name, std::string operation)
143 {
144  if (ocl::kernel_exists(kernel_name))
145  return ocl::get_kernel(kernel_name);
146 
147  std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
148 
149  source.append("inline DATATYPE operation(DATATYPE element)\n{\n");
150  source.append(operation);
151  source.append("\n}\n");
152 
153  source.append(
154  R"(
155  __kernel void KERNEL_NAME(
156  __global DATATYPE* vec, int size, int vec_offset,
157  __global DATATYPE* result, int result_offset)
158  {
159  int i = get_global_id(0);
160 
161  if (i<size)
162  result[i+result_offset] = operation(vec[i+vec_offset]);
163  }
164  )"
165  );
166 
167  viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
168 
169  kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
170 
171  return kernel;
172 }
173 
187 template <class T>
188 viennacl::ocl::kernel& generate_two_arg_elementwise_kernel(
189  std::string kernel_name, std::string operation)
190 {
191  if (ocl::kernel_exists(kernel_name))
192  return ocl::get_kernel(kernel_name);
193 
194  std::string source = ocl::generate_kernel_preamble<T>(kernel_name);
195 
196  source.append("inline DATATYPE operation(DATATYPE element1, DATATYPE element2)\n{\n");
197  source.append(operation);
198  source.append("\n}\n");
199 
200  source.append(
201  R"(
202  __kernel void KERNEL_NAME(
203  __global DATATYPE* vec1, int size, int vec1_offset,
204  __global DATATYPE* vec2, int vec2_offset,
205  __global DATATYPE* result, int result_offset)
206  {
207  int i = get_global_id(0);
208 
209  if (i<size)
210  result[i+result_offset] =
211  operation(vec1[i+vec1_offset], vec2[i+vec2_offset]);
212  }
213  )"
214  );
215 
216  viennacl::ocl::kernel& kernel = ocl::compile_kernel(kernel_name, source);
217 
218  kernel.local_work_size(0, OCL_WORK_GROUP_SIZE_1D);
219 
220  return kernel;
221 }
222 
234 inline std::string replace_all(std::string str, const std::string& from, const std::string& to)
235 {
236  size_t start_pos=0;
237  while ((start_pos=str.find(from, start_pos))!=std::string::npos)
238  {
239  str.replace(start_pos, from.length(), to);
240  start_pos+=to.length(); // Handles case where 'to' is a substring of 'from'
241  }
242  return str;
243 }
244 
245 #if defined(HAVE_CXX0X) || defined(HAVE_CXX11)
246 
262 inline std::string format(const char* str, std::initializer_list<shogun::linalg::ocl::Parameter> params)
263 {
264  std::string fmt(str);
265  for (auto i=params.begin(); i!=params.end(); ++i)
266  fmt=replace_all(fmt, "{"+i->m_name+"}", *i);
267  return fmt.append("\n");
268 }
269 #endif // defined(HAVE_CXX0X) || defined(HAVE_CXX11)
270 
271 } // ocl
272 
273 } // implementation
274 
275 } // linalg
276 
277 } // shogun
278 
279 #endif // HAVE_VIENNACL
280 
281 #endif // __OPENCL_UTIL_H__
all of classes and functions are contained in the shogun namespace
Definition: class_list.h:18

SHOGUN 机器学习工具包 - 项目文档