Compute Library
 21.05
CLCompileContext.cpp
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2020-2021 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(), _is_wbsm_supported()
141 {
142 }
143 
144 CLCompileContext::CLCompileContext(cl::Context context, const cl::Device &device)
145  : _context(), _device(), _programs_map(), _built_programs_map(), _is_wbsm_supported()
146 {
147  _context = std::move(context);
148  _device = CLDevice(device);
149  _is_wbsm_supported = get_wbsm_support_info(device);
150 }
151 
152 Kernel CLCompileContext::create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source,
153  const std::string &kernel_path, const StringSet &build_options_set, bool is_binary) const
154 {
155  const std::string build_options = generate_build_options(build_options_set, kernel_path);
156  const std::string built_program_name = program_name + "_" + build_options;
157  auto built_program_it = _built_programs_map.find(built_program_name);
158  cl::Program cl_program;
159 
160  if(_built_programs_map.end() != built_program_it)
161  {
162  // If program has been built, retrieve to create kernel from it
163  cl_program = built_program_it->second;
164  }
165  else
166  {
167  Program program = load_program(program_name, program_source, is_binary);
168 
169  // Build program
170  cl_program = program.build(build_options);
171 
172  // Add built program to internal map
173  _built_programs_map.emplace(built_program_name, cl_program);
174  }
175 
176  // Create and return kernel
177  return Kernel(kernel_name, cl_program);
178 }
179 
180 const Program &CLCompileContext::load_program(const std::string &program_name, const std::string &program_source, bool is_binary) const
181 {
182  const auto program_it = _programs_map.find(program_name);
183 
184  if(program_it != _programs_map.end())
185  {
186  return program_it->second;
187  }
188 
189  Program program;
190 
191 #ifdef EMBEDDED_KERNELS
192  ARM_COMPUTE_UNUSED(is_binary);
193  program = Program(_context, program_name, program_source);
194 #else /* EMBEDDED_KERNELS */
195  if(is_binary)
196  {
197  program = Program(_context, _device.cl_device(), program_name, std::vector<unsigned char>(program_source.begin(), program_source.end()));
198  }
199  else
200  {
201  program = Program(_context, program_name, program_source);
202  }
203 #endif /* EMBEDDED_KERNELS */
204 
205  // Insert program to program map
206  const auto new_program = _programs_map.emplace(program_name, std::move(program));
207 
208  return new_program.first->second;
209 }
210 
211 void CLCompileContext::set_context(cl::Context context)
212 {
213  _context = std::move(context);
214  if(_context.get() != nullptr)
215  {
216  const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
217 
218  if(!cl_devices.empty())
219  {
220  _device = CLDevice(cl_devices[0]);
221  }
222  }
223 }
224 
225 std::string CLCompileContext::generate_build_options(const StringSet &build_options_set, const std::string &kernel_path) const
226 {
227  std::string concat_str;
228 
229 #if defined(ARM_COMPUTE_DEBUG_ENABLED)
230  // Enable debug properties in CL kernels
231  concat_str += " -DARM_COMPUTE_DEBUG_ENABLED";
232 #endif // defined(ARM_COMPUTE_DEBUG_ENABLED)
233 
234  GPUTarget gpu_arch = get_arch_from_target(_device.target());
235  concat_str += " -DGPU_ARCH=" + support::cpp11::to_string(
236  static_cast<std::underlying_type<GPUTarget>::type>(gpu_arch));
237 
238  if(_device.supported("cl_khr_fp16"))
239  {
240  concat_str += " -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
241  }
242 
243  if(_device.supported("cl_arm_integer_dot_product_int8"))
244  {
245  concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ENABLED=1 ";
246  }
247 
248  if(_device.supported("cl_arm_integer_dot_product_accumulate_int8"))
249  {
250  concat_str += " -DARM_COMPUTE_OPENCL_DOT8_ACC_ENABLED=1 ";
251  }
252 
253  if(_device.version() == CLVersion::CL20)
254  {
255  concat_str += " -cl-std=CL2.0 ";
256  }
257  else if(_device.supported("cl_arm_non_uniform_work_group_size"))
258  {
259  concat_str += " -cl-arm-non-uniform-work-group-size ";
260  }
261  else
262  {
263  ARM_COMPUTE_ERROR("Non uniform workgroup size is not supported!!");
264  }
265 
266  std::string build_options = stringify_set(build_options_set, kernel_path) + concat_str;
267 
268  return build_options;
269 }
270 
272 {
273  return _device.supported("cl_khr_fp16");
274 }
275 
276 std::string CLCompileContext::stringify_set(const StringSet &s, const std::string &kernel_path) const
277 {
278  std::string concat_set;
279 #ifndef EMBEDDED_KERNELS
280  concat_set += "-I" + kernel_path + " ";
281 #else /* EMBEDDED_KERNELS */
282  ARM_COMPUTE_UNUSED(kernel_path);
283 #endif /* EMBEDDED_KERNELS */
284 
285  // Concatenate set
286  for(const auto &el : s)
287  {
288  concat_set += " " + el;
289  }
290 
291  return concat_set;
292 }
293 
294 void CLCompileContext::add_built_program(const std::string &built_program_name, const cl::Program &program) const
295 {
296  _built_programs_map.emplace(built_program_name, program);
297 }
298 
300 {
301  _programs_map.clear();
302  _built_programs_map.clear();
303 }
304 
305 const std::map<std::string, cl::Program> &CLCompileContext::get_built_programs() const
306 {
307  return _built_programs_map;
308 }
309 
311 {
312  return _context;
313 }
314 
315 const cl::Device &CLCompileContext::get_device() const
316 {
317  return _device.cl_device();
318 }
319 
320 void CLCompileContext::set_device(cl::Device device)
321 {
322  _device = std::move(device);
323  _is_wbsm_supported = get_wbsm_support_info(device);
324 }
325 
327 {
328  GPUTarget _target = get_target_from_device(_device.cl_device());
329  cl::NDRange default_range;
330 
331  switch(_target)
332  {
333  case GPUTarget::MIDGARD:
334  case GPUTarget::T600:
335  case GPUTarget::T700:
336  case GPUTarget::T800:
337  default_range = cl::NDRange(128u, 1);
338  break;
339  default:
340  default_range = cl::NullRange;
341  }
342 
343  return default_range;
344 }
345 
347 {
348  return _device.supported("cl_khr_int64_base_atomics");
349 }
350 
352 {
353  return _is_wbsm_supported;
354 }
355 
356 size_t CLCompileContext::max_local_workgroup_size(const cl::Kernel &kernel) const
357 {
358  size_t result;
359 
360  size_t err = kernel.getWorkGroupInfo(_device.cl_device(), CL_KERNEL_WORK_GROUP_SIZE, &result);
361  ARM_COMPUTE_ERROR_ON_MSG(err != 0, "clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
362  ARM_COMPUTE_UNUSED(err);
363 
364  return result;
365 }
366 
368 {
369  return _device.device_version();
370 }
371 
373 {
374  return _device.compute_units();
375 }
376 } // namespace arm_compute
Definition: ICLTensor.h:33
std::string device_version() const
Returns the device version as a string.
Definition: CLDevice.h:130
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.
bool get_wbsm_support_info(const cl::Device &device)
Definition: CLHelpers.cpp:419
CLBuildOptions()
Default constructor.
GPUTarget get_arch_from_target(GPUTarget target)
Helper function to get the GPU arch.
Definition: GPUTarget.cpp:189
decltype(strategy::transforms) typedef type
std::string get_device_version() const
Return the device version.
Copyright (c) 2017-2021 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:121
#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:94
#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:43
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:141
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.
const char * name
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:103
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.
Wrapper to configure the Khronos OpenCL C++ header.
const cl::Device & cl_device() const
Returns the underlying cl device object.
Definition: CLDevice.h:112
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.