1 /* 2 * Copyright (c) 2020-2022 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 */ 24 #ifndef ARM_COMPUTE_CLCOMPILECONTEXT_H 25 #define ARM_COMPUTE_CLCOMPILECONTEXT_H 26 27 #include "arm_compute/core/CL/CLDevice.h" 28 #include "arm_compute/core/CL/OpenCL.h" 29 30 #include <map> 31 #include <set> 32 #include <string> 33 #include <utility> 34 35 namespace arm_compute 36 { 37 /** Build options */ 38 class CLBuildOptions final 39 { 40 using StringSet = std::set<std::string>; 41 42 public: 43 /** Default constructor. */ 44 CLBuildOptions(); 45 /** Adds option to the existing build option list 46 * 47 * @param[in] option Option to add 48 */ 49 void add_option(std::string option); 50 /** Adds option if a given condition is true; 51 * 52 * @param[in] cond Condition to check 53 * @param[in] option Option to add if condition is true 54 */ 55 void add_option_if(bool cond, std::string option); 56 /** Adds first option if condition is true else the second one 57 * 58 * @param[in] cond Condition to check 59 * @param[in] option_true Option to add if condition is true 60 * @param[in] option_false Option to add if condition is false 61 */ 62 void add_option_if_else(bool cond, std::string option_true, std::string option_false); 63 /** Appends given build options to the current's objects options. 64 * 65 * @param[in] options Build options to append 66 */ 67 void add_options(const StringSet &options); 68 /** Appends given build options to the current's objects options if a given condition is true. 69 * 70 * @param[in] cond Condition to check 71 * @param[in] options Option to add if condition is true 72 */ 73 void add_options_if(bool cond, const StringSet &options); 74 /** Gets the current options list set 75 * 76 * @return Build options set 77 */ 78 const StringSet &options() const; 79 80 bool operator==(const CLBuildOptions &other) const; 81 82 private: 83 StringSet _build_opts; /**< Build options set */ 84 }; 85 86 /** Program class */ 87 class Program final 88 { 89 public: 90 /** Default constructor. */ 91 Program(); 92 /** Construct program from source file. 93 * 94 * @param[in] context CL context used to create the program. 95 * @param[in] name Program name. 96 * @param[in] source Program source. 97 */ 98 Program(cl::Context context, std::string name, std::string source); 99 /** Construct program from binary file. 100 * 101 * @param[in] context CL context used to create the program. 102 * @param[in] device CL device for which the programs are created. 103 * @param[in] name Program name. 104 * @param[in] binary Program binary. 105 */ 106 Program(cl::Context context, cl::Device device, std::string name, std::vector<unsigned char> binary); 107 /** Default Copy Constructor. */ 108 Program(const Program &) = default; 109 /** Default Move Constructor. */ 110 Program(Program &&) = default; 111 /** Default copy assignment operator */ 112 Program &operator=(const Program &) = default; 113 /** Default move assignment operator */ 114 Program &operator=(Program &&) = default; 115 /** Returns program name. 116 * 117 * @return Program's name. 118 */ name()119 std::string name() const 120 { 121 return _name; 122 } 123 /** Returns program binary data. 124 * 125 * @return Program's binary data. 126 */ binary()127 const std::vector<unsigned char> &binary() const 128 { 129 return _binary; 130 } 131 /** User-defined conversion to the underlying CL program. 132 * 133 * @return The CL program object. 134 */ 135 explicit operator cl::Program() const; 136 /** Build the given CL program. 137 * 138 * @param[in] program The CL program to build. 139 * @param[in] build_options Options to build the CL program. 140 * 141 * @return True if the CL program builds successfully. 142 */ 143 static bool build(const cl::Program &program, const std::string &build_options = ""); 144 /** Build the underlying CL program. 145 * 146 * @param[in] build_options Options used to build the CL program. 147 * 148 * @return A reference to itself. 149 */ 150 cl::Program build(const std::string &build_options = "") const; 151 152 private: 153 cl::Context _context; /**< Underlying CL context. */ 154 cl::Device _device; /**< CL device for which the programs are created. */ 155 bool _is_binary; /**< Create program from binary? */ 156 std::string _name; /**< Program name. */ 157 std::string _source; /**< Source code for the program. */ 158 std::vector<unsigned char> _binary; /**< Binary from which to create the program. */ 159 }; 160 161 /** Kernel class */ 162 class Kernel final 163 { 164 public: 165 /** Default Constructor. */ 166 Kernel(); 167 /** Default Copy Constructor. */ 168 Kernel(const Kernel &) = default; 169 /** Default Move Constructor. */ 170 Kernel(Kernel &&) = default; 171 /** Default copy assignment operator */ 172 Kernel &operator=(const Kernel &) = default; 173 /** Default move assignment operator */ 174 Kernel &operator=(Kernel &&) = default; 175 /** Constructor. 176 * 177 * @param[in] name Kernel name. 178 * @param[in] program Built program. 179 */ 180 Kernel(std::string name, const cl::Program &program); 181 /** Returns kernel name. 182 * 183 * @return Kernel's name. 184 */ name()185 std::string name() const 186 { 187 return _name; 188 } 189 /** Returns OpenCL kernel. 190 * 191 * @return OpenCL Kernel. 192 */ Kernel()193 explicit operator cl::Kernel() const 194 { 195 return _kernel; 196 } 197 198 private: 199 std::string _name; /**< Kernel name */ 200 cl::Kernel _kernel; /**< OpenCL Kernel */ 201 }; 202 203 /** CLCompileContext class */ 204 class CLCompileContext final 205 { 206 using StringSet = std::set<std::string>; 207 208 public: 209 /** Constructor */ 210 CLCompileContext(); 211 /** Constructor 212 * 213 * @param[in] context A CL context. 214 * @param[in] device A CL device. 215 * */ 216 CLCompileContext(cl::Context context, const cl::Device &device); 217 218 /** Accessor for the associated CL context. 219 * 220 * @return A CL context. 221 */ 222 cl::Context &context(); 223 224 /** Sets the CL context used to create programs. 225 * 226 * @note Setting the context also resets the device to the 227 * first one available in the new context. 228 * 229 * @param[in] context A CL context. 230 */ 231 void set_context(cl::Context context); 232 233 /** Gets the CL device for which the programs are created. */ 234 const cl::Device &get_device() const; 235 236 /** Sets the CL device for which the programs are created. 237 * 238 * @param[in] device A CL device. 239 */ 240 void set_device(cl::Device device); 241 242 /** Creates an OpenCL kernel. 243 * 244 * @param[in] kernel_name Kernel name. 245 * @param[in] program_name Program name. 246 * @param[in] program_source Program source. 247 * @param[in] kernel_path CL kernel path. 248 * @param[in] build_options_set Kernel build options as a set. 249 * @param[in] is_binary Flag to indicate if the program source is binary. 250 * 251 * @return The created kernel. 252 */ 253 Kernel create_kernel(const std::string &kernel_name, const std::string &program_name, const std::string &program_source, 254 const std::string &kernel_path, const StringSet &build_options_set, bool is_binary) const; 255 256 /** Clear the library's cache of binary programs 257 */ 258 void clear_programs_cache(); 259 260 /** Access the cache of built OpenCL programs */ 261 const std::map<std::string, cl::Program> &get_built_programs() const; 262 263 /** Add a new built program to the cache 264 * 265 * @param[in] built_program_name Name of the program 266 * @param[in] program Built program to add to the cache 267 */ 268 void add_built_program(const std::string &built_program_name, const cl::Program &program) const; 269 270 /** Returns true if FP16 is supported by the CL device 271 * 272 * @return true if the CL device supports FP16 273 */ 274 bool fp16_supported() const; 275 276 /** Return the maximum number of compute units in the device 277 * 278 * @return The content of CL_DEVICE_MAX_COMPUTE_UNITS 279 */ 280 cl_uint get_num_compute_units() const; 281 /** Find the maximum number of local work items in a workgroup can be supported for the kernel. 282 * 283 */ 284 size_t max_local_workgroup_size(const cl::Kernel &kernel) const; 285 /** Return the default NDRange for the device. 286 * 287 */ 288 cl::NDRange default_ndrange() const; 289 /** Return the device version 290 * 291 * @return The content of CL_DEVICE_VERSION 292 */ 293 std::string get_device_version() const; 294 295 /** Returns true if int64_base_atomics extension is supported by the CL device 296 * 297 * @return true if the CL device supports int64_base_atomics extension 298 */ 299 bool int64_base_atomics_supported() const; 300 301 /* Returns true if the workgroup batch size modifier parameter is supported on the cl device 302 * 303 * @return true if the workgroup batch size modifier parameter is supported, false otherwise 304 */ 305 bool is_wbsm_supported() const; 306 307 /** Return the DDK version. If the DDK version cannot be detected, return -1. 308 * 309 * @return The DDK version. 310 */ 311 int32_t get_ddk_version() const; 312 313 /** Return the Gpu target of the associated device 314 * 315 * @return GPUTarget 316 */ 317 GPUTarget get_gpu_target() const; 318 319 private: 320 /** Load program and its dependencies. 321 * 322 * @param[in] program_name Name of the program to load. 323 * @param[in] program_source Source of the program. 324 * @param[in] is_binary Flag to indicate if the program source is binary. 325 */ 326 const Program &load_program(const std::string &program_name, const std::string &program_source, bool is_binary) const; 327 328 /** Generates the build options given a string of user defined ones 329 * 330 * @param[in] build_options User defined build options 331 * @param[in] kernel_path Path of the CL kernels 332 * 333 * @return Generated build options 334 */ 335 std::string generate_build_options(const StringSet &build_options, const std::string &kernel_path) const; 336 337 /** Concatenates contents of a set into a single string. 338 * 339 * @param[in] s Input set to concatenate. 340 * @param[in] kernel_path Path of the CL kernels 341 * 342 * @return Concatenated string. 343 */ 344 std::string stringify_set(const StringSet &s, const std::string &kernel_path) const; 345 346 cl::Context _context; /**< Underlying CL context. */ 347 CLDevice _device; /**< Underlying CL device. */ 348 mutable std::map<std::string, const Program> _programs_map; /**< Map with all already loaded program data. */ 349 mutable std::map<std::string, cl::Program> _built_programs_map; /**< Map with all already built program data. */ 350 bool _is_wbsm_supported; /**< Support of worksize batch size modifier support boolean*/ 351 }; 352 } // namespace arm_compute 353 #endif /* ARM_COMPUTE_CLCOMPILECONTEXT_H */ 354