2 * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
4 * Licensed under the Apache License, Version 2.0 (the "License");
5 * you may not use this file except in compliance with the License.
6 * You may obtain a copy of the License at
8 * http://www.apache.org/licenses/LICENSE-2.0
10 * Unless required by applicable law or agreed to in writing, software
11 * distributed under the License is distributed on an "AS IS" BASIS,
12 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 * See the License for the specific language governing permissions and
14 * limitations under the License.
18 * Copyright (c) 2016-2018 ARM Limited.
20 * SPDX-License-Identifier: MIT
22 * Permission is hereby granted, free of charge, to any person obtaining a copy
23 * of this software and associated documentation files (the "Software"), to
24 * deal in the Software without restriction, including without limitation the
25 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
26 * sell copies of the Software, and to permit persons to whom the Software is
27 * furnished to do so, subject to the following conditions:
29 * The above copyright notice and this permission notice shall be included in all
30 * copies or substantial portions of the Software.
32 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
33 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
34 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
35 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
36 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
37 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
42 * @file CLKernelLibraryEx.h
43 * @ingroup COM_AI_RUNTIME
44 * @brief This file is a cloned version of CLKernelLibrary.h in ACL. This file defines
45 * an interface for CLKernelLibrary.cpp which adds more OpenCL kernels on top of ACL.
48 #ifndef __ARM_COMPUTE_CLKERNELLIBRARY_EX_H__
49 #define __ARM_COMPUTE_CLKERNELLIBRARY_EX_H__
51 #include "arm_compute/core/CL/OpenCL.h"
62 * @brief Class to build OpenCL kernels added from nnfw
64 class CLKernelLibraryEx
66 using StringSet = std::set<std::string>;
70 * @brief Construct a new CLKernelLibraryEx object
76 * @brief Prevent instances of this class from being copied.
78 CLKernelLibraryEx(const CLKernelLibraryEx &) = delete;
81 * @brief Prevent instances of this class from being copied.
83 const CLKernelLibraryEx &operator=(const CLKernelLibraryEx &) = delete;
86 * @brief Get the KernelLibrary singleton.
87 * @return The KernelLibrary instance
89 static CLKernelLibraryEx &get();
92 * @brief Initialise the kernel library.
93 * @param[in] kernel_path Path of the directory from which kernel sources are loaded.
94 * @param[in] context CL context used to create programs.
95 * @param[in] device CL device for which the programs are created.
98 void init(std::string kernel_path, cl::Context context, cl::Device device)
100 _kernel_path = std::move(kernel_path);
101 _context = std::move(context);
102 _device = std::move(device);
106 * @brief Set the path that the kernels reside in.
107 * @param[in] kernel_path Path of the directory from which kernel sources are loaded.
110 void set_kernel_path(const std::string &kernel_path) { _kernel_path = kernel_path; };
113 * @brief Get the path that the kernels reside in.
114 * @return the path of kernel files
116 std::string get_kernel_path() { return _kernel_path; };
119 * @brief Get the source of the selected program.
120 * @param[in] program_name Program name.
121 * @return Source of the selected program.
123 std::string get_program_source(const std::string &program_name);
126 * @brief Set the CL context used to create programs.
127 * @note Setting the context also resets the device to the
128 * first one available in the new context.
129 * @param[in] context A CL context.
132 void set_context(cl::Context context)
134 _context = std::move(context);
135 if (_context.get() == nullptr)
137 _device = cl::Device();
141 const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
143 if (cl_devices.empty())
145 _device = cl::Device();
149 _device = cl_devices[0];
155 * @brief Return associated CL context.
156 * @return A CL context.
158 cl::Context &context() { return _context; }
161 * @brief Set the CL device for which the programs are created.
162 * @param[in] device A CL device.
165 void set_device(cl::Device device) { _device = std::move(device); }
168 * @brief Gets the CL device for which the programs are created.
169 * @return A CL device.
171 cl::Device &get_device() { return _device; }
174 * @brief Return the device version
175 * @return The content of CL_DEVICE_VERSION
177 std::string get_device_version();
180 * @brief Create a kernel from the kernel library.
181 * @param[in] kernel_name Kernel name.
182 * @param[in] build_options_set Kernel build options as a set.
183 * @return The created kernel.
185 Kernel create_kernel(const std::string &kernel_name,
186 const StringSet &build_options_set = {}) const;
189 * @brief Find the maximum number of local work items in a workgroup can be supported for the
191 * @param[in] kernel kernel object
194 size_t max_local_workgroup_size(const cl::Kernel &kernel) const;
196 * @brief Return the default NDRange for the device.
197 * @return default NDRangeof the device
199 cl::NDRange default_ndrange() const;
202 * @brief Clear the library's cache of binary programs
205 void clear_programs_cache()
207 _programs_map.clear();
208 _built_programs_map.clear();
212 * @brief Access the cache of built OpenCL programs
213 * @return program map data structure of which key is name of kernel and value is
214 * kerel source name. (*.cl)
216 const std::map<std::string, cl::Program> &get_built_programs() const
218 return _built_programs_map;
222 * @brief Add a new built program to the cache
223 * @param[in] built_program_name Name of the program
224 * @param[in] program Built program to add to the cache
227 void add_built_program(const std::string &built_program_name, cl::Program program);
230 * @brief Returns true if FP16 is supported by the CL device
231 * @return true if the CL device supports FP16
233 bool fp16_supported() const;
236 * @brief Returns true if int64_base_atomics extension is supported by the CL device
237 * @return true if the CL device supports int64_base_atomics extension
239 bool int64_base_atomics_supported() const;
243 * @brief Load program and its dependencies.
244 * @param[in] program_name Name of the program to load.
246 const Program &load_program(const std::string &program_name) const;
248 * @brief Concatenates contents of a set into a single string.
249 * @param[in] s Input set to concatenate.
250 * @return Concatenated string.
252 std::string stringify_set(const StringSet &s) const;
254 cl::Context _context; /**< Underlying CL context. */
255 cl::Device _device; /**< Underlying CL device. */
256 std::string _kernel_path; /**< Path to the kernels folder. */
257 mutable std::map<std::string, const Program>
258 _programs_map; /**< Map with all already loaded program data. */
259 mutable std::map<std::string, cl::Program>
260 _built_programs_map; /**< Map with all already built program data. */
261 static const std::map<std::string, std::string>
262 _kernel_program_map; /**< Map that associates kernel names with programs. */
263 static const std::map<std::string, std::string>
264 _program_source_map; /**< Contains sources for all programs.
265 Used for compile-time kernel inclusion. >*/
268 #endif /* __ARM_COMPUTE_CLKERNELLIBRARY_EX_H__ */