blob: 0b59ec8a7131412c5fabbc4d072cf42495f6de49 [file] [log] [blame]
Anthony Barbier6ff3b192017-09-04 18:44:23 +01001/*
Michele Di Giorgiod9eaf612020-07-08 11:12:57 +01002 * Copyright (c) 2016-2020 Arm Limited.
Anthony Barbier6ff3b192017-09-04 18:44:23 +01003 *
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
steniu0134702472017-07-11 09:22:58 +010026#include "arm_compute/core/CL/CLHelpers.h"
Anthony Barbier6ff3b192017-09-04 18:44:23 +010027#include "arm_compute/core/Error.h"
28#include "arm_compute/core/Utils.h"
Matthew Bentham758b5ba2020-03-05 23:37:48 +000029#include "support/StringSupport.h"
Anthony Barbier6ff3b192017-09-04 18:44:23 +010030
steniu015f910722017-08-23 10:15:22 +010031#include <algorithm>
Anthony Barbier6ff3b192017-09-04 18:44:23 +010032#include <fstream>
33#include <iostream>
34#include <utility>
35#include <vector>
36
37using namespace arm_compute;
Anthony Barbier6ff3b192017-09-04 18:44:23 +010038const std::map<std::string, std::string> CLKernelLibrary::_kernel_program_map =
39{
40 { "absdiff", "absdiff.cl" },
41 { "accumulate", "accumulate.cl" },
42 { "accumulate_squared", "accumulate.cl" },
43 { "accumulate_weighted", "accumulate.cl" },
44 { "activation_layer", "activation_layer.cl" },
Manuel Bottini30dbeef2019-06-26 16:23:03 +010045 { "activation_layer_quant", "activation_layer_quant.cl" },
46 { "activation_layer_quant_f32", "activation_layer_quant.cl" },
Manuel Bottini7b9998d2019-10-21 17:59:07 +010047 { "arg_min_max_x", "arg_min_max.cl" },
48 { "arg_min_max_y", "arg_min_max.cl" },
49 { "arg_min_max_z", "arg_min_max.cl" },
50 { "arg_min_max_w", "arg_min_max.cl" },
Michalis Spyrouf1addb62018-09-11 11:16:47 +010051 { "batch_to_space_nchw", "batch_to_space.cl" },
52 { "batch_to_space_static_nchw", "batch_to_space.cl" },
53 { "batch_to_space_nhwc", "batch_to_space.cl" },
54 { "batch_to_space_static_nhwc", "batch_to_space.cl" },
Michele Di Giorgiobf3c6622018-03-08 11:52:27 +000055 { "batchnormalization_layer_nchw", "batchnormalization_layer.cl" },
56 { "batchnormalization_layer_nhwc", "batchnormalization_layer.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +010057 { "bitwise_or", "bitwise_op.cl" },
58 { "bitwise_and", "bitwise_op.cl" },
59 { "bitwise_xor", "bitwise_op.cl" },
60 { "bitwise_not", "bitwise_op.cl" },
giuros01c04a0e82018-10-03 12:44:35 +010061 { "bounding_box_transform", "bounding_box_transform.cl" },
Michele Di Giorgio4aff98f2019-08-28 16:27:26 +010062 { "bounding_box_transform_quantized", "bounding_box_transform_quantized.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +010063 { "channel_combine_NV", "channel_combine.cl" },
64 { "channel_combine_RGB888", "channel_combine.cl" },
65 { "channel_combine_RGBA8888", "channel_combine.cl" },
66 { "channel_combine_UYVY422", "channel_combine.cl" },
67 { "channel_combine_YUYV422", "channel_combine.cl" },
Michele Di Giorgio72175632018-05-01 16:52:00 +010068 { "channel_shuffle_nchw", "channel_shuffle.cl" },
Gian Marco Iodice8bab0ee2018-09-13 11:51:56 +010069 { "channel_shuffle_nhwc", "channel_shuffle.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +010070 { "channel_extract_NV12", "channel_extract.cl" },
71 { "channel_extract_NV21", "channel_extract.cl" },
72 { "channel_extract_RGB888", "channel_extract.cl" },
73 { "channel_extract_RGBA8888", "channel_extract.cl" },
74 { "channel_extract_UYVY422", "channel_extract.cl" },
75 { "channel_extract_YUYV422", "channel_extract.cl" },
76 { "combine_gradients_L1", "canny.cl" },
77 { "combine_gradients_L2", "canny.cl" },
Georgios Pinitas7900a9e2018-11-23 11:44:58 +000078 { "compare_equal", "comparisons.cl" },
79 { "compare_equal_quantized", "comparisons.cl" },
80 { "compare_notequal", "comparisons.cl" },
81 { "compare_notequal_quantized", "comparisons.cl" },
82 { "compare_greater", "comparisons.cl" },
83 { "compare_greater_quantized", "comparisons.cl" },
84 { "compare_greaterequal", "comparisons.cl" },
85 { "compare_greaterequal_quantized", "comparisons.cl" },
86 { "compare_less", "comparisons.cl" },
87 { "compare_less_quantized", "comparisons.cl" },
88 { "compare_lessequal", "comparisons.cl" },
89 { "compare_lessequal_quantized", "comparisons.cl" },
Vidhya Sudhan Loganathan338595b2019-06-28 14:09:53 +010090 { "concatenate", "concatenate.cl" },
Michalis Spyrou55b3d122018-05-09 09:59:23 +010091 { "concatenate_width", "concatenate.cl" },
Pablo Tello6a14adb2019-03-05 17:33:08 +000092 { "concatenate_height", "concatenate.cl" },
Michele Di Giorgio27400b92018-11-01 13:44:05 +000093 { "concatenate_width_x2", "concatenate.cl" },
94 { "concatenate_width_x4", "concatenate.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +010095 { "convolution_rectangle", "convolution_rectangle.cl" },
Gian Marco76faef82018-01-29 12:15:32 +000096 { "col2im", "col2im.cl" },
Giorgio Arena657bdb32018-04-26 18:52:01 +010097 { "convert_depth_down", "depth_convert.cl" },
98 { "convert_depth_up", "depth_convert.cl" },
99 { "convert_fc_weights", "convert_fc_weights.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100100 { "convolution3x3_static", "convolution3x3.cl" },
101 { "convolution5x5_static", "convolution5x5.cl" },
102 { "convolution7x7_static", "convolution7x7.cl" },
103 { "convolution9x9_static", "convolution9x9.cl" },
104 { "convolution_separable1x5_static", "convolution5x5.cl" },
105 { "convolution_separable5x1_static", "convolution5x5.cl" },
106 { "convolution_separable1x7_static", "convolution7x7.cl" },
107 { "convolution_separable7x1_static", "convolution7x7.cl" },
108 { "convolution_separable1x9_static", "convolution9x9.cl" },
109 { "convolution_separable9x1_static", "convolution9x9.cl" },
Michalis Spyrou5c8e05c2018-03-22 11:56:01 +0000110 { "copy_tensor", "copy_tensor.cl" },
Giuseppe Rossinid7647d42018-07-17 18:13:13 +0100111 { "copy_pad_tensor", "copy_tensor.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100112 { "copy_plane", "channel_extract.cl" },
113 { "copy_planes_3p", "channel_combine.cl" },
114 { "copy_to_keypoint", "fast_corners.cl" },
George Wort894066d2019-02-15 15:12:52 +0000115 { "crop_tensor", "crop_tensor.cl" },
giuros0146a49a02019-04-01 13:50:22 +0100116 { "deconvolution_reshape", "deconvolution_layer.cl" },
Michalis Spyrou780db4e2017-11-23 09:49:51 +0000117 { "deconvolution_upsample", "deconvolution_layer.cl" },
Giorgio Arena93a690e2017-08-01 16:09:33 +0100118 { "depthwise_convolution_3x3", "depthwise_convolution.cl" },
Michele Di Giorgio933fe862018-02-19 15:42:12 +0000119 { "depthwise_convolution_3x3_f16", "depthwise_convolution.cl" },
Giorgio Arenad051e972018-06-20 11:46:42 +0100120 { "depthwise_convolution_3x3_nhwc", "depthwise_convolution.cl" },
121 { "depthwise_convolution_3x3_nhwc_stride1", "depthwise_convolution.cl" },
Gian Marco Iodice9285adb2019-09-05 16:10:27 +0100122 { "dwc_MxN_native_fp_nhwc", "depthwise_convolution.cl" },
Michele Di Giorgioa046e162019-10-08 09:36:26 +0100123 { "dwc_MxN_native_quantized8_nhwc", "depthwise_convolution_quantized.cl" },
Michele Di Giorgiodf4cf572019-10-09 15:32:39 +0100124 { "dwc_3x3_native_quantized8_nchw", "depthwise_convolution_quantized.cl" },
125 { "dwc_3x3_native_quantized8_dot8_nchw", "depthwise_convolution_quantized.cl" },
126 { "dwc_3x3_reshaped_quantized8_nhwc", "depthwise_convolution_quantized.cl" },
127 { "dwc_3x3_reshaped_quantized8_stride1_nhwc", "depthwise_convolution_quantized.cl" },
128 { "dwc_3x3_reshaped_quantized8_dot8_stride1_nhwc", "depthwise_convolution_quantized.cl" },
Michalis Spyrou649962c2019-05-22 11:11:55 +0100129 { "depth_to_space_nchw", "depth_to_space.cl" },
130 { "depth_to_space_nhwc", "depth_to_space.cl" },
Michele Di Giorgio3ebef322018-02-21 10:02:58 +0000131 { "depthwise_convolution_3x3_stridex1_stridey1_bifrost_f16", "depthwise_convolution.cl" },
132 { "depthwise_convolution_3x3_stridex2_stridey2_bifrost_f16", "depthwise_convolution.cl" },
133 { "depthwise_convolution_3x3_stridex1_stridey1_bifrost_f32", "depthwise_convolution.cl" },
134 { "depthwise_convolution_3x3_stridex2_stridey2_bifrost_f32", "depthwise_convolution.cl" },
giuros016d109962019-01-07 17:47:19 +0000135 { "depthwise_convolution_reshape_weights", "depthwise_convolution.cl" },
Michele Di Giorgio56dd7262017-07-27 09:53:49 +0100136 { "dequantization_layer", "dequantization_layer.cl" },
Michalis Spyrou3f632f32019-08-22 16:52:00 +0100137 { "dequantization_layer_per_channel_nhwc", "dequantization_layer.cl" },
138 { "dequantization_layer_per_channel_nchw", "dequantization_layer.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100139 { "derivative", "derivative.cl" },
140 { "dilate", "dilate.cl" },
SiCong Lic51b72f2017-07-28 14:46:20 +0100141 { "direct_convolution1x1", "direct_convolution1x1.cl" },
Pablo Tello3d319462018-06-21 15:13:17 +0100142 { "direct_convolution1x1_nhwc", "direct_convolution1x1.cl" },
Gian Marco Iodice1c8409d2017-09-06 17:24:25 +0100143 { "direct_convolution1x1_f32_bifrost", "direct_convolution1x1.cl" },
SiCong Lic51b72f2017-07-28 14:46:20 +0100144 { "direct_convolution3x3", "direct_convolution3x3.cl" },
Pablo Tello3d319462018-06-21 15:13:17 +0100145 { "direct_convolution3x3_nhwc", "direct_convolution3x3.cl" },
Gian Marco Iodice1246b632017-08-16 18:38:32 +0100146 { "direct_convolution3x3_f32_bifrost", "direct_convolution3x3.cl" },
steniu01db006682017-08-09 16:26:22 +0100147 { "direct_convolution5x5", "direct_convolution5x5.cl" },
Pablo Tello3d319462018-06-21 15:13:17 +0100148 { "direct_convolution5x5_nhwc", "direct_convolution5x5.cl" },
Gian Marco Iodice1246b632017-08-16 18:38:32 +0100149 { "direct_convolution5x5_f32_bifrost", "direct_convolution5x5.cl" },
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100150 { "direct_convolution_quantized", "direct_convolution_quantized.cl" },
Michalis Spyrou45091732019-05-13 17:41:01 +0100151 { "direct_convolution9x9_nhwc", "direct_convolution9x9.cl" },
giuros01164a2722018-11-20 18:34:46 +0000152 { "elementwise_operation_ADD", "elementwise_operation.cl" },
153 { "elementwise_operation_SUB", "elementwise_operation.cl" },
154 { "elementwise_operation_MAX", "elementwise_operation.cl" },
155 { "elementwise_operation_MIN", "elementwise_operation.cl" },
156 { "elementwise_operation_DIV", "elementwise_operation.cl" },
157 { "elementwise_operation_SQUARED_DIFF", "elementwise_operation.cl" },
Usama Arif52c54f62019-05-14 10:22:36 +0100158 { "elementwise_operation_POWER", "elementwise_operation.cl" },
giuros011e6e1b82019-05-14 16:12:53 +0100159 { "elementwise_operation_PRELU", "elementwise_operation.cl" },
giuros01164a2722018-11-20 18:34:46 +0000160 { "elementwise_operation_ADD_quantized", "elementwise_operation_quantized.cl" },
161 { "elementwise_operation_SUB_quantized", "elementwise_operation_quantized.cl" },
162 { "elementwise_operation_MAX_quantized", "elementwise_operation_quantized.cl" },
163 { "elementwise_operation_MIN_quantized", "elementwise_operation_quantized.cl" },
164 { "elementwise_operation_DIV_quantized", "elementwise_operation_quantized.cl" },
165 { "elementwise_operation_SQUARED_DIFF_quantized", "elementwise_operation_quantized.cl" },
giuros011e6e1b82019-05-14 16:12:53 +0100166 { "elementwise_operation_PRELU_quantized", "elementwise_operation_quantized.cl" },
Michalis Spyroue9362622018-11-23 17:41:37 +0000167 { "elementwise_unary", "elementwise_unary.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100168 { "erode", "erode.cl" },
169 { "fast_corners", "fast_corners.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000170 { "fft_digit_reverse_axis_0", "fft_digit_reverse.cl" },
171 { "fft_digit_reverse_axis_1", "fft_digit_reverse.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000172 { "fft_radix_2_first_stage_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000173 { "fft_radix_2_first_stage_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000174 { "fft_radix_2_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000175 { "fft_radix_2_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000176 { "fft_radix_3_first_stage_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000177 { "fft_radix_3_first_stage_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000178 { "fft_radix_3_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000179 { "fft_radix_3_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000180 { "fft_radix_4_first_stage_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000181 { "fft_radix_4_first_stage_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000182 { "fft_radix_4_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000183 { "fft_radix_4_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000184 { "fft_radix_5_first_stage_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000185 { "fft_radix_5_first_stage_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000186 { "fft_radix_5_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000187 { "fft_radix_5_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000188 { "fft_radix_7_first_stage_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000189 { "fft_radix_7_first_stage_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000190 { "fft_radix_7_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000191 { "fft_radix_7_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000192 { "fft_radix_8_first_stage_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000193 { "fft_radix_8_first_stage_axis_1", "fft.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000194 { "fft_radix_8_axis_0", "fft.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000195 { "fft_radix_8_axis_1", "fft.cl" },
196 { "fft_scale_conj", "fft_scale.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100197 { "fill_image_borders_constant", "fill_border.cl" },
198 { "fill_image_borders_replicate", "fill_border.cl" },
199 { "finalize", "optical_flow_pyramid_lk.cl" },
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000200 { "flatten", "flatten.cl" },
Georgios Pinitasd8e765b2017-08-02 13:44:33 +0100201 { "floor_layer", "floor.cl" },
Manuel Bottini2732cca2019-05-28 11:44:41 +0100202 { "fuse_batchnormalization_layer", "batchnormalization_layer.cl" },
Manuel Bottini8529bd62018-11-21 11:53:04 +0000203 { "gather", "gather.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100204 { "gaussian1x5_sub_x", "gaussian_pyramid.cl" },
205 { "gaussian5x1_sub_y", "gaussian_pyramid.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100206 { "gemm_ma_f16", "gemm.cl" },
207 { "gemm_ma_f32", "gemm.cl" },
Giorgio Arena9fe41442017-08-23 16:36:24 +0100208 { "gemm_mv", "gemv.cl" },
Georgios Pinitasde5a1cc2018-02-02 12:52:07 +0000209 { "gemm_mv_quantized", "gemv.cl" },
Gian Marco Iodiceedfa9f42017-08-15 11:45:22 +0100210 { "gemm_mm_interleaved_transposed_f16", "gemm.cl" },
Vidhya Sudhan Loganathan38d93bd2018-11-20 15:38:13 +0000211 { "gemm_mm_interleaved_transposed_f16_acc32", "gemm.cl" },
Gian Marco Iodicebb36a8e2018-04-19 12:05:08 +0100212 { "gemm_mm_interleaved_transposed_f16_bifrost", "gemm.cl" },
213 { "gemm_mm_interleaved_transposed_f32", "gemm.cl" },
Gian Marco Iodiceedfa9f42017-08-15 11:45:22 +0100214 { "gemm_mm_interleaved_transposed_f32_bifrost", "gemm.cl" },
Gian Marco Iodiceedfa9f42017-08-15 11:45:22 +0100215 { "gemm_mm_floating_point", "gemm.cl" },
Gian Marco Iodicefd683112018-04-17 09:52:44 +0100216 { "gemm_mm_floating_point_f16_bifrost", "gemm.cl" },
Vidhya Sudhan Loganathana25d16c2018-11-16 11:33:12 +0000217 { "gemm_mm_floating_point_f16_bifrost_acc32", "gemm.cl" },
Anton Lokhmotov3e80c7f2017-11-20 11:02:10 +0000218 { "gemm_mm_floating_point_f32_bifrost", "gemm.cl" },
219 { "gemm_mm_floating_point_f32_bifrost_1000", "gemm.cl" },
giuros01b3204e72019-04-01 13:50:22 +0100220 { "gemm_mm_native", "gemm.cl" },
Gian Marco Iodicebf9731e2018-12-12 10:18:04 +0000221 { "gemm_mm_reshaped_lhs_nt_rhs_t", "gemm.cl" },
Gian Marco Iodicee3a849a2020-06-10 17:59:30 +0100222 { "gemm_mm_reshaped_lhs_nt_rhs_t_texture", "gemm.cl" },
Giorgio Arenaae99b6e2019-08-01 14:22:12 +0100223 { "gemm_mm_reshaped_lhs_t_rhs_nt", "gemm.cl" },
Gian Marco Iodicee3a849a2020-06-10 17:59:30 +0100224 { "gemm_mm_reshaped_lhs_t_rhs_nt_texture", "gemm.cl" },
Gian Marco Iodiceba5e0962019-03-11 12:17:44 +0000225 { "gemm_mm_reshaped_only_rhs_nt", "gemm.cl" },
Gian Marco Iodice781cba72020-06-19 16:56:57 +0100226 { "gemm_mm_reshaped_only_rhs_nt_texture", "gemm.cl" },
Gian Marco Iodiceadc53952019-02-15 11:10:31 +0000227 { "gemm_mm_reshaped_only_rhs_t", "gemm.cl" },
Gian Marco Iodice781cba72020-06-19 16:56:57 +0100228 { "gemm_mm_reshaped_only_rhs_t_texture", "gemm.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100229 { "gemm_lc_vm_f32", "gemm.cl" },
Gian Marco Iodice5ba5e092018-12-06 17:13:09 +0000230 { "gemm_reshape_lhs_matrix_nt", "gemm.cl" },
Gian Marco Iodice08ddd7b2018-12-19 10:01:18 +0000231 { "gemm_reshape_lhs_matrix_t", "gemm.cl" },
Gian Marco Iodice3b0a2652018-12-07 11:18:09 +0000232 { "gemm_reshape_rhs_matrix_nt", "gemm.cl" },
233 { "gemm_reshape_rhs_matrix_t", "gemm.cl" },
Gian Marco05288a22017-11-21 10:57:50 +0000234 { "gemmlowp_matrix_a_reduction", "gemmlowp.cl" },
Gian Marco Iodice4b908652018-10-18 10:21:02 +0100235 { "gemmlowp_matrix_a_reduction_dot8", "gemmlowp.cl" },
Gian Marco05288a22017-11-21 10:57:50 +0000236 { "gemmlowp_matrix_b_reduction", "gemmlowp.cl" },
Gian Marco Iodicee7510622019-06-03 17:28:17 +0100237 { "gemmlowp_mm_native", "gemmlowp.cl" },
Gian Marco Iodicedb63b9c2019-01-17 09:47:04 +0000238 { "gemmlowp_mm_reshaped_lhs_nt_rhs_t", "gemmlowp.cl" },
Gian Marco Iodice62251f72019-03-11 16:07:12 +0000239 { "gemmlowp_mm_reshaped_only_rhs_t", "gemmlowp.cl" },
Michele Di Giorgiob54ba282020-01-14 15:31:55 +0000240 { "gemmlowp_mm_reshaped_only_rhs_t_fused_output_stage_fixedpoint", "gemmlowp.cl" },
Gian Marco05288a22017-11-21 10:57:50 +0000241 { "gemmlowp_offset_contribution", "gemmlowp.cl" },
Gian Marco Iodice4b908652018-10-18 10:21:02 +0100242 { "gemmlowp_offset_contribution_quantize_down", "gemmlowp.cl" },
243 { "gemmlowp_offset_contribution_quantize_down_fixedpoint", "gemmlowp.cl" },
Gian Marco05288a22017-11-21 10:57:50 +0000244 { "gemmlowp_output_stage_quantize_down", "gemmlowp.cl" },
Gian Marco58c57942017-11-28 09:10:03 +0000245 { "gemmlowp_output_stage_quantize_down_fixedpoint", "gemmlowp.cl" },
Manuel Bottini9c9b70b2019-07-01 17:35:56 +0100246 { "gemmlowp_output_stage_quantize_down_fixedpoint_qsymm16", "gemmlowp.cl" },
Georgios Pinitas51e53a32018-10-22 13:49:08 +0100247 { "gemmlowp_output_stage_quantize_down_float", "gemmlowp.cl" },
Manuel Bottini5209be52019-02-13 16:34:56 +0000248 { "generate_proposals_compute_all_anchors", "generate_proposals.cl" },
Michele Di Giorgio6b612f52019-09-05 12:30:22 +0100249 { "generate_proposals_compute_all_anchors_quantized", "generate_proposals_quantized.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100250 { "harris_score_3x3", "harris_corners.cl" },
251 { "harris_score_5x5", "harris_corners.cl" },
252 { "harris_score_7x7", "harris_corners.cl" },
253 { "hist_border_kernel", "histogram.cl" },
254 { "hist_border_kernel_fixed", "histogram.cl" },
255 { "hist_local_kernel", "histogram.cl" },
256 { "hist_local_kernel_fixed", "histogram.cl" },
257 { "hog_block_normalization", "hog.cl" },
258 { "hog_detector", "hog.cl" },
259 { "hog_orientation_binning", "hog.cl" },
260 { "hysteresis", "canny.cl" },
Gian Marco Iodice215b4ea2018-06-28 16:29:29 +0100261 { "im2col1x1_stridex1_nchw", "im2col.cl" },
262 { "im2col3x3_nchw", "im2col.cl" },
263 { "im2col5x5_nchw", "im2col.cl" },
264 { "im2col11x11_padx0_pady0_nchw", "im2col.cl" },
265 { "im2col_generic_nchw", "im2col.cl" },
266 { "im2col_generic_padx0_pady0_nchw", "im2col.cl" },
Pablo Tello4a626a72018-04-04 10:01:14 +0100267 { "im2col3x3_nhwc", "im2col.cl" },
Gian Marco Iodicebf9731e2018-12-12 10:18:04 +0000268 { "im2col9x9_nhwc", "im2col.cl" },
Pablo Tello4a626a72018-04-04 10:01:14 +0100269 { "im2col_generic_nhwc", "im2col.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100270 { "init_level", "optical_flow_pyramid_lk.cl" },
271 { "init_level_max", "optical_flow_pyramid_lk.cl" },
272 { "init_level_max_initial_estimate", "optical_flow_pyramid_lk.cl" },
Manuel Bottini79f88e62019-09-18 15:02:53 +0100273 { "instance_normalization", "instance_normalization.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100274 { "integral_horizontal", "integral_image.cl" },
275 { "integral_vertical", "integral_image.cl" },
276 { "IYUV_to_NV12_bt709", "color_convert.cl" },
277 { "IYUV_to_RGB888_bt709", "color_convert.cl" },
278 { "IYUV_to_RGBA8888_bt709", "color_convert.cl" },
279 { "IYUV_to_YUV444_bt709", "color_convert.cl" },
Michalis Spyrou5538d342018-11-14 08:10:13 +0000280 { "l2_normalize_x", "l2_normalize.cl" },
281 { "l2_normalize_y", "l2_normalize.cl" },
282 { "l2_normalize_z", "l2_normalize.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100283 { "lktracker_stage0", "optical_flow_pyramid_lk.cl" },
284 { "lktracker_stage1", "optical_flow_pyramid_lk.cl" },
285 { "magnitude_phase", "magnitude_phase.cl" },
Gian Marco Iodice4d81d752020-07-14 15:05:31 +0100286 { "max_unpooling_layer_2", "unpooling_layer.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100287 { "mean_stddev_accumulate", "mean_stddev.cl" },
Michele Di Giorgio5b48ad72019-06-04 18:43:35 +0100288 { "mean_stddev_normalization", "mean_stddev_normalization.cl" },
Giuseppe Rossinid7647d42018-07-17 18:13:13 +0100289 { "memset", "memset.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100290 { "minmax", "minmaxloc.cl" },
291 { "minmax_border", "minmaxloc.cl" },
Michele Di Giorgio56dd7262017-07-27 09:53:49 +0100292 { "minmax_layer", "minmax_layer.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100293 { "minmaxloc", "minmaxloc.cl" },
294 { "non_linear_filter_box3x3", "non_linear_filter3x3.cl" },
295 { "non_linear_filter_cross3x3", "non_linear_filter3x3.cl" },
296 { "non_linear_filter_disk3x3", "non_linear_filter3x3.cl" },
297 { "non_linear_filter_box5x5", "non_linear_filter5x5.cl" },
298 { "non_linear_filter_cross5x5", "non_linear_filter5x5.cl" },
299 { "non_linear_filter_disk5x5", "non_linear_filter5x5.cl" },
300 { "non_max_suppression", "nonmax.cl" },
301 { "normalization_layer_cross_map", "normalization_layer.cl" },
Michele Di Giorgio9d3a8312018-11-20 12:31:24 +0000302 { "normalization_layer_in_map_nchw", "normalization_layer.cl" },
303 { "normalization_layer_in_map_nhwc", "normalization_layer.cl" },
Michele Di Giorgiob57be0d2018-08-31 16:26:25 +0100304 { "normalize_planar_yuv_layer_nchw", "normalize_planar_yuv_layer.cl" },
305 { "normalize_planar_yuv_layer_nhwc", "normalize_planar_yuv_layer.cl" },
Michele Di Giorgiod63dfa22018-09-12 10:18:54 +0100306 { "normalize_planar_yuv_layer_q8_nchw", "normalize_planar_yuv_layer_quantized.cl" },
307 { "normalize_planar_yuv_layer_q8_nhwc", "normalize_planar_yuv_layer_quantized.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100308 { "NV12_to_IYUV_bt709", "color_convert.cl" },
309 { "NV12_to_RGB888_bt709", "color_convert.cl" },
310 { "NV12_to_RGBA8888_bt709", "color_convert.cl" },
311 { "NV12_to_YUV444_bt709", "color_convert.cl" },
312 { "NV21_to_IYUV_bt709", "color_convert.cl" },
313 { "NV21_to_RGB888_bt709", "color_convert.cl" },
314 { "NV21_to_RGBA8888_bt709", "color_convert.cl" },
315 { "NV21_to_YUV444_bt709", "color_convert.cl" },
Giorgio Arena5c4a8e92019-08-28 17:55:07 +0100316 { "pad_layer_constant", "pad_layer.cl" },
317 { "pad_layer_symmetric_reflect", "pad_layer.cl" },
shubhame1a4e372019-01-07 21:37:55 +0530318 { "permute", "permute.cl" },
Georgios Pinitas8be91482019-03-26 17:23:28 +0000319 { "pixelwise_mul_complex", "pixelwise_mul_float.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100320 { "pixelwise_mul_float", "pixelwise_mul_float.cl" },
321 { "pixelwise_mul_int", "pixelwise_mul_int.cl" },
Georgios Pinitasbf28a3c2018-09-18 14:34:48 +0100322 { "pixelwise_mul_quantized", "pixelwise_mul_int.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100323 { "pooling_layer_2", "pooling_layer.cl" },
324 { "pooling_layer_3", "pooling_layer.cl" },
Anton Lokhmotovaf6204c2017-11-08 09:34:19 +0000325 { "pooling_layer_optimized_3", "pooling_layer.cl" },
Georgios Pinitasce093142017-06-19 16:11:53 +0100326 { "pooling_layer_7", "pooling_layer.cl" },
Michalis Spyroue74b2012018-04-18 09:49:16 +0100327 { "pooling_layer_MxN_nchw", "pooling_layer.cl" },
328 { "pooling_layer_MxN_nhwc", "pooling_layer.cl" },
Sheri Zhang801bbcb2020-08-03 20:11:56 +0100329 { "pooling_layer_2_nhwc_indices_fp32", "pooling_layer.cl" },
330 { "pooling_layer_2_nhwc_indices_fp16", "pooling_layer.cl" },
331 { "pooling_layer_2_nchw_indices_fp32", "pooling_layer.cl" },
332 { "pooling_layer_2_nchw_indices_fp16", "pooling_layer.cl" },
Michalis Spyroue74b2012018-04-18 09:49:16 +0100333 { "pooling_layer_MxN_quantized_nhwc", "pooling_layer_quantized.cl" },
334 { "pooling_layer_MxN_quantized_nchw", "pooling_layer_quantized.cl" },
Michalis Spyrou6c7c38e2018-08-29 16:28:11 +0100335 { "prior_box_layer_nchw", "prior_box_layer.cl" },
Sheri Zhangb18252d2020-04-07 11:04:57 +0100336 { "qlstm_layer_normalization", "qlstm_layer_normalization.cl" },
Michele Di Giorgio56dd7262017-07-27 09:53:49 +0100337 { "quantization_layer", "quantization_layer.cl" },
Vidhya Sudhan Loganathan5e96be72018-12-18 14:17:00 +0000338 { "range", "range.cl" },
339 { "range_quantized", "range.cl" },
Michalis Spyrou7e9391b2018-10-05 14:49:28 +0100340 { "reduction_operation_x", "reduction_operation.cl" },
Michalis Spyrou7930db42018-11-22 17:36:28 +0000341 { "reduction_operation_non_parallel_x", "reduction_operation.cl" },
Michalis Spyrou7e9391b2018-10-05 14:49:28 +0100342 { "reduction_operation_y", "reduction_operation.cl" },
343 { "reduction_operation_z", "reduction_operation.cl" },
344 { "reduction_operation_w", "reduction_operation.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100345 { "remap_nearest_neighbour", "remap.cl" },
346 { "remap_bilinear", "remap.cl" },
Gian Marco Iodice477531c2018-08-21 17:53:38 +0100347 { "reorg_layer_nchw", "reorg_layer.cl" },
348 { "reorg_layer_nhwc", "reorg_layer.cl" },
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +0100349 { "reshape_layer", "reshape_layer.cl" },
Gian Marco Iodice215b4ea2018-06-28 16:29:29 +0100350 { "reshape_to_columns", "convolution_layer.cl" },
Michele Di Giorgio5daeffd2018-11-26 10:01:15 +0000351 { "reverse", "reverse.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100352 { "RGB888_to_IYUV_bt709", "color_convert.cl" },
353 { "RGB888_to_NV12_bt709", "color_convert.cl" },
354 { "RGB888_to_RGBA8888_bt709", "color_convert.cl" },
Manuel Bottiniacaf21d2018-09-26 17:38:19 +0100355 { "RGB888_to_U8_bt709", "color_convert.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100356 { "RGB888_to_YUV444_bt709", "color_convert.cl" },
357 { "RGBA8888_to_IYUV_bt709", "color_convert.cl" },
358 { "RGBA8888_to_NV12_bt709", "color_convert.cl" },
359 { "RGBA8888_to_RGB888_bt709", "color_convert.cl" },
360 { "RGBA8888_to_YUV444_bt709", "color_convert.cl" },
giuros0118870812018-09-13 09:31:40 +0100361 { "roi_align_layer", "roi_align_layer.cl" },
Michele Di Giorgio578a9fc2019-08-23 11:49:04 +0100362 { "roi_align_layer_quantized", "roi_align_layer_quantized.cl" },
SiCong Li3e363692017-07-04 15:02:10 +0100363 { "roi_pooling_layer", "roi_pooling_layer.cl" },
Michalis Spyrou46da23f2018-04-10 13:41:30 +0100364 { "scale_nearest_neighbour_nchw", "scale.cl" },
365 { "scale_nearest_neighbour_nhwc", "scale.cl" },
366 { "scale_bilinear_nchw", "scale.cl" },
367 { "scale_bilinear_nhwc", "scale.cl" },
Michalis Spyrou17220e22018-09-12 13:35:38 +0100368 { "scale_bilinear_quantized_nchw", "scale_quantized.cl" },
369 { "scale_bilinear_quantized_nhwc", "scale_quantized.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100370 { "scharr3x3", "scharr_filter.cl" },
Georgios Pinitasaaa27182018-11-21 16:32:15 +0000371 { "select_same_rank", "select.cl" },
372 { "select_different_rank_2", "select.cl" },
373 { "select_different_rank_n", "select.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100374 { "sobel3x3", "sobel_filter.cl" },
375 { "sobel_separable5x1", "sobel_filter.cl" },
376 { "sobel_separable1x5", "sobel_filter.cl" },
377 { "sobel_separable7x1", "sobel_filter.cl" },
378 { "sobel_separable1x7", "sobel_filter.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100379 { "softmax_layer_norm", "softmax_layer.cl" },
Chunosovf450caa2017-11-08 16:09:35 +0700380 { "softmax_layer_norm_quantized", "softmax_layer_quantized.cl" },
Giorgio Arena4402cb92018-02-15 13:37:40 +0000381 { "softmax_layer_max_shift_exp_sum_quantized_serial", "softmax_layer_quantized.cl" },
382 { "softmax_layer_max_shift_exp_sum_quantized_parallel", "softmax_layer_quantized.cl" },
Chunosovd6afedc2017-11-06 22:09:45 +0700383 { "softmax_layer_max_shift_exp_sum_serial", "softmax_layer.cl" },
Michalis Spyrou13a51e12018-09-18 13:09:30 +0100384 { "space_to_batch_nchw", "space_to_batch.cl" },
385 { "space_to_batch_static_nchw", "space_to_batch.cl" },
386 { "space_to_batch_nhwc", "space_to_batch.cl" },
387 { "space_to_batch_static_nhwc", "space_to_batch.cl" },
Michalis Spyroud69b3b22019-05-29 17:03:38 +0100388 { "space_to_depth_nchw", "space_to_depth.cl" },
389 { "space_to_depth_nhwc", "space_to_depth.cl" },
Chunosovd6afedc2017-11-06 22:09:45 +0700390 { "softmax_layer_max_shift_exp_sum_parallel", "softmax_layer.cl" },
Gian Marco Iodice8aa985e2018-11-27 15:58:08 +0000391 { "stack_layer", "stack_layer.cl" },
Georgios Pinitasc1a72452018-08-24 11:25:32 +0100392 { "strided_slice", "slice_ops.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100393 { "suppress_non_maximum", "canny.cl" },
394 { "tablelookup_U8", "tablelookup.cl" },
395 { "tablelookup_S16", "tablelookup.cl" },
396 { "threshold_binary", "threshold.cl" },
397 { "threshold_range", "threshold.cl" },
giuros013175fcf2018-11-21 09:59:17 +0000398 { "tile", "tile.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100399 { "transpose", "transpose.cl" },
400 { "UYVY422_to_IYUV_bt709", "color_convert.cl" },
401 { "UYVY422_to_NV12_bt709", "color_convert.cl" },
402 { "UYVY422_to_RGB888_bt709", "color_convert.cl" },
403 { "UYVY422_to_RGBA8888_bt709", "color_convert.cl" },
Michalis Spyrouceb889e2018-09-17 18:24:41 +0100404 { "upsample_layer_nchw", "upsample_layer.cl" },
405 { "upsample_layer_nhwc", "upsample_layer.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100406 { "warp_affine_nearest_neighbour", "warp_affine.cl" },
407 { "warp_affine_bilinear", "warp_affine.cl" },
408 { "warp_perspective_nearest_neighbour", "warp_perspective.cl" },
409 { "warp_perspective_bilinear", "warp_perspective.cl" },
Giorgio Arenaa50e5e02018-07-02 13:42:23 +0100410 { "winograd_filter_transform_2x2_3x3_nchw", "winograd_filter_transform.cl" },
411 { "winograd_filter_transform_2x1_3x1_nchw", "winograd_filter_transform.cl" },
412 { "winograd_filter_transform_1x2_1x3_nchw", "winograd_filter_transform.cl" },
413 { "winograd_filter_transform_4x4_3x3_nchw", "winograd_filter_transform.cl" },
414 { "winograd_filter_transform_4x1_3x1_nchw", "winograd_filter_transform.cl" },
415 { "winograd_filter_transform_1x4_1x3_nchw", "winograd_filter_transform.cl" },
416 { "winograd_filter_transform_4x4_5x5_nchw", "winograd_filter_transform.cl" },
Gian Marco Iodice876be2a2018-07-03 12:22:09 +0100417 { "winograd_filter_transform_4x1_5x1_nchw", "winograd_filter_transform.cl" },
418 { "winograd_filter_transform_1x4_1x5_nchw", "winograd_filter_transform.cl" },
Giorgio Arena149fdf32018-07-04 17:03:33 +0100419 { "winograd_filter_transform_4x1_3x1_nhwc", "winograd_filter_transform.cl" },
420 { "winograd_filter_transform_1x4_1x3_nhwc", "winograd_filter_transform.cl" },
Giorgio Arenaa50e5e02018-07-02 13:42:23 +0100421 { "winograd_filter_transform_4x4_3x3_nhwc", "winograd_filter_transform.cl" },
422 { "winograd_filter_transform_4x4_5x5_nhwc", "winograd_filter_transform.cl" },
Gian Marco Iodiced28b7512018-07-06 12:59:28 +0100423 { "winograd_filter_transform_4x1_5x1_nhwc", "winograd_filter_transform.cl" },
424 { "winograd_filter_transform_1x4_1x5_nhwc", "winograd_filter_transform.cl" },
Michele Di Giorgio881c6842019-02-27 14:26:51 +0000425 { "winograd_filter_transform_2x2_7x7_nhwc", "winograd_filter_transform.cl" },
426 { "winograd_filter_transform_2x1_7x1_nhwc", "winograd_filter_transform.cl" },
427 { "winograd_filter_transform_1x2_1x7_nhwc", "winograd_filter_transform.cl" },
Giorgio Arenaa50e5e02018-07-02 13:42:23 +0100428 { "winograd_input_transform_2x2_3x3_stepz1_nchw", "winograd_input_transform.cl" },
429 { "winograd_input_transform_2x2_3x3_stepz2_nchw", "winograd_input_transform.cl" },
430 { "winograd_input_transform_2x1_3x1_stepz1_nchw", "winograd_input_transform.cl" },
431 { "winograd_input_transform_2x1_3x1_stepz2_nchw", "winograd_input_transform.cl" },
432 { "winograd_input_transform_1x2_1x3_stepz1_nchw", "winograd_input_transform.cl" },
433 { "winograd_input_transform_1x2_1x3_stepz2_nchw", "winograd_input_transform.cl" },
434 { "winograd_input_transform_4x4_3x3_stepz1_nchw", "winograd_input_transform.cl" },
435 { "winograd_input_transform_4x1_3x1_stepz1_nchw", "winograd_input_transform.cl" },
436 { "winograd_input_transform_1x4_1x3_stepz1_nchw", "winograd_input_transform.cl" },
437 { "winograd_input_transform_4x4_5x5_stepz1_nchw", "winograd_input_transform.cl" },
Gian Marco Iodice876be2a2018-07-03 12:22:09 +0100438 { "winograd_input_transform_4x1_5x1_stepz1_nchw", "winograd_input_transform.cl" },
439 { "winograd_input_transform_1x4_1x5_stepz1_nchw", "winograd_input_transform.cl" },
Giorgio Arena149fdf32018-07-04 17:03:33 +0100440 { "winograd_input_transform_4x1_3x1_stepz1_nhwc", "winograd_input_transform.cl" },
441 { "winograd_input_transform_1x4_1x3_stepz1_nhwc", "winograd_input_transform.cl" },
Giorgio Arenaa50e5e02018-07-02 13:42:23 +0100442 { "winograd_input_transform_4x4_3x3_stepz1_nhwc", "winograd_input_transform.cl" },
443 { "winograd_input_transform_4x4_5x5_stepz1_nhwc", "winograd_input_transform.cl" },
Gian Marco Iodiced28b7512018-07-06 12:59:28 +0100444 { "winograd_input_transform_4x1_5x1_stepz1_nhwc", "winograd_input_transform.cl" },
445 { "winograd_input_transform_1x4_1x5_stepz1_nhwc", "winograd_input_transform.cl" },
Michele Di Giorgiof955d512019-02-27 14:26:51 +0000446 { "winograd_input_transform_2x2_7x7_stepz1_nhwc", "winograd_input_transform.cl" },
447 { "winograd_input_transform_2x1_7x1_stepz1_nhwc", "winograd_input_transform.cl" },
448 { "winograd_input_transform_1x2_1x7_stepz1_nhwc", "winograd_input_transform.cl" },
Giorgio Arenaa50e5e02018-07-02 13:42:23 +0100449 { "winograd_output_transform_2x2_3x3_nchw", "winograd_output_transform.cl" },
450 { "winograd_output_transform_2x1_3x1_nchw", "winograd_output_transform.cl" },
451 { "winograd_output_transform_1x2_1x3_nchw", "winograd_output_transform.cl" },
452 { "winograd_output_transform_4x4_3x3_nchw", "winograd_output_transform.cl" },
453 { "winograd_output_transform_4x1_3x1_nchw", "winograd_output_transform.cl" },
454 { "winograd_output_transform_1x4_1x3_nchw", "winograd_output_transform.cl" },
455 { "winograd_output_transform_4x4_5x5_nchw", "winograd_output_transform.cl" },
Gian Marco Iodice876be2a2018-07-03 12:22:09 +0100456 { "winograd_output_transform_4x1_5x1_nchw", "winograd_output_transform.cl" },
457 { "winograd_output_transform_1x4_1x5_nchw", "winograd_output_transform.cl" },
Giorgio Arena149fdf32018-07-04 17:03:33 +0100458 { "winograd_output_transform_4x1_3x1_nhwc", "winograd_output_transform.cl" },
459 { "winograd_output_transform_1x4_1x3_nhwc", "winograd_output_transform.cl" },
Giorgio Arenaa50e5e02018-07-02 13:42:23 +0100460 { "winograd_output_transform_4x4_3x3_nhwc", "winograd_output_transform.cl" },
461 { "winograd_output_transform_4x4_5x5_nhwc", "winograd_output_transform.cl" },
Gian Marco Iodiced28b7512018-07-06 12:59:28 +0100462 { "winograd_output_transform_4x1_5x1_nhwc", "winograd_output_transform.cl" },
463 { "winograd_output_transform_1x4_1x5_nhwc", "winograd_output_transform.cl" },
giuros013bfacb22019-04-01 12:07:02 +0100464 { "winograd_output_transform_2x2_7x7_nhwc", "winograd_output_transform.cl" },
465 { "winograd_output_transform_2x1_7x1_nhwc", "winograd_output_transform.cl" },
466 { "winograd_output_transform_1x2_1x7_nhwc", "winograd_output_transform.cl" },
Giorgio Arena73023022018-09-04 14:55:55 +0100467 { "yolo_layer_nchw", "yolo_layer.cl" },
468 { "yolo_layer_nhwc", "yolo_layer.cl" },
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100469 { "YUYV422_to_IYUV_bt709", "color_convert.cl" },
470 { "YUYV422_to_NV12_bt709", "color_convert.cl" },
471 { "YUYV422_to_RGB888_bt709", "color_convert.cl" },
472 { "YUYV422_to_RGBA8888_bt709", "color_convert.cl" },
473};
474
475const std::map<std::string, std::string> CLKernelLibrary::_program_source_map =
476{
477#ifdef EMBEDDED_KERNELS
478 {
479 "absdiff.cl",
480#include "./cl_kernels/absdiff.clembed"
481 },
482 {
483 "accumulate.cl",
484#include "./cl_kernels/accumulate.clembed"
485 },
486 {
487 "activation_layer.cl",
488#include "./cl_kernels/activation_layer.clembed"
489 },
490 {
Manuel Bottini30dbeef2019-06-26 16:23:03 +0100491 "activation_layer_quant.cl",
492#include "./cl_kernels/activation_layer_quant.clembed"
Michel Iwaniec00633802017-10-12 14:14:15 +0100493 },
494 {
Manuel Bottini7b9998d2019-10-21 17:59:07 +0100495 "arg_min_max.cl",
496#include "./cl_kernels/arg_min_max.clembed"
497 },
498 {
Michalis Spyrou6a8d3b62018-08-31 10:07:09 +0100499 "batch_to_space.cl",
500#include "./cl_kernels/batch_to_space.clembed"
501 },
502 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100503 "bitwise_op.cl",
504#include "./cl_kernels/bitwise_op.clembed"
505 },
506 {
giuros01c04a0e82018-10-03 12:44:35 +0100507 "bounding_box_transform.cl",
508#include "./cl_kernels/bounding_box_transform.clembed"
509 },
510 {
Michele Di Giorgio4aff98f2019-08-28 16:27:26 +0100511 "bounding_box_transform_quantized.cl",
512#include "./cl_kernels/bounding_box_transform_quantized.clembed"
513 },
514 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100515 "canny.cl",
516#include "./cl_kernels/canny.clembed"
517 },
518 {
519 "channel_combine.cl",
520#include "./cl_kernels/channel_combine.clembed"
521 },
522 {
523 "channel_extract.cl",
524#include "./cl_kernels/channel_extract.clembed"
525 },
526 {
Michele Di Giorgio72175632018-05-01 16:52:00 +0100527 "channel_shuffle.cl",
528#include "./cl_kernels/channel_shuffle.clembed"
529 },
530 {
Gian Marco76faef82018-01-29 12:15:32 +0000531 "col2im.cl",
532#include "./cl_kernels/col2im.clembed"
533 },
534 {
Georgios Pinitas7900a9e2018-11-23 11:44:58 +0000535 "comparisons.cl",
536#include "./cl_kernels/comparisons.clembed"
537 },
538 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100539 "concatenate.cl",
540#include "./cl_kernels/concatenate.clembed"
541 },
542 {
543 "color_convert.cl",
544#include "./cl_kernels/color_convert.clembed"
545 },
546 {
Giorgio Arena657bdb32018-04-26 18:52:01 +0100547 "convert_fc_weights.cl",
548#include "./cl_kernels/convert_fc_weights.clembed"
549 },
550 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100551 "convolution3x3.cl",
552#include "./cl_kernels/convolution3x3.clembed"
553 },
554 {
555 "convolution5x5.cl",
556#include "./cl_kernels/convolution5x5.clembed"
557 },
558 {
559 "convolution7x7.cl",
560#include "./cl_kernels/convolution7x7.clembed"
561 },
562 {
563 "convolution9x9.cl",
564#include "./cl_kernels/convolution9x9.clembed"
565 },
566 {
567 "convolution_layer.cl",
568#include "./cl_kernels/convolution_layer.clembed"
569 },
570 {
571 "convolution_rectangle.cl",
572#include "./cl_kernels/convolution_rectangle.clembed"
573 },
574 {
Michalis Spyrou5c8e05c2018-03-22 11:56:01 +0000575 "copy_tensor.cl",
576#include "./cl_kernels/copy_tensor.clembed"
577 },
578 {
George Wort894066d2019-02-15 15:12:52 +0000579 "crop_tensor.cl",
580#include "./cl_kernels/crop_tensor.clembed"
581 },
582 {
Michalis Spyrouceb889e2018-09-17 18:24:41 +0100583 "upsample_layer.cl",
584#include "./cl_kernels/upsample_layer.clembed"
585 },
586 {
Michalis Spyrou780db4e2017-11-23 09:49:51 +0000587 "deconvolution_layer.cl",
588#include "./cl_kernels/deconvolution_layer.clembed"
589 },
590 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100591 "depth_convert.cl",
592#include "./cl_kernels/depth_convert.clembed"
593 },
594 {
Michalis Spyrou649962c2019-05-22 11:11:55 +0100595 "depth_to_space.cl",
596#include "./cl_kernels/depth_to_space.clembed"
597 },
598 {
Giorgio Arena93a690e2017-08-01 16:09:33 +0100599 "depthwise_convolution.cl",
600#include "./cl_kernels/depthwise_convolution.clembed"
601 },
602 {
Dmitry Savenkod7295b72017-11-20 22:00:08 +0700603 "depthwise_convolution_quantized.cl",
604#include "./cl_kernels/depthwise_convolution_quantized.clembed"
605 },
606 {
Michele Di Giorgio56dd7262017-07-27 09:53:49 +0100607 "dequantization_layer.cl",
608#include "./cl_kernels/dequantization_layer.clembed"
609 },
610 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100611 "derivative.cl",
612#include "./cl_kernels/derivative.clembed"
613 },
614 {
615 "dilate.cl",
616#include "./cl_kernels/dilate.clembed"
617 },
618 {
SiCong Lic51b72f2017-07-28 14:46:20 +0100619 "direct_convolution1x1.cl",
620#include "./cl_kernels/direct_convolution1x1.clembed"
621 },
622 {
623 "direct_convolution3x3.cl",
624#include "./cl_kernels/direct_convolution3x3.clembed"
steniu0127b386c2017-07-18 17:37:43 +0100625 },
626 {
steniu01db006682017-08-09 16:26:22 +0100627 "direct_convolution5x5.cl",
628#include "./cl_kernels/direct_convolution5x5.clembed"
629 },
630 {
Sang-Hoon Parkab5b1a22019-10-15 09:29:13 +0100631 "direct_convolution_quantized.cl",
632#include "./cl_kernels/direct_convolution_quantized.clembed"
Chunosovd621bca2017-11-03 17:33:15 +0700633 },
634 {
Michalis Spyrou45091732019-05-13 17:41:01 +0100635 "direct_convolution9x9.cl",
636#include "./cl_kernels/direct_convolution9x9.clembed"
637 },
638 {
giuros01164a2722018-11-20 18:34:46 +0000639 "elementwise_operation.cl",
640#include "./cl_kernels/elementwise_operation.clembed"
641 },
642 {
643 "elementwise_operation_quantized.cl",
644#include "./cl_kernels/elementwise_operation_quantized.clembed"
645 },
646 {
Michalis Spyroue9362622018-11-23 17:41:37 +0000647 "elementwise_unary.cl",
648#include "./cl_kernels/elementwise_unary.clembed"
649 },
650 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100651 "erode.cl",
652#include "./cl_kernels/erode.clembed"
653 },
654 {
655 "fast_corners.cl",
656#include "./cl_kernels/fast_corners.clembed"
657 },
658 {
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000659 "fft.cl",
660#include "./cl_kernels/fft.clembed"
Gian Marco Iodice215b4ea2018-06-28 16:29:29 +0100661 },
662 {
Georgios Pinitas8be91482019-03-26 17:23:28 +0000663 "fft_digit_reverse.cl",
664#include "./cl_kernels/fft_digit_reverse.clembed"
665 },
666 {
667 "fft_scale.cl",
668#include "./cl_kernels/fft_scale.clembed"
669 },
670 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100671 "fill_border.cl",
672#include "./cl_kernels/fill_border.clembed"
673 },
674 {
Georgios Pinitas0bc78492019-03-18 20:07:37 +0000675 "flatten.cl",
676#include "./cl_kernels/flatten.clembed"
677 },
678 {
Georgios Pinitasd8e765b2017-08-02 13:44:33 +0100679 "floor.cl",
680#include "./cl_kernels/floor.clembed"
681 },
682 {
Manuel Bottini8529bd62018-11-21 11:53:04 +0000683 "gather.cl",
684#include "./cl_kernels/gather.clembed"
685 },
686 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100687 "gaussian_pyramid.cl",
688#include "./cl_kernels/gaussian_pyramid.clembed"
689 },
690 {
691 "gemm.cl",
692#include "./cl_kernels/gemm.clembed"
693 },
694 {
Gian Marco05288a22017-11-21 10:57:50 +0000695 "gemmlowp.cl",
696#include "./cl_kernels/gemmlowp.clembed"
697 },
698 {
Giorgio Arena9fe41442017-08-23 16:36:24 +0100699 "gemv.cl",
700#include "./cl_kernels/gemv.clembed"
701 },
702 {
Manuel Bottini5209be52019-02-13 16:34:56 +0000703 "generate_proposals.cl",
704#include "./cl_kernels/generate_proposals.clembed"
705 },
706 {
Michele Di Giorgio6b612f52019-09-05 12:30:22 +0100707 "generate_proposals_quantized.cl",
708#include "./cl_kernels/generate_proposals_quantized.clembed"
709 },
710 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100711 "harris_corners.cl",
712#include "./cl_kernels/harris_corners.clembed"
713 },
714 {
715 "helpers.h",
716#include "./cl_kernels/helpers.hembed"
717 },
718 {
Chunosovd621bca2017-11-03 17:33:15 +0700719 "helpers_asymm.h",
720#include "./cl_kernels/helpers_asymm.hembed"
721 },
722 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100723 "histogram.cl",
724#include "./cl_kernels/histogram.clembed"
725 },
726 {
727 "hog.cl",
728#include "./cl_kernels/hog.clembed"
729 },
730 {
Gian Marco76faef82018-01-29 12:15:32 +0000731 "im2col.cl",
732#include "./cl_kernels/im2col.clembed"
733 },
734 {
Manuel Bottini79f88e62019-09-18 15:02:53 +0100735 "instance_normalization.cl",
736#include "./cl_kernels/instance_normalization.clembed"
737 },
738 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100739 "integral_image.cl",
740#include "./cl_kernels/integral_image.clembed"
741 },
742 {
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100743 "l2_normalize.cl",
744#include "./cl_kernels/l2_normalize.clembed"
745 },
746 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100747 "magnitude_phase.cl",
748#include "./cl_kernels/magnitude_phase.clembed"
749 },
750 {
751 "mean_stddev.cl",
752#include "./cl_kernels/mean_stddev.clembed"
753 },
754 {
Michele Di Giorgio5b48ad72019-06-04 18:43:35 +0100755 "mean_stddev_normalization.cl",
756#include "./cl_kernels/mean_stddev_normalization.clembed"
757 },
758 {
Giuseppe Rossinid7647d42018-07-17 18:13:13 +0100759 "memset.cl",
760#include "./cl_kernels/memset.clembed"
761 },
762 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100763 "minmaxloc.cl",
764#include "./cl_kernels/minmaxloc.clembed"
765 },
766 {
Michele Di Giorgio56dd7262017-07-27 09:53:49 +0100767 "minmax_layer.cl",
768#include "./cl_kernels/minmax_layer.clembed"
769 },
770 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100771 "non_linear_filter3x3.cl",
772#include "./cl_kernels/non_linear_filter3x3.clembed"
773 },
774 {
775 "non_linear_filter5x5.cl",
776#include "./cl_kernels/non_linear_filter5x5.clembed"
777 },
778 {
779 "non_linear_filter_helpers.h",
780#include "./cl_kernels/non_linear_filter_helpers.hembed"
781 },
782 {
783 "nonmax.cl",
784#include "./cl_kernels/nonmax.clembed"
785 },
786 {
787 "normalization_layer.cl",
788#include "./cl_kernels/normalization_layer.clembed"
789 },
790 {
Michele Di Giorgiob57be0d2018-08-31 16:26:25 +0100791 "normalize_planar_yuv_layer.cl",
792#include "./cl_kernels/normalize_planar_yuv_layer.clembed"
793 },
794 {
Michele Di Giorgiod63dfa22018-09-12 10:18:54 +0100795 "normalize_planar_yuv_layer_quantized.cl",
796#include "./cl_kernels/normalize_planar_yuv_layer_quantized.clembed"
797 },
798 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100799 "batchnormalization_layer.cl",
800#include "./cl_kernels/batchnormalization_layer.clembed"
801 },
802 {
803 "optical_flow_pyramid_lk.cl",
804#include "./cl_kernels/optical_flow_pyramid_lk.clembed"
805 },
806 {
Giorgio Arena205eed82019-08-14 10:13:50 +0100807 "pad_layer.cl",
808#include "./cl_kernels/pad_layer.clembed"
809 },
810 {
Michalis Spyrou5237e012018-01-17 09:40:27 +0000811 "permute.cl",
812#include "./cl_kernels/permute.clembed"
813 },
814 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100815 "pixelwise_mul_float.cl",
816#include "./cl_kernels/pixelwise_mul_float.clembed"
817 },
818 {
819 "pixelwise_mul_int.cl",
820#include "./cl_kernels/pixelwise_mul_int.clembed"
821 },
822 {
823 "pooling_layer.cl",
824#include "./cl_kernels/pooling_layer.clembed"
825 },
826 {
Anton Lokhmotovaf6204c2017-11-08 09:34:19 +0000827 "pooling_layer_quantized.cl",
828#include "./cl_kernels/pooling_layer_quantized.clembed"
829 },
830 {
Michalis Spyrou6c7c38e2018-08-29 16:28:11 +0100831 "prior_box_layer.cl",
832#include "./cl_kernels/prior_box_layer.clembed"
833 },
834 {
Sheri Zhangb18252d2020-04-07 11:04:57 +0100835 "qlstm_layer_normalization.cl",
836#include "./cl_kernels/qlstm_layer_normalization.clembed"
837 },
838 {
Michele Di Giorgio56dd7262017-07-27 09:53:49 +0100839 "quantization_layer.cl",
840#include "./cl_kernels/quantization_layer.clembed"
841 },
842 {
Vidhya Sudhan Loganathan5e96be72018-12-18 14:17:00 +0000843 "range.cl",
844#include "./cl_kernels/range.clembed"
845 },
846 {
Michalis Spyrou04f089c2017-08-08 17:42:38 +0100847 "reduction_operation.cl",
848#include "./cl_kernels/reduction_operation.clembed"
849 },
850 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100851 "remap.cl",
852#include "./cl_kernels/remap.clembed"
853 },
854 {
Gian Marco Iodice477531c2018-08-21 17:53:38 +0100855 "reorg_layer.cl",
856#include "./cl_kernels/reorg_layer.clembed"
857 },
858 {
Georgios Pinitas5ee66ea2017-09-07 17:29:16 +0100859 "reshape_layer.cl",
860#include "./cl_kernels/reshape_layer.clembed"
861 },
862 {
Michele Di Giorgio5daeffd2018-11-26 10:01:15 +0000863 "reverse.cl",
864#include "./cl_kernels/reverse.clembed"
865 },
866 {
giuros0118870812018-09-13 09:31:40 +0100867 "roi_align_layer.cl",
868#include "./cl_kernels/roi_align_layer.clembed"
869 },
870 {
Michele Di Giorgio578a9fc2019-08-23 11:49:04 +0100871 "roi_align_layer_quantized.cl",
872#include "./cl_kernels/roi_align_layer_quantized.clembed"
873 },
874 {
SiCong Li3e363692017-07-04 15:02:10 +0100875 "roi_pooling_layer.cl",
876#include "./cl_kernels/roi_pooling_layer.clembed"
877 },
878 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100879 "scale.cl",
880#include "./cl_kernels/scale.clembed"
881 },
882 {
Michalis Spyrou17220e22018-09-12 13:35:38 +0100883 "scale_quantized.cl",
884#include "./cl_kernels/scale_quantized.clembed"
885 },
886 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100887 "scharr_filter.cl",
888#include "./cl_kernels/scharr_filter.clembed"
889 },
890 {
Georgios Pinitasaaa27182018-11-21 16:32:15 +0000891 "select.cl",
892#include "./cl_kernels/select.clembed"
893 },
894 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100895 "sobel_filter.cl",
896#include "./cl_kernels/sobel_filter.clembed"
897 },
898 {
899 "softmax_layer.cl",
900#include "./cl_kernels/softmax_layer.clembed"
901 },
902 {
Chunosovf450caa2017-11-08 16:09:35 +0700903 "softmax_layer_quantized.cl",
904#include "./cl_kernels/softmax_layer_quantized.clembed"
905 },
906 {
Georgios Pinitasc1a72452018-08-24 11:25:32 +0100907 "slice_ops.cl",
908#include "./cl_kernels/slice_ops.clembed"
Georgios Pinitas77589b52018-08-21 14:41:35 +0100909 },
910 {
Michalis Spyrou16934a52018-08-21 18:03:58 +0100911 "space_to_batch.cl",
912#include "./cl_kernels/space_to_batch.clembed"
913 },
914 {
Michalis Spyroud69b3b22019-05-29 17:03:38 +0100915 "space_to_depth.cl",
916#include "./cl_kernels/space_to_depth.clembed"
917 },
918 {
Gian Marco Iodice8aa985e2018-11-27 15:58:08 +0000919 "stack_layer.cl",
920#include "./cl_kernels/stack_layer.clembed"
921 },
922 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100923 "tablelookup.cl",
924#include "./cl_kernels/tablelookup.clembed"
925 },
926 {
927 "threshold.cl",
928#include "./cl_kernels/threshold.clembed"
929 },
930 {
giuros013175fcf2018-11-21 09:59:17 +0000931 "tile.cl",
932#include "./cl_kernels/tile.clembed"
933 },
934 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100935 "transpose.cl",
936#include "./cl_kernels/transpose.clembed"
937 },
938 {
939 "types.h",
940#include "./cl_kernels/types.hembed"
941 },
942 {
Gian Marco Iodice4d81d752020-07-14 15:05:31 +0100943 "unpooling_layer.cl",
944#include "./cl_kernels/unpooling_layer.clembed"
945 },
946 {
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100947 "warp_affine.cl",
948#include "./cl_kernels/warp_affine.clembed"
949 },
950 {
951 "warp_helpers.h",
952#include "./cl_kernels/warp_helpers.hembed"
953 },
954 {
955 "warp_perspective.cl",
956#include "./cl_kernels/warp_perspective.clembed"
Michalis Spyroud7e82812017-06-20 15:00:14 +0100957 },
Giorgio Arena1f9ca1d2018-03-01 11:13:45 +0000958 {
Giorgio Arenaa50e5e02018-07-02 13:42:23 +0100959 "winograd_filter_transform.cl",
960#include "./cl_kernels/winograd_filter_transform.clembed"
961 },
962 {
963 "winograd_input_transform.cl",
964#include "./cl_kernels/winograd_input_transform.clembed"
965 },
966 {
967 "winograd_output_transform.cl",
968#include "./cl_kernels/winograd_output_transform.clembed"
Giorgio Arena1f9ca1d2018-03-01 11:13:45 +0000969 },
Giorgio Arena73023022018-09-04 14:55:55 +0100970 {
971 "yolo_layer.cl",
972#include "./cl_kernels/yolo_layer.clembed"
973 },
Anthony Barbierac69aa12017-07-03 17:39:37 +0100974#endif /* EMBEDDED_KERNELS */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100975};
976
977CLKernelLibrary::CLKernelLibrary()
Michalis Spyrou11d49182020-03-26 10:31:32 +0000978 : _compile_context(), _kernel_path()
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100979{
Anthony Barbierecb1c622018-04-17 11:45:10 +0100980 opencl_is_available(); // Make sure the OpenCL symbols are initialised *before* the CLKernelLibrary is built
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100981}
982
983CLKernelLibrary &CLKernelLibrary::get()
984{
985 static CLKernelLibrary _kernel_library;
986 return _kernel_library;
987}
988
Michalis Spyrou11d49182020-03-26 10:31:32 +0000989Kernel CLKernelLibrary::create_kernel(const std::string &kernel_name, const std::set<std::string> &build_options_set) const
990{
991 const std::string program_name = get_program_name(kernel_name);
992 auto program = get_program(program_name);
993
994 return _compile_context.create_kernel(kernel_name, program_name, program.first, _kernel_path, build_options_set, program.second);
995}
996
997std::string CLKernelLibrary::get_program_name(const std::string &kernel_name) const
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100998{
999 // Find which program contains the kernel
1000 auto kernel_program_it = _kernel_program_map.find(kernel_name);
1001
1002 if(_kernel_program_map.end() == kernel_program_it)
1003 {
Michalis Spyrou7c60c992019-10-10 14:33:47 +01001004 ARM_COMPUTE_ERROR_VAR("Kernel %s not found in the CLKernelLibrary", kernel_name.c_str());
Anthony Barbier6ff3b192017-09-04 18:44:23 +01001005 }
steniu0134702472017-07-11 09:22:58 +01001006
Michalis Spyrou11d49182020-03-26 10:31:32 +00001007 const std::string program_name = kernel_program_it->second;
Georgios Pinitasdf473ea2018-05-31 18:53:52 +01001008
Michalis Spyrou11d49182020-03-26 10:31:32 +00001009 return program_name;
Anthony Barbier6ff3b192017-09-04 18:44:23 +01001010}
1011
Pablo Tellodb8485a2019-09-24 11:03:47 +01001012void CLKernelLibrary::init(std::string kernel_path, cl::Context context, cl::Device device)
1013{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001014 _compile_context = CLCompileContext(context, device);
1015 _kernel_path = kernel_path;
Pablo Tellodb8485a2019-09-24 11:03:47 +01001016}
1017
1018void CLKernelLibrary::set_kernel_path(const std::string &kernel_path)
1019{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001020 _kernel_path = std::move(kernel_path);
Pablo Tellodb8485a2019-09-24 11:03:47 +01001021}
1022
1023cl::Context &CLKernelLibrary::context()
1024{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001025 return _compile_context.context();
Pablo Tellodb8485a2019-09-24 11:03:47 +01001026}
1027
Michalis Spyrou11d49182020-03-26 10:31:32 +00001028const cl::Device &CLKernelLibrary::get_device()
Pablo Tellodb8485a2019-09-24 11:03:47 +01001029{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001030 return _compile_context.get_device();
Pablo Tellodb8485a2019-09-24 11:03:47 +01001031}
1032
1033void CLKernelLibrary::set_device(cl::Device device)
1034{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001035 _compile_context.set_device(device);
1036}
1037
1038void CLKernelLibrary::set_context(cl::Context context)
1039{
1040 _compile_context.set_context(context);
Pablo Tellodb8485a2019-09-24 11:03:47 +01001041}
1042
1043std::string CLKernelLibrary::get_kernel_path()
1044{
1045 return _kernel_path;
1046}
1047
1048void CLKernelLibrary::clear_programs_cache()
1049{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001050 _compile_context.clear_programs_cache();
Pablo Tellodb8485a2019-09-24 11:03:47 +01001051}
1052
1053const std::map<std::string, cl::Program> &CLKernelLibrary::get_built_programs() const
1054{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001055 return _compile_context.get_built_programs();
Pablo Tellodb8485a2019-09-24 11:03:47 +01001056}
1057
giuros0146a49a02019-04-01 13:50:22 +01001058void CLKernelLibrary::add_built_program(const std::string &built_program_name, const cl::Program &program)
Anthony Barbier7da55aa2018-04-13 16:58:43 +01001059{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001060 _compile_context.add_built_program(built_program_name, program);
Anthony Barbier7da55aa2018-04-13 16:58:43 +01001061}
1062
Vidhya Sudhan Loganathanf1f49062018-05-25 13:21:26 +01001063bool CLKernelLibrary::fp16_supported() const
1064{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001065 return _compile_context.fp16_supported();
Vidhya Sudhan Loganathanf1f49062018-05-25 13:21:26 +01001066}
1067
Vidhya Sudhan Loganathan76c85642018-05-25 13:53:02 +01001068bool CLKernelLibrary::int64_base_atomics_supported() const
1069{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001070 return _compile_context.int64_base_atomics_supported();
Vidhya Sudhan Loganathan76c85642018-05-25 13:53:02 +01001071}
1072
Michalis Spyrou11d49182020-03-26 10:31:32 +00001073std::pair<std::string, bool> CLKernelLibrary::get_program(const std::string &program_name) const
Anthony Barbier6ff3b192017-09-04 18:44:23 +01001074{
Anthony Barbier6ff3b192017-09-04 18:44:23 +01001075#ifdef EMBEDDED_KERNELS
1076 const auto program_source_it = _program_source_map.find(program_name);
1077
Michalis Spyroud7e82812017-06-20 15:00:14 +01001078 if(program_source_it == _program_source_map.end())
1079 {
Michalis Spyrou7c60c992019-10-10 14:33:47 +01001080 ARM_COMPUTE_ERROR_VAR("Embedded program for %s does not exist.", program_name.c_str());
Michalis Spyroud7e82812017-06-20 15:00:14 +01001081 }
1082
Michalis Spyrou11d49182020-03-26 10:31:32 +00001083 return std::make_pair(program_source_it->second, false);
1084#else /* EMBEDDED_KERNELS */
1085 // Check for binary
1086 std::string source_name = _kernel_path + program_name;
1087 std::string binary_name = source_name + "bin";
1088 std::string program_source{};
1089 bool is_binary = false;
1090
1091 if(std::ifstream(binary_name).is_open())
1092 {
1093 program_source = read_file(binary_name, true);
1094 is_binary = true;
1095 }
1096 else if(std::ifstream(source_name).is_open())
1097 {
1098 program_source = read_file(source_name, false);
1099 }
1100 else
1101 {
1102 ARM_COMPUTE_ERROR_VAR("Kernel file %s does not exist.", source_name.c_str());
1103 }
1104
1105 return std::make_pair(program_source, is_binary);
1106#endif /* EMBEDDED_KERNELS */
Michalis Spyroud7e82812017-06-20 15:00:14 +01001107}
steniu015f910722017-08-23 10:15:22 +01001108
Abel Bernabeu5a6e0532017-09-28 09:53:45 +01001109size_t CLKernelLibrary::max_local_workgroup_size(const cl::Kernel &kernel) const
steniu015f910722017-08-23 10:15:22 +01001110{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001111 return _compile_context.max_local_workgroup_size(kernel);
steniu015f910722017-08-23 10:15:22 +01001112}
1113
Abel Bernabeu5a6e0532017-09-28 09:53:45 +01001114cl::NDRange CLKernelLibrary::default_ndrange() const
steniu015f910722017-08-23 10:15:22 +01001115{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001116 return _compile_context.default_ndrange();
steniu015f910722017-08-23 10:15:22 +01001117}
Anthony Barbier847864d2018-03-07 11:35:53 +00001118
1119std::string CLKernelLibrary::get_device_version()
1120{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001121 return _compile_context.get_device_version();
Anthony Barbier847864d2018-03-07 11:35:53 +00001122}
Giorgio Arena5d42b462019-07-26 15:54:20 +01001123
1124cl_uint CLKernelLibrary::get_num_compute_units()
1125{
Michalis Spyrou11d49182020-03-26 10:31:32 +00001126 return _compile_context.get_num_compute_units();
1127}
1128
1129CLCompileContext &CLKernelLibrary::get_compile_context()
1130{
1131 return _compile_context;
Michele Di Giorgio578a9fc2019-08-23 11:49:04 +01001132}