compute/ARMComputeEx/arm_compute/core/CL/CLKernelLibraryEx.h

   1 /*
   2  * Copyright (c) 2018 Samsung Electronics Co., Ltd. All Rights Reserved
   3  *
   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
   7  *
   8  *      http://www.apache.org/licenses/LICENSE-2.0
   9  *
  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.
  15  */
  16
  17 /*
  18  * Copyright (c) 2016-2018 ARM Limited.
  19  *
  20  * SPDX-License-Identifier: MIT
  21  *
  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:
  28  *
  29  * The above copyright notice and this permission notice shall be included in all
  30  * copies or substantial portions of the Software.
  31  *
  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
  38  * SOFTWARE.
  39  */
  40
  41 /**
  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.
  46  */
  47
  48 #ifndef __ARM_COMPUTE_CLKERNELLIBRARY_EX_H__
  49 #define __ARM_COMPUTE_CLKERNELLIBRARY_EX_H__
  50
  51 #include "arm_compute/core/CL/OpenCL.h"
  52
  53 #include <map>
  54 #include <set>
  55 #include <string>
  56 #include <utility>
  57
  58 namespace arm_compute
  59 {
  60
  61 /**
  62  * @brief Class to build OpenCL kernels added from nnfw
  63  * */
  64 class CLKernelLibraryEx
  65 {
  66   using StringSet = std::set<std::string>;
  67
  68 private:
  69   /**
  70    * @brief Construct a new CLKernelLibraryEx object
  71    */
  72   CLKernelLibraryEx();
  73
  74 public:
  75   /**
  76    * @brief Prevent instances of this class from being copied.
  77    */
  78   CLKernelLibraryEx(const CLKernelLibraryEx &) = delete;
  79
  80   /**
  81    * @brief Prevent instances of this class from being copied.
  82    */
  83   const CLKernelLibraryEx &operator=(const CLKernelLibraryEx &) = delete;
  84
  85   /**
  86    * @brief Get the KernelLibrary singleton.
  87    * @return The KernelLibrary instance
  88    */
  89   static CLKernelLibraryEx &get();
  90
  91   /**
  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.
  96    * @return N/A
  97    */
  98   void init(std::string kernel_path, cl::Context context, cl::Device device)
  99   {
 100     _kernel_path = std::move(kernel_path);
 101     _context = std::move(context);
 102     _device = std::move(device);
 103   }
 104
 105   /**
 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.
 108    * @return N/A
 109    */
 110   void set_kernel_path(const std::string &kernel_path) { _kernel_path = kernel_path; };
 111
 112   /**
 113    * @brief Get the path that the kernels reside in.
 114    * @return the path of kernel files
 115    */
 116   std::string get_kernel_path() { return _kernel_path; };
 117
 118   /**
 119    * @brief Get the source of the selected program.
 120    * @param[in] program_name Program name.
 121    * @return Source of the selected program.
 122    */
 123   std::string get_program_source(const std::string &program_name);
 124
 125   /**
 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.
 130    * @return N/A
 131    */
 132   void set_context(cl::Context context)
 133   {
 134     _context = std::move(context);
 135     if (_context.get() == nullptr)
 136     {
 137       _device = cl::Device();
 138     }
 139     else
 140     {
 141       const auto cl_devices = _context.getInfo<CL_CONTEXT_DEVICES>();
 142
 143       if (cl_devices.empty())
 144       {
 145         _device = cl::Device();
 146       }
 147       else
 148       {
 149         _device = cl_devices[0];
 150       }
 151     }
 152   }
 153
 154   /**
 155    * @brief Return associated CL context.
 156    * @return A CL context.
 157    */
 158   cl::Context &context() { return _context; }
 159
 160   /**
 161    * @brief Set the CL device for which the programs are created.
 162    * @param[in] device A CL device.
 163    * @return N/A
 164    */
 165   void set_device(cl::Device device) { _device = std::move(device); }
 166
 167   /**
 168    * @brief Gets the CL device for which the programs are created.
 169    * @return A CL device.
 170    */
 171   cl::Device &get_device() { return _device; }
 172
 173   /**
 174    * @brief Return the device version
 175    * @return The content of CL_DEVICE_VERSION
 176    */
 177   std::string get_device_version();
 178
 179   /**
 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.
 184    */
 185   Kernel create_kernel(const std::string &kernel_name,
 186                        const StringSet &build_options_set = {}) const;
 187
 188   /**
 189    * @brief Find the maximum number of local work items in a workgroup can be supported for the
 190    * kernel.
 191    * @param[in] kernel       kernel object
 192    */
 193
 194   size_t max_local_workgroup_size(const cl::Kernel &kernel) const;
 195   /**
 196    * @brief Return the default NDRange for the device.
 197    * @return default NDRangeof the device
 198    */
 199   cl::NDRange default_ndrange() const;
 200
 201   /**
 202    * @brief Clear the library's cache of binary programs
 203    * @return N/A
 204    */
 205   void clear_programs_cache()
 206   {
 207     _programs_map.clear();
 208     _built_programs_map.clear();
 209   }
 210
 211   /**
 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)
 215    */
 216   const std::map<std::string, cl::Program> &get_built_programs() const
 217   {
 218     return _built_programs_map;
 219   }
 220
 221   /**
 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
 225    * @return N/A
 226    */
 227   void add_built_program(const std::string &built_program_name, cl::Program program);
 228
 229   /**
 230    * @brief Returns true if FP16 is supported by the CL device
 231    * @return true if the CL device supports FP16
 232    */
 233   bool fp16_supported() const;
 234
 235   /**
 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
 238    */
 239   bool int64_base_atomics_supported() const;
 240
 241 private:
 242   /**
 243    * @brief Load program and its dependencies.
 244    * @param[in] program_name Name of the program to load.
 245    */
 246   const Program &load_program(const std::string &program_name) const;
 247   /**
 248    * @brief Concatenates contents of a set into a single string.
 249    * @param[in] s Input set to concatenate.
 250    * @return Concatenated string.
 251    */
 252   std::string stringify_set(const StringSet &s) const;
 253
 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. >*/
 266 };
 267 }
 268 #endif /* __ARM_COMPUTE_CLKERNELLIBRARY_EX_H__ */