compute/ARMComputeEx/src/core/CL/CLKernelLibrary.cpp

   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 #include "arm_compute/core/CL/CLKernelLibrary.h"
  42 #include "arm_compute/core/CL/CLKernelLibraryEx.h"
  43
  44 #include "arm_compute/core/CL/CLHelpers.h"
  45 #include "arm_compute/core/Error.h"
  46 #include "arm_compute/core/Utils.h"
  47
  48 #include <algorithm>
  49 #include <fstream>
  50 #include <iostream>
  51 #include <utility>
  52 #include <vector>
  53
  54 using namespace arm_compute;
  55
  56 const std::map<std::string, std::string> CLKernelLibraryEx::_kernel_program_map = {
  57   // ARMComputeEx kernels
  58   {"arg_min_max_ex_x", "arg_min_max_ex.cl"},
  59   {"arg_min_max_ex_y", "arg_min_max_ex.cl"},
  60   {"arg_min_max_ex_z", "arg_min_max_ex.cl"},
  61   {"arg_min_max_ex_w", "arg_min_max_ex.cl"},
  62   {"binary_logical_op", "binary_logical_op.cl"},
  63   {"cast_bool", "cast.cl"},
  64   {"embedding_lookup", "embedding_lookup.cl"},
  65   {"gather_ex", "gather_ex.cl"},
  66   {"gather_ex_1d", "gather_ex.cl"},
  67   {"gather_ex_1d_out", "gather_ex.cl"},
  68   {"gemmlowp_mm_midgard_ex", "gemmlowp_ex.cl"},
  69   {"gemm_accumulate_biases", "gemm.cl"},
  70   {"hashtable_lookup", "hashtable_lookup.cl"},
  71   {"instance_normalization_ex", "instance_normalization_ex.cl"},
  72   {"memset", "memset.cl"},
  73   {"multiply_scale_factor", "multiply_scale_factor.cl"},
  74   {"neg_tensor", "neg_tensor.cl"},
  75   {"one_hot", "one_hot.cl"},
  76   {"one_hot_only_on_value", "one_hot.cl"},
  77   {"pad_layer_constant", "pad_layer.cl"},
  78   {"pad_layer_symmetric_reflect", "pad_layer.cl"},
  79   {"quantization_symm8", "quantization_symm8.cl"},
  80   {"reduce_min_max", "reduce_operation.cl"},
  81   {"reduce_sum_mean", "reduce_operation.cl"},
  82   {"topkv2_init", "topkv2.cl"},
  83   {"topkv2_find_first_negative", "topkv2.cl"},
  84   {"topkv2_reorder_negatives", "topkv2.cl"},
  85   {"topkv2_store", "topkv2.cl"},
  86   {"topkv2_quicksort", "topkv2_quicksort.cl"},
  87   {"scale_factor_symm8", "scale_factor.cl"},
  88 };
  89
  90 const std::map<std::string, std::string> CLKernelLibraryEx::_program_source_map = {
  91 #ifdef EMBEDDED_KERNELS
  92   {
  93     "activation_float_helpers.h",
  94 #include "./cl_kernels/activation_float_helpers.hembed"
  95   },
  96   {
  97     "arg_min_max_ex.cl",
  98 #include "./cl_kernels/arg_min_max_ex.clembed"
  99   },
 100   {
 101     "binary_logical_op.cl",
 102 #include "./cl_kernels/binary_logical_op.clembed"
 103   },
 104   {
 105     "cast.cl",
 106 #include "./cl_kernels/cast.clembed"
 107   },
 108   {
 109     "embedding_lookup.cl",
 110 #include "./cl_kernels/embedding_lookup.clembed"
 111   },
 112   {
 113     "gather_ex.cl",
 114 #include "./cl_kernels/gather_ex.clembed"
 115   },
 116   {
 117     "gemmlowp_ex.cl",
 118 #include "./cl_kernels/gemmlowp_ex.clembed"
 119   },
 120   {
 121     "gemm_helpers.h",
 122 #include "./cl_kernels/gemm_helpers.hembed"
 123   },
 124   {
 125     "hashtable_lookup.cl",
 126 #include "./cl_kernels/hashtable_lookup.clembed"
 127   },
 128   {
 129     "helpers.h",
 130 #include "./cl_kernels/helpers.hembed"
 131   },
 132   {
 133     "helpers_asymm.h",
 134 #include "./cl_kernels/helpers_asymm.hembed"
 135   },
 136   {
 137     "instance_normalization_ex.cl",
 138 #include "./cl_kernels/instance_normalization_ex.clembed"
 139   },
 140   {
 141     "gemm.cl",
 142 #include "./cl_kernels/gemm.clembed"
 143   },
 144   {
 145     "memset.cl",
 146 #include "./cl_kernels/memset.clembed"
 147   },
 148   {
 149     "multiply_scale_factor.cl",
 150 #include "./cl_kernels/multiply_scale_factor.clembed"
 151   },
 152   {
 153     "neg_tensor.cl",
 154 #include "./cl_kernels/neg_tensor.clembed"
 155   },
 156   {
 157     "one_hot.cl",
 158 #include "./cl_kernels/one_hot.clembed"
 159   },
 160   {
 161     "pad_layer.cl",
 162 #include "./cl_kernels/pad_layer.clembed"
 163   },
 164   {
 165     "quantization_symm8.cl",
 166 #include "./cl_kernels/quantization_symm8.clembed"
 167   },
 168   {
 169     "reduce_operation.cl",
 170 #include "./cl_kernels/reduce_operation.clembed"
 171   },
 172   {
 173     "repeat.h",
 174 #include "./cl_kernels/repeat.hembed"
 175   },
 176   {
 177     "scale_factor.cl",
 178 #include "./cl_kernels/scale_factor.clembed"
 179   },
 180   {
 181     "topkv2.cl",
 182 #include "./cl_kernels/topkv2.clembed"
 183   },
 184   {
 185     "topkv2_quicksort.cl",
 186 #include "./cl_kernels/topkv2_quicksort.clembed"
 187   },
 188
 189 #endif /* EMBEDDED_KERNELS */
 190 };
 191
 192 CLKernelLibraryEx::CLKernelLibraryEx()
 193   : _context(), _device(), _kernel_path("."), _programs_map(), _built_programs_map()
 194 {
 195   opencl_is_available(); // Make sure the OpenCL symbols are initialised *before* the
 196                          // CLKernelLibraryEx is built
 197 }
 198
 199 CLKernelLibraryEx &CLKernelLibraryEx::get()
 200 {
 201   static CLKernelLibraryEx _kernel_library;
 202   return _kernel_library;
 203 }
 204
 205 Kernel CLKernelLibraryEx::create_kernel(const std::string &kernel_name,
 206                                         const StringSet &build_options_set) const
 207 {
 208   // Find which program contains the kernel
 209   auto kernel_program_it = _kernel_program_map.find(kernel_name);
 210
 211   if (_kernel_program_map.end() == kernel_program_it)
 212   {
 213     ARM_COMPUTE_ERROR_VAR("Kernel %s not found in the CLKernelLibrary", kernel_name.c_str());
 214   }
 215   std::string concat_str;
 216
 217   if (fp16_supported())
 218   {
 219     concat_str += " -DARM_COMPUTE_OPENCL_FP16_ENABLED=1 ";
 220   }
 221
 222   if (get_cl_version(_device) == CLVersion::CL20)
 223   {
 224     concat_str += " -cl-std=CL2.0 ";
 225   }
 226   else if (arm_non_uniform_workgroup_supported(_device))
 227   {
 228     concat_str += " -cl-arm-non-uniform-work-group-size ";
 229   }
 230   else
 231   {
 232     ARM_COMPUTE_ERROR("Non uniform workgroup size is not supported!!");
 233   }
 234
 235   // Check if the program has been built before with same build options.
 236   const std::string program_name = kernel_program_it->second;
 237   const std::string build_options = stringify_set(build_options_set) + concat_str;
 238
 239   const std::string built_program_name = program_name + "_" + build_options;
 240   auto built_program_it = _built_programs_map.find(built_program_name);
 241
 242   cl::Program cl_program;
 243
 244   if (_built_programs_map.end() != built_program_it)
 245   {
 246     // If program has been built, retrieve to create kernel from it
 247     cl_program = built_program_it->second;
 248   }
 249   else
 250   {
 251     // Get program
 252     Program program = load_program(program_name);
 253
 254     // Build program
 255     cl_program = program.build(build_options);
 256
 257     // Add built program to internal map
 258     _built_programs_map.emplace(built_program_name, cl_program);
 259   }
 260
 261   // Create and return kernel
 262   return Kernel(kernel_name, cl_program);
 263 }
 264
 265 void CLKernelLibraryEx::add_built_program(const std::string &built_program_name,
 266                                           cl::Program program)
 267 {
 268   _built_programs_map.emplace(built_program_name, program);
 269 }
 270
 271 bool CLKernelLibraryEx::fp16_supported() const { return ::fp16_supported(_device); }
 272
 273 bool CLKernelLibraryEx::int64_base_atomics_supported() const
 274 {
 275   return device_supports_extension(_device, "cl_khr_int64_base_atomics");
 276 }
 277
 278 const Program &CLKernelLibraryEx::load_program(const std::string &program_name) const
 279 {
 280   const auto program_it = _programs_map.find(program_name);
 281
 282   if (program_it != _programs_map.end())
 283   {
 284     return program_it->second;
 285   }
 286
 287   Program program;
 288
 289 #ifdef EMBEDDED_KERNELS
 290   const auto program_source_it = _program_source_map.find(program_name);
 291
 292   if (_program_source_map.end() == program_source_it)
 293   {
 294     ARM_COMPUTE_ERROR_VAR("Embedded program for %s does not exist.", program_name.c_str());
 295   }
 296
 297   program = Program(_context, program_name, program_source_it->second);
 298 #else  /* EMBEDDED_KERNELS */
 299   // Check for binary
 300   std::string source_name = _kernel_path + program_name;
 301   std::string binary_name = source_name + "bin";
 302
 303   if (std::ifstream(binary_name).is_open())
 304   {
 305     const std::string program_binary = read_file(binary_name, true);
 306     program = Program(_context, _device, program_name,
 307                       std::vector<unsigned char>(program_binary.begin(), program_binary.end()));
 308   }
 309   else if (std::ifstream(source_name).is_open())
 310   {
 311     program = Program(_context, program_name, read_file(source_name, false));
 312   }
 313   else
 314   {
 315     ARM_COMPUTE_ERROR_VAR("Kernel file %s does not exist.", source_name.c_str());
 316   }
 317 #endif /* EMBEDDED_KERNELS */
 318
 319   // Insert program to program map
 320   const auto new_program = _programs_map.emplace(program_name, std::move(program));
 321
 322   return new_program.first->second;
 323 }
 324
 325 std::string CLKernelLibraryEx::stringify_set(const StringSet &s) const
 326 {
 327   std::string concat_set;
 328
 329 #ifndef EMBEDDED_KERNELS
 330   concat_set += "-I" + _kernel_path + " ";
 331 #endif /* EMBEDDED_KERNELS */
 332
 333   // Concatenate set
 334   for (const auto &el : s)
 335   {
 336     concat_set += " " + el;
 337   }
 338
 339   return concat_set;
 340 }
 341
 342 std::string CLKernelLibraryEx::get_program_source(const std::string &program_name)
 343 {
 344   const auto program_source_it = _program_source_map.find(program_name);
 345
 346   if (program_source_it == _program_source_map.end())
 347   {
 348     ARM_COMPUTE_ERROR_VAR("Embedded program for %s does not exist.", program_name.c_str());
 349   }
 350
 351   return program_source_it->second;
 352 }
 353
 354 size_t CLKernelLibraryEx::max_local_workgroup_size(const cl::Kernel &kernel) const
 355 {
 356   size_t result;
 357
 358   size_t err = kernel.getWorkGroupInfo(_device, CL_KERNEL_WORK_GROUP_SIZE, &result);
 359   ARM_COMPUTE_ERROR_ON_MSG(
 360     err != 0,
 361     "clGetKernelWorkGroupInfo failed to return the maximum workgroup size for the kernel");
 362   ARM_COMPUTE_UNUSED(err);
 363
 364   return result;
 365 }
 366
 367 cl::NDRange CLKernelLibraryEx::default_ndrange() const
 368 {
 369   //    GPUTarget   _target = get_target_from_device(_device);
 370   cl::Device device = cl::Device::getDefault();
 371   GPUTarget _target = get_target_from_device(device);
 372   cl::NDRange default_range;
 373
 374   switch (_target)
 375   {
 376     case GPUTarget::MIDGARD:
 377     case GPUTarget::T600:
 378     case GPUTarget::T700:
 379     case GPUTarget::T800:
 380       default_range = cl::NDRange(128u, 1);
 381       break;
 382     default:
 383       default_range = cl::NullRange;
 384   }
 385
 386   return default_range;
 387 }
 388
 389 std::string CLKernelLibraryEx::get_device_version() { return _device.getInfo<CL_DEVICE_VERSION>(); }