blob: 0afb7e5e0e0473cb1aa0531c9622696962dcd68a [file] [log] [blame]
Michalis Spyrou11d49182020-03-26 10:31:32 +00001/*
Michele Di Giorgiod9eaf612020-07-08 11:12:57 +01002 * Copyright (c) 2020 Arm Limited.
Michalis Spyrou11d49182020-03-26 10:31:32 +00003 *
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 */
24#include "arm_compute/core/CL/CLCompileContext.h"
25#include "arm_compute/core/CL/OpenCL.h"
26
27#include "arm_compute/core/CL/CLHelpers.h"
28#include "arm_compute/core/Error.h"
29#include "arm_compute/core/Utils.h"
30#include "support/StringSupport.h"
31
32namespace arm_compute
33{
34CLBuildOptions::CLBuildOptions()
35 : _build_opts()
36{
37}
38
39void CLBuildOptions::add_option(std::string option)
40{
41 _build_opts.emplace(std::move(option));
42}
43
44void CLBuildOptions::add_option_if(bool cond, std::string option)
45{
46 if(cond)
47 {
48 add_option(std::move(option));
49 }
50}
51
52void 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
57void CLBuildOptions::add_options(const StringSet &options)
58{
59 _build_opts.insert(options.begin(), options.end());
60}
61
62void CLBuildOptions::add_options_if(bool cond, const StringSet &options)
63{
64 if(cond)
65 {
66 add_options(options);
67 }
68}
69
70const CLBuildOptions::StringSet &CLBuildOptions::options() const
71{
72 return _build_opts;
73}
74
75Program::Program()
76 : _context(), _device(), _is_binary(false), _name(), _source(), _binary()
77{
78}
79
80Program::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
85Program::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
90Program::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
102bool 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
122cl::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
129Kernel::Kernel()
130 : _name(), _kernel()
131{
132}
133
134Kernel::Kernel(std::string name, const cl::Program &program)
135 : _name(std::move(name)),
136 _kernel(cl::Kernel(program, _name.c_str()))
137{
138}
139CLCompileContext::CLCompileContext()
140 : _context(), _device(), _programs_map(), _built_programs_map()
141{
142}
143
144CLCompileContext::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
151Kernel 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
Michalis Spyrou12910f22020-04-14 13:54:45 +0100172 _built_programs_map.emplace(built_program_name, cl_program);
Michalis Spyrou11d49182020-03-26 10:31:32 +0000173 }
174
175 // Create and return kernel
176 return Kernel(kernel_name, cl_program);
177}
178
179const 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
210void 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
224std::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
270bool CLCompileContext::fp16_supported() const
271{
272 return _device.supported("cl_khr_fp16");
273}
274
275std::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
293void 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
298void CLCompileContext::clear_programs_cache()
299{
300 _programs_map.clear();
301 _built_programs_map.clear();
302}
303
304const std::map<std::string, cl::Program> &CLCompileContext::get_built_programs() const
305{
306 return _built_programs_map;
307}
308
309cl::Context &CLCompileContext::context()
310{
311 return _context;
312}
313
314const cl::Device &CLCompileContext::get_device() const
315{
316 return _device.cl_device();
317}
318
319void CLCompileContext::set_device(cl::Device device)
320{
321 _device = std::move(device);
322}
323
324cl::NDRange CLCompileContext::default_ndrange() const
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
344bool CLCompileContext::int64_base_atomics_supported() const
345{
346 return _device.supported("cl_khr_int64_base_atomics");
347}
348
349size_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
360std::string CLCompileContext::get_device_version() const
361{
362 return _device.device_version();
363}
364
365cl_uint CLCompileContext::get_num_compute_units() const
366{
367 return _device.compute_units();
368}
369} // namespace arm_compute