Compute Library
 20.08
CLCompileContext.cpp
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2020 Arm Limited.
3  *
4  * SPDX-License-Identifier: MIT
5  *
6  * Permission is hereby granted, free of charge, to any person obtaining a copy
7  * of this software and associated documentation files (the "Software"), to
8  * deal in the Software without restriction, including without limitation the
9  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10  * sell copies of the Software, and to permit persons to whom the Software is
11  * furnished to do so, subject to the following conditions:
12  *
13  * The above copyright notice and this permission notice shall be included in all
14  * copies or substantial portions of the Software.
15  *
16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22  * SOFTWARE.
23  */
26 
28 #include "arm_compute/core/Error.h"
29 #include "arm_compute/core/Utils.h"
30 #include "support/StringSupport.h"
31 
32 namespace arm_compute
33 {
35  : _build_opts()
36 {
37 }
38 
39 void CLBuildOptions::add_option(std::string option)
40 {
41  _build_opts.emplace(std::move(option));
42 }
43 
44 void CLBuildOptions::add_option_if(bool cond, std::string option)
45 {
46  if(cond)
47  {
48  add_option(std::move(option));
49  }
50 }
51 
52 void CLBuildOptions::add_option_if_else(bool cond, std::string option_true, std::string option_false)
53 {
54  (cond) ? add_option(std::move(option_true)) : add_option(std::move(option_false));
55 }
56 
57 void CLBuildOptions::add_options(const StringSet &options)
58 {
59  _build_opts.insert(options.begin(), options.end());
60 }
61 
62 void CLBuildOptions::add_options_if(bool cond, const StringSet &options)
63 {
64  if(cond)
65  {
67  }
68 }
69 
70 const CLBuildOptions::StringSet &CLBuildOptions::options() const
71 {
72  return _build_opts;
73 }
74 
76  : _context(), _device(), _is_binary(false), _name(), _source(), _binary()
77 {
78 }
79 
80 Program::Program(cl::Context context, std::string name, std::string source)
81  : _context(std::move(context)), _device(), _is_binary(false), _name(std::move(name)), _source(std::move(source)), _binary()
82 {
83 }
84 
85 Program::Program(cl::Context context, cl::Device device, std::string name, std::vector<unsigned char> binary)
86  : _context(std::move(context)), _device(std::move(device)), _is_binary(true), _name(std::move(name)), _source(), _binary(std::move(binary))
87 {
88 }
89 
90 Program::operator cl::Program() const
91 {
92  if(_is_binary)
93  {
94  return cl::Program(_context, { _device }, { _binary });
95  }
96  else
97  {
98  return cl::Program(_context, _source, false);
99  }
100 }
101 
102 bool Program::build(const cl::Program &program, const std::string &build_options)
103 {
104  try
105  {
106  return program.build(build_options.c_str()) == CL_SUCCESS;
107  }
108  catch(const cl::Error &e)
109  {
110  cl_int err = CL_SUCCESS;
111  const auto build_info = program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(&err);
112 
113  for(auto &pair : build_info)
114  {
115  std::cerr << pair.second << std::endl;
116  }
117 
118  return false;
119  }
120 }
121 
122 cl::Program Program::build(const std::string &build_options) const
123 {
124  cl::Program cl_program = static_cast<cl::Program>(*this);
125  build(cl_program, build_options);
126  return cl_program;
127 }
128 
130  : _name(), _kernel()
131 {
132 }
133 
134 Kernel::Kernel(std::string name, const cl::Program &program)
135  : _name(std::move(name)),
136  _kernel(cl::Kernel(program, _name.c_str()))
137 {
138 }
140  : _context(), _device(), _programs_map(), _built_programs_map()
141 {
142 }
143 
144 CLCompileContext::CLCompileContext(cl::Context context, const cl::Device &device)
145  : _context(), _device(), _programs_map(), _built_programs_map()
146 {
147  _context = std::move(context);
148  _device = CLDevice(device);
149 }
150 
151 Kernel CLCompileContext::create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source,
152  const std::string &kernel_path, const StringSet &build_options_set, bool is_binary) const
153 {
154  const std::string build_options = generate_build_options(build_options_set, kernel_path);
155  const std::string built_program_name = program_name + "_" + build_options;
156  auto built_program_it = _built_programs_map.find(built_program_name);
157  cl::Program cl_program;
158 
159  if(_built_programs_map.end() != built_program_it)
160  {
161  // If program has been built, retrieve to create kernel from it
162  cl_program = built_program_it->second;
163  }
164  else
165  {
166  Program program = load_program(program_name, program_source, is_binary);
167 
168  // Build program
169  cl_program = program.build(build_options);
170 
171  // Add built program to internal map
172  _built_programs_map.emplace(built_program_name, cl_program);
173  }
174 
175  // Create and return kernel
176  return Kernel(kernel_name, cl_program);
177 }
178 
179 const Program &CLCompileContext::load_program(const std::string &program_name, const std::string &program_source, bool is_binary) const
180 {
181  const auto program_it = _programs_map.find(program_name);
182 
183  if(program_it != _programs_map.end())
184  {
185  return program_it->second;
186  }
187 
188  Program program;
189 
190 #ifdef EMBEDDED_KERNELS
191  ARM_COMPUTE_UNUSED(is_binary);
192  program = Program(_context, program_name, program_source);
193 #else /* EMBEDDED_KERNELS */
194  if(is_binary)
195  {
196  program = Program(_context, _device.cl_device(), program_name, std::vector<unsigned char>(program_source.begin(), program_source.end()));
197  }
198  else
199  {
200  program = Program(_context, program_name, program_source);
201  }
202 #endif /* EMBEDDED_KERNELS */
203 
204  // Insert program to program map
205  const auto new_program = _programs_map.emplace(program_name, std::move(program));
206 
207  return new_program.first->second;
208 }
209 
210 void CLCompileContext::set_context(cl::Context context)
211 {
212  _context = std::move(context);
213  if(_context.get() != nullptr)
214  {
215  const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
216 
217  if(!cl_devices.empty())
218  {
219  _device = CLDevice(cl_devices[0]);
220  }
221  }
222 }
223 
224 std::string CLCompileContext::generate_build_options(const StringSet &build_options_set, const std::string &kernel_path) const
225 {
226  std::string concat_str;
227 
228 #if defined(ARM_COMPUTE_DEBUG_ENABLED)
229  // Enable debug properties in CL kernels
230  concat_str += " -DARM_COMPUTE_DEBUG_ENABLED";
231 #endif // defined(ARM_COMPUTE_DEBUG_ENABLED)
232 
233  GPUTarget gpu_arch = get_arch_from_target(_device.target());
234  concat_str += " -DGPU_ARCH=" + support::cpp11::to_string(
235  static_cast<std::underlying_type<GPUTarget>::type>(gpu_arch));
236 
237  if(_device.supported("cl_khr_fp16"))
238  {
239  concat_str += " -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
240  }
241 
242  if(_device.supported("cl_arm_integer_dot_product_int8"))
243  {
244  concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 ";
245  }
246 
247  if(_device.supported("cl_arm_integer_dot_product_accumulate_int8"))
248  {
249  concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED=1 ";
250  }
251 
252  if(_device.version() == CLVersion::CL20)
253  {
254  concat_str += " -cl-std=CL2.0 ";
255  }
256  else if(_device.supported("cl_arm_non_uniform_work_group_size"))
257  {
258  concat_str += " -cl-arm-non-uniform-work-group-size ";
259  }
260  else
261  {
262  ARM_COMPUTE_ERROR("Non uniform workgroup size is not supported!!");
263  }
264 
265  std::string build_options = stringify_set(build_options_set, kernel_path) + concat_str;
266 
267  return build_options;
268 }
269 
271 {
272  return _device.supported("cl_khr_fp16");
273 }
274 
275 std::string CLCompileContext::stringify_set(const StringSet &s, const std::string &kernel_path) const
276 {
277  std::string concat_set;
278 #ifndef EMBEDDED_KERNELS
279  concat_set += "-I" + kernel_path + " ";
280 #else /* EMBEDDED_KERNELS */
281  ARM_COMPUTE_UNUSED(kernel_path);
282 #endif /* EMBEDDED_KERNELS */
283 
284  // Concatenate set
285  for(const auto &el : s)
286  {
287  concat_set += " " + el;
288  }
289 
290  return concat_set;
291 }
292 
293 void CLCompileContext::add_built_program(const std::string &built_program_name, const cl::Program &program) const
294 {
295  _built_programs_map.emplace(built_program_name, program);
296 }
297 
299 {
300  _programs_map.clear();
301  _built_programs_map.clear();
302 }
303 
304 const std::map<std::string, cl::Program> &CLCompileContext::get_built_programs() const
305 {
306  return _built_programs_map;
307 }
308 
310 {
311  return _context;
312 }
313 
314 const cl::Device &CLCompileContext::get_device() const
315 {
316  return _device.cl_device();
317 }
318 
319 void CLCompileContext::set_device(cl::Device device)
320 {
321  _device = std::move(device);
322 }
323 
325 {
326  GPUTarget _target = get_target_from_device(_device.cl_device());
327  cl::NDRange default_range;
328 
329  switch(_target)
330  {
331  case GPUTarget::MIDGARD:
332  case GPUTarget::T600:
333  case GPUTarget::T700:
334  case GPUTarget::T800:
335  default_range = cl::NDRange(128u, 1);
336  break;
337  default:
338  default_range = cl::NullRange;
339  }
340 
341  return default_range;
342 }
343 
345 {
346  return _device.supported("cl_khr_int64_base_atomics");
347 }
348 
349 size_t CLCompileContext::max_local_workgroup_size(const cl::Kernel &kernel) const
350 {
351  size_t result;
352 
353  size_t err = kernel.getWorkGroupInfo(_device.cl_device(), CL_KERNEL_WORK_GROUP_SIZE, &result);
354  ARM_COMPUTE_ERROR_ON_MSG(err != 0, "clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
355  ARM_COMPUTE_UNUSED(err);
356 
357  return result;
358 }
359 
361 {
362  return _device.device_version();
363 }
364 
366 {
367  return _device.compute_units();
368 }
369 } // namespace arm_compute
std::string device_version() const
Returns the device version as a string.
Definition: CLDevice.h:129
void set_device(cl::Device device)
Sets the CL device for which the programs are created.
const cl::Device & get_device() const
Gets the CL device for which the programs are created.
cl::Context & context()
Accessor for the associated CL context.
Kernel()
Default Constructor.
const StringSet & options() const
Gets the current options list set.
#define ARM_COMPUTE_ERROR(msg)
Print the given message then throw an std::runtime_error.
Definition: Error.h:352
std::string to_string(T &&value)
Convert integer and float values to string.
CLBuildOptions()
Default constructor.
GPUTarget get_arch_from_target(GPUTarget target)
Helper function to get the GPU arch.
Definition: GPUTarget.cpp:189
std::string get_device_version() const
Return the device version.
Copyright (c) 2017-2020 Arm Limited.
void add_option(std::string option)
Adds option to the existing build option list.
void add_options(const StringSet &options)
Appends given build options to the current's objects options.
CLVersion version() const
Returns the device's CL version.
Definition: CLDevice.h:120
#define ARM_COMPUTE_UNUSED(...)
To avoid unused variables warnings.
Definition: Error.h:152
cl_uint get_num_compute_units() const
Return the maximum number of compute units in the device.
const GPUTarget & target() const
Returns the GPU target of the cl device.
Definition: CLDevice.h:93
#define ARM_COMPUTE_ERROR_ON_MSG(cond, msg)
Definition: Error.h:456
std::set< std::string > build_options
std::string kernel_name
OpenCL device type class.
Definition: CLDevice.h:42
static bool build(const cl::Program &program, const std::string &build_options="")
Build the given CL program.
bool supported(const std::string &extension) const override
Check if extensions on a device are supported.
Definition: CLDevice.h:140
cl::NDRange default_ndrange() const
Return the default NDRange for the device.
void add_option_if(bool cond, std::string option)
Adds option if a given condition is true;.
GPUTarget get_target_from_device(const cl::Device &device)
Helper function to get the GPU target from CL device.
Definition: CLHelpers.cpp:221
Kernel create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source, const std::string &kernel_path, const StringSet &build_options_set, bool is_binary) const
Creates an OpenCL kernel.
void clear_programs_cache()
Clear the library's cache of binary programs.
Program()
Default constructor.
bool int64_base_atomics_supported() const
Returns true if int64_base_atomics extension is supported by the CL device.
bool fp16_supported() const
Returns true if FP16 is supported by the CL device.
size_t compute_units() const
Returns the number of compute units available.
Definition: CLDevice.h:102
const std::map< std::string, cl::Program > & get_built_programs() const
Access the cache of built OpenCL programs.
GPUTarget
Available GPU Targets.
Definition: GPUTarget.h:34
size_t max_local_workgroup_size(const cl::Kernel &kernel) const
Find the maximum number of local work items in a workgroup can be supported for the kernel.
const cl::Device & cl_device() const
Returns the underlying cl device object.
Definition: CLDevice.h:111
void add_built_program(const std::string &built_program_name, const cl::Program &program) const
Add a new built program to the cache.
void set_context(cl::Context context)
Sets the CL context used to create programs.
void add_options_if(bool cond, const StringSet &options)
Appends given build options to the current's objects options if a given condition is true.
void add_option_if_else(bool cond, std::string option_true, std::string option_false)
Adds first option if condition is true else the second one.