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