1 /*
2 * Copyright (c) 2016-2021 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 #include "arm_compute/core/CL/CLKernelLibrary.h"
25 #include "arm_compute/core/Error.h"
26 #include "src/gpu/cl/ClKernelLibrary.h"
27 #include <algorithm>
28 #include <array>
29 #include <fstream>
30 #include <utility>
31 #include <vector>
32 namespace arm_compute
33 {
CLKernelLibrary()34 CLKernelLibrary::CLKernelLibrary()
35 : _compile_context()
36 {
37 opencl_is_available(); // Make sure the OpenCL symbols are initialised *before* the CLKernelLibrary is built
38 }
get()39 CLKernelLibrary &CLKernelLibrary::get()
40 {
41 static CLKernelLibrary _kernel_library;
42 return _kernel_library;
43 }
create_kernel(const std::string & kernel_name,const std::set<std::string> & build_options_set) const44 Kernel CLKernelLibrary::create_kernel(const std::string &kernel_name, const std::set<std::string> &build_options_set) const
45 {
46 const opencl::ClKernelLibrary &klib = opencl::ClKernelLibrary::get();
47 const std::string program_name = klib.program_name(kernel_name);
48 auto program = klib.program(program_name);
49 const std::string &kernel_path = CLKernelLibrary::get().get_kernel_path();
50 return _compile_context.create_kernel(kernel_name, program_name, program.program, kernel_path, build_options_set, program.is_binary);
51 }
get_program_name(const std::string & kernel_name) const52 std::string CLKernelLibrary::get_program_name(const std::string &kernel_name) const
53 {
54 return opencl::ClKernelLibrary::get().program_name(kernel_name);
55 }
init(std::string kernel_path,cl::Context context,cl::Device device)56 void CLKernelLibrary::init(std::string kernel_path, cl::Context context, cl::Device device)
57 {
58 _compile_context = CLCompileContext(context, device);
59 opencl::ClKernelLibrary::get().set_kernel_path(kernel_path);
60 }
set_kernel_path(const std::string & kernel_path)61 void CLKernelLibrary::set_kernel_path(const std::string &kernel_path)
62 {
63 opencl::ClKernelLibrary::get().set_kernel_path(kernel_path);
64 }
context()65 cl::Context &CLKernelLibrary::context()
66 {
67 return _compile_context.context();
68 }
get_device()69 const cl::Device &CLKernelLibrary::get_device()
70 {
71 return _compile_context.get_device();
72 }
set_device(cl::Device device)73 void CLKernelLibrary::set_device(cl::Device device)
74 {
75 _compile_context.set_device(device);
76 }
set_context(cl::Context context)77 void CLKernelLibrary::set_context(cl::Context context)
78 {
79 _compile_context.set_context(context);
80 }
get_kernel_path()81 std::string CLKernelLibrary::get_kernel_path()
82 {
83 return opencl::ClKernelLibrary::get().kernel_path();
84 }
clear_programs_cache()85 void CLKernelLibrary::clear_programs_cache()
86 {
87 _compile_context.clear_programs_cache();
88 }
get_built_programs() const89 const std::map<std::string, cl::Program> &CLKernelLibrary::get_built_programs() const
90 {
91 return _compile_context.get_built_programs();
92 }
add_built_program(const std::string & built_program_name,const cl::Program & program)93 void CLKernelLibrary::add_built_program(const std::string &built_program_name, const cl::Program &program)
94 {
95 _compile_context.add_built_program(built_program_name, program);
96 }
fp16_supported() const97 bool CLKernelLibrary::fp16_supported() const
98 {
99 return _compile_context.fp16_supported();
100 }
int64_base_atomics_supported() const101 bool CLKernelLibrary::int64_base_atomics_supported() const
102 {
103 return _compile_context.int64_base_atomics_supported();
104 }
is_wbsm_supported()105 bool CLKernelLibrary::is_wbsm_supported()
106 {
107 return _compile_context.is_wbsm_supported();
108 }
get_program(const std::string & program_name) const109 std::pair<std::string, bool> CLKernelLibrary::get_program(const std::string &program_name) const
110 {
111 auto program_info = opencl::ClKernelLibrary::get().program(program_name);
112 return std::make_pair(std::move(program_info.program), program_info.is_binary);
113 }
max_local_workgroup_size(const cl::Kernel & kernel) const114 size_t CLKernelLibrary::max_local_workgroup_size(const cl::Kernel &kernel) const
115 {
116 return _compile_context.max_local_workgroup_size(kernel);
117 }
default_ndrange() const118 cl::NDRange CLKernelLibrary::default_ndrange() const
119 {
120 return _compile_context.default_ndrange();
121 }
get_device_version()122 std::string CLKernelLibrary::get_device_version()
123 {
124 return _compile_context.get_device_version();
125 }
get_num_compute_units()126 cl_uint CLKernelLibrary::get_num_compute_units()
127 {
128 return _compile_context.get_num_compute_units();
129 }
get_compile_context()130 CLCompileContext &CLKernelLibrary::get_compile_context()
131 {
132 return _compile_context;
133 }
134 } // namespace arm_compute