From 17220e2eb49e75b85f2b802489a44b8019997c25 Mon Sep 17 00:00:00 2001 From: Michalis Spyrou Date: Wed, 12 Sep 2018 13:35:38 +0100 Subject: COMPMID-1507 Add support for QASYMM8 in CLScaleKernel Change-Id: I4a32e47e6d9152633668cf0e14db88fc8c26f7ea Reviewed-on: https://eu-gerrit-1.euhpc.arm.com/148167 Tested-by: bsgcomp Reviewed-by: Michele DiGiorgio --- src/core/CL/CLKernelLibrary.cpp | 6 + src/core/CL/cl_kernels/scale_quantized.cl | 169 ++++++++++++++++++++++++ src/core/CL/cl_kernels/warp_helpers_quantized.h | 138 +++++++++++++++++++ src/core/CL/kernels/CLScaleKernel.cpp | 15 ++- 4 files changed, 325 insertions(+), 3 deletions(-) create mode 100644 src/core/CL/cl_kernels/scale_quantized.cl create mode 100644 src/core/CL/cl_kernels/warp_helpers_quantized.h (limited to 'src/core/CL') diff --git a/src/core/CL/CLKernelLibrary.cpp b/src/core/CL/CLKernelLibrary.cpp index 8f5e81bae9..392fbfefb0 100644 --- a/src/core/CL/CLKernelLibrary.cpp +++ b/src/core/CL/CLKernelLibrary.cpp @@ -356,6 +356,8 @@ const std::map CLKernelLibrary::_kernel_program_map = { "scale_nearest_neighbour_nhwc", "scale.cl" }, { "scale_bilinear_nchw", "scale.cl" }, { "scale_bilinear_nhwc", "scale.cl" }, + { "scale_bilinear_quantized_nchw", "scale_quantized.cl" }, + { "scale_bilinear_quantized_nhwc", "scale_quantized.cl" }, { "scharr3x3", "scharr_filter.cl" }, { "sobel3x3", "sobel_filter.cl" }, { "sobel_separable5x1", "sobel_filter.cl" }, @@ -743,6 +745,10 @@ const std::map CLKernelLibrary::_program_source_map = { "scale.cl", #include "./cl_kernels/scale.clembed" + }, + { + "scale_quantized.cl", +#include "./cl_kernels/scale_quantized.clembed" }, { "scharr_filter.cl", diff --git a/src/core/CL/cl_kernels/scale_quantized.cl b/src/core/CL/cl_kernels/scale_quantized.cl new file mode 100644 index 0000000000..3211e7efa1 --- /dev/null +++ b/src/core/CL/cl_kernels/scale_quantized.cl @@ -0,0 +1,169 @@ +/* + * Copyright (c) 2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers_asymm.h" +#include "warp_helpers_quantized.h" + +/** Transforms four 2D coordinates. This is used to map the output coordinates to the input coordinates. + * + * @param[in] coord 2D coordinates to transform. + * @param[in] scale input/output scale ratio + * + * @return a float8 containing 4 2D transformed values in the input image. + */ +inline const float8 transform_bilinear_quantized(const float2 coord, const float2 scale) +{ + const float4 in_x_coords = (float4)(coord.s0, 1 + coord.s0, 2 + coord.s0, 3 + coord.s0); +#ifdef SAMPLING_POLICY_TOP_LEFT + const float4 new_x = in_x_coords * (float4)(scale.s0); + const float4 new_y = (float4)(coord.s1 * scale.s1); + return (float8)(new_x.s0, new_y.s0, new_x.s1, new_y.s1, new_x.s2, new_y.s2, new_x.s3, new_y.s3); +#elif SAMPLING_POLICY_CENTER + const float4 new_x = (in_x_coords + ((float4)(0.5f))) * (float4)(scale.s0) - (float4)(0.5f); + const float4 new_y = (float4)((coord.s1 + 0.5f) * scale.s1 - 0.5f); + return (float8)(new_x.s0, new_y.s0, new_x.s1, new_y.s1, new_x.s2, new_y.s2, new_x.s3, new_y.s3); +#else /* SAMPLING_POLICY */ +#error("Unsupported sampling policy"); +#endif /* SAMPLING_POLICY */ +} + +/** Performs an affine transformation on an image interpolating with the BILINEAR method. + * + * @note Sampling policy to used is passed as -DSAMPLING_POLICY_(TYPE) e.g. -DSAMPLING_POLICY_TOP_LEFT + * @note Scale value for QASYMM8 data type to used is passed as -DSCALE= e.g. -DSCALE=0.5 + * @note Offset value for QASYMM8 data type to used is passed as -DOFFSET= e.g. -DOFFSET=1 + * + * @param[in] in_ptr Pointer to the source image. Supported data types: QASYMM8. + * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] in_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] in_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] in_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] in_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] out_ptr Pointer to the destination image. Supported data types: U8, S16. (Must be the same as the input) + * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] out_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] out_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[in] input_width Input image width + * @param[in] input_height Input image height + * @param[in] scale_x The scale factor along x dimension + * @param[in] scale_y The scale factor along y dimension + */ +__kernel void scale_bilinear_quantized_nchw( + IMAGE_DECLARATION(in), + IMAGE_DECLARATION(out), + const float input_width, + const float input_height, + const float scale_x, + const float scale_y) +{ + Image in = CONVERT_TO_IMAGE_STRUCT_NO_STEP(in); + Image out = CONVERT_TO_IMAGE_STRUCT(out); + const float2 r = (float2)(scale_x, scale_y); + const float8 tc = transform_bilinear_quantized(get_current_coords_quantized(), r); + vstore4(bilinear_interpolate_with_border_quantized(&in, tc, input_width, input_height, BORDER_SIZE, SCALE, OFFSET), 0, (__global DATA_TYPE *)out.ptr); +} + +/** Performs scale on an image interpolating with the BILINEAR method. (NHWC) + * + * @note Sampling policy to be used is passed as -DSAMPLING_POLICY_(TYPE) e.g. -DSAMPLING_POLICY_TOP_LEFT + * @note Scale value for QASYMM8 data type to used is passed as -DSCALE= e.g. -DSCALE=0.5 + * @note Offset value for QASYMM8 data type to used is passed as -DOFFSET= e.g. -DOFFSET=1 + * @note If border mode replicate is used, is should be passed as -DBORDER_MODE_REPLICATE + * + * @param[in] in_ptr Pointer to the source image. Supported data types: QASYMM8. + * @param[in] in_stride_x Stride of the source image in X dimension (in bytes) + * @param[in] in_step_x src_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] in_stride_y Stride of the source image in Y dimension (in bytes) + * @param[in] in_step_y src_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] in_stride_z Stride of the source image in Z dimension (in bytes) + * @param[in] in_step_z src_stride_z * number of elements along Z processed per workitem(in bytes) + * @param[in] in_offset_first_element_in_bytes The offset of the first element in the source image + * @param[out] out_ptr Pointer to the destination image. Supported data types: same as @p in_ptr + * @param[in] out_stride_x Stride of the destination image in X dimension (in bytes) + * @param[in] out_step_x dst_stride_x * number of elements along X processed per workitem(in bytes) + * @param[in] out_stride_y Stride of the destination image in Y dimension (in bytes) + * @param[in] out_step_y dst_stride_y * number of elements along Y processed per workitem(in bytes) + * @param[in] out_stride_z Stride of the destination image in Z dimension (in bytes) + * @param[in] out_step_z dst_stride_y * number of elements along Z processed per workitem(in bytes) + * @param[in] out_offset_first_element_in_bytes The offset of the first element in the destination image + * @param[in] input_width Input image width + * @param[in] input_height Input image height + * @param[in] scale_x The scale factor along x dimension + * @param[in] scale_y The scale factor along y dimension + */ +__kernel void scale_bilinear_quantized_nhwc( + TENSOR3D_DECLARATION(in), + TENSOR3D_DECLARATION(out), + const float input_width, + const float input_height, + const float scale_x, + const float scale_y) +{ + Tensor3D in = CONVERT_TO_TENSOR3D_STRUCT_NO_STEP(in); + Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(out); + +#ifdef SAMPLING_POLICY_TOP_LEFT + const float new_x = get_global_id(1) * scale_x; + const float new_y = get_global_id(2) * scale_y; +#elif SAMPLING_POLICY_CENTER + const float new_x = (get_global_id(1) + 0.5f) * scale_x - 0.5f; + const float new_y = (get_global_id(2) + 0.5f) * scale_y - 0.5f; +#else /* SAMPLING_POLICY */ +#error("Unsupported sampling policy"); +#endif /* SAMPLING_POLICY */ + + const float new_xf = floor(new_x); + const float new_yf = floor(new_y); + float clamped_x = clamp(new_xf, 0.0f, input_width - 1); + float clamped_x1 = clamp(new_xf + 1, 0.0f, input_width - 1); + float clamped_x_ = clamped_x; + float clamped_x1_ = clamped_x1; + const float clamped_y = clamp(new_yf, 0.0f, input_height - 1); + const float clamped_y1 = clamp(new_yf + 1, 0.0f, input_height - 1); + +#ifndef BORDER_MODE_REPLICATE + clamped_x1 = select(clamped_x1, 0.0f - BORDER_SIZE, new_yf + 1 < 0.f || new_yf + 1 > input_height - 1 || new_xf + 1 < 0.f || new_xf + 1 > input_width - 1); + clamped_x_ = select(clamped_x_, 0.0f - BORDER_SIZE, new_yf + 1 > input_height - 1 || new_xf < 0.f || new_xf > input_width - 1); + clamped_x = select(clamped_x, 0.0f - BORDER_SIZE, new_yf < 0.f || new_yf > input_height - 1 || new_xf < 0.f || new_xf > input_width - 1); + clamped_x1_ = select(clamped_x1_, 0.0f - BORDER_SIZE, new_xf + 1 < 0.f || new_xf + 1 > input_width - 1 || new_yf < 0.f || new_yf > input_height - 1); +#endif /* BORDER_MODE_REPLICATE */ + + int4 ins = (int4)(*((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x), convert_int(clamped_y))), + *((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x1_), convert_int(clamped_y))), + *((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x_), convert_int(clamped_y1))), + *((__global DATA_TYPE *)tensor3D_offset(&in, get_global_id(0), convert_int(clamped_x1), convert_int(clamped_y1)))); + + const float a = new_x - new_xf; + const float b = 1.f - a; + const float a1 = new_y - new_yf; + const float b1 = 1.f - a1; + const float4 insf32 = convert_float4(ins - (int4)OFFSET) * (float4)SCALE; + + const float fr = ((insf32.s0 * b * b1) + (insf32.s1 * a * b1) + (insf32.s2 * b * a1) + (insf32.s3 * a * a1)); + + uchar res = convert_uchar_sat(convert_int_sat_rtp(fr / SCALE) + OFFSET); + + *((__global DATA_TYPE *)out.ptr) = res; +} diff --git a/src/core/CL/cl_kernels/warp_helpers_quantized.h b/src/core/CL/cl_kernels/warp_helpers_quantized.h new file mode 100644 index 0000000000..48d6faef73 --- /dev/null +++ b/src/core/CL/cl_kernels/warp_helpers_quantized.h @@ -0,0 +1,138 @@ +/* + * Copyright (c) 2018 ARM Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ +#include "helpers_asymm.h" + +/** Clamps the given coordinates to the borders according to the border size. + * + * @param[in] coords Vector of 2D coordinates to clamp. Even positions are X coords, odd positions are Y coords. + * @param[in] width Width of the image + * @param[in] height Height of the image + * @param[in] border_size Border size of the image + * + */ +inline const float8 clamp_to_border_with_size_quantized(float8 coords, const float width, const float height, const float border_size) +{ + const float4 clamped_x = clamp(coords.even, 0.0f - border_size, width - 1 + border_size); + const float4 clamped_y = clamp(coords.odd, 0.0f - border_size, height - 1 + border_size); + return (float8)(clamped_x.s0, clamped_y.s0, clamped_x.s1, clamped_y.s1, clamped_x.s2, clamped_y.s2, clamped_x.s3, clamped_y.s3); +} + +/* FIXME(COMPMID-682): Clamp border properly in UNDEFINED border mode in Warp, Scale, Remap */ +/** Clamps the given coordinates to the borders. + * + * @param[in] coords Vector of 2D coordinates to clamp. Even positions are X coords, odd positions are Y coords. + * @param[in] width Width of the image + * @param[in] height Height of the image + * + */ +inline const float8 clamp_to_border_quantized(float8 coords, const float width, const float height) +{ + return clamp_to_border_with_size_quantized(coords, width, height, 1); +} + +/** Given a texel coordinates this function will return the following array of coordinates: + * [ P, right neighbour, below neighbour, below right neighbour ] + * + * @note No checks to see if the coordinates are out of the image are done here. + * + * @param[in] coord Input coordinates + * + * @return vector of 8 floats with the coordinates, even positions are x and odd y. + */ +inline const float8 get_neighbour_coords_quantized(const float2 coord) +{ + return (float8)(/*tl*/ coord.s0, coord.s1, /*tr*/ coord.s0 + 1, coord.s1, /*bl*/ coord.s0, coord.s1 + 1, /*br*/ coord.s0 + 1, coord.s1 + 1); +} + +/** Returns the current thread coordinates. */ +inline const float2 get_current_coords_quantized() +{ + return (float2)(get_global_id(0) * 4, get_global_id(1)); +} + +/** Computes the bilinear interpolation for each set of coordinates in the vector coords and returns the values + * + * @param[in] in Pointer to the source image. + * @param[in] coords Vector of four 2D coordinates. Even pos is x and odd y. + * @param[in] width Width of the image + * @param[in] height Height of the image + * @param[in] border_size Border size + * @param[in] scale Scale value + * @param[in] offset_qasymm Offset value + */ +inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate_with_border_quantized(const Image *in, const float8 coords, const float width, const float height, const float border_size, + const float scale, const int offset_qasymm) +{ + // If any of the 4 texels is out of the image's boundaries we use the border value (REPLICATE or CONSTANT) for any texel out of the image. + + // Sets the 4x4 coordinates for each of the four input texels + const float8 fc = floor(coords); + const float16 c1 = (float16)( + clamp_to_border_with_size_quantized(get_neighbour_coords_quantized((float2)(fc.s0, fc.s1)), width, height, border_size), + clamp_to_border_with_size_quantized(get_neighbour_coords_quantized((float2)(fc.s2, fc.s3)), width, height, border_size)); + const float16 c2 = (float16)( + clamp_to_border_with_size_quantized(get_neighbour_coords_quantized((float2)(fc.s4, fc.s5)), width, height, border_size), + clamp_to_border_with_size_quantized(get_neighbour_coords_quantized((float2)(fc.s6, fc.s7)), width, height, border_size)); + + // Loads the values from the input image + const int16 t = (int16)( + /* tl, tr, bl, br */ + * ((__global DATA_TYPE *)offset(in, c1.s0, c1.s1)), *((__global DATA_TYPE *)offset(in, c1.s2, c1.s3)), + *((__global DATA_TYPE *)offset(in, c1.s4, c1.s5)), *((__global DATA_TYPE *)offset(in, c1.s6, c1.s7)), + *((__global DATA_TYPE *)offset(in, c1.s8, c1.s9)), *((__global DATA_TYPE *)offset(in, c1.sa, c1.sb)), + *((__global DATA_TYPE *)offset(in, c1.sc, c1.sd)), *((__global DATA_TYPE *)offset(in, c1.se, c1.sf)), + *((__global DATA_TYPE *)offset(in, c2.s0, c2.s1)), *((__global DATA_TYPE *)offset(in, c2.s2, c2.s3)), + *((__global DATA_TYPE *)offset(in, c2.s4, c2.s5)), *((__global DATA_TYPE *)offset(in, c2.s6, c2.s7)), + *((__global DATA_TYPE *)offset(in, c2.s8, c2.s9)), *((__global DATA_TYPE *)offset(in, c2.sa, c2.sb)), + *((__global DATA_TYPE *)offset(in, c2.sc, c2.sd)), *((__global DATA_TYPE *)offset(in, c2.se, c2.sf))); + + const float16 inf32 = convert_float16(t - (int16)offset_qasymm) * (float16)scale; + + const float8 a = coords - fc; + const float8 b = ((float8)(1.f)) - a; + const float4 fr = (float4)( + ((inf32.s0 * b.s0 * b.s1) + (inf32.s1 * a.s0 * b.s1) + (inf32.s2 * b.s0 * a.s1) + (inf32.s3 * a.s0 * a.s1)), + ((inf32.s4 * b.s2 * b.s3) + (inf32.s5 * a.s2 * b.s3) + (inf32.s6 * b.s2 * a.s3) + (inf32.s7 * a.s2 * a.s3)), + ((inf32.s8 * b.s4 * b.s5) + (inf32.s9 * a.s4 * b.s5) + (inf32.sa * b.s4 * a.s5) + (inf32.sb * a.s4 * a.s5)), + ((inf32.sc * b.s6 * b.s7) + (inf32.sd * a.s6 * b.s7) + (inf32.se * b.s6 * a.s7) + (inf32.sf * a.s6 * a.s7))); + + const uchar4 res = convert_uchar4_sat(convert_int4_sat_rtp(fr / scale) + offset_qasymm); + + return res; +} + +/* FIXME(COMPMID-682): Clamp border properly in UNDEFINED border mode in Warp, Scale, Remap */ +/** Computes the bilinear interpolation for each set of coordinates in the vector coords and returns the values + * + * @param[in] in Pointer to the source image. + * @param[in] coords Vector of four 2D coordinates. Even pos is x and odd y. + * @param[in] width Width of the image + * @param[in] height Height of the image + * @param[in] scale Scale value + * @param[in] offset_qasymm Offset value + */ +inline const VEC_DATA_TYPE(DATA_TYPE, 4) bilinear_interpolate_quantized(const Image *in, const float8 coords, const float width, const float height, const float scale, const int offset_qasymm) +{ + return bilinear_interpolate_with_border_quantized(in, coords, width, height, 1, scale, offset_qasymm); +} diff --git a/src/core/CL/kernels/CLScaleKernel.cpp b/src/core/CL/kernels/CLScaleKernel.cpp index d56d6f7da8..ce6c016154 100644 --- a/src/core/CL/kernels/CLScaleKernel.cpp +++ b/src/core/CL/kernels/CLScaleKernel.cpp @@ -62,7 +62,7 @@ inline std::pair calculate_scale_factors(const ITensorInfo &input, Status validate_arguments(const ITensorInfo *input, const ITensorInfo *output, InterpolationPolicy policy) { ARM_COMPUTE_RETURN_ERROR_ON_F16_UNSUPPORTED(input); - ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::U8, DataType::S16, DataType::F16, DataType::F32); + ARM_COMPUTE_RETURN_ERROR_ON_DATA_TYPE_CHANNEL_NOT_IN(input, 1, DataType::QASYMM8, DataType::U8, DataType::S16, DataType::F16, DataType::F32); ARM_COMPUTE_RETURN_ERROR_ON_NULLPTR(output); ARM_COMPUTE_RETURN_ERROR_ON_MISMATCHING_DATA_TYPES(input, output); ARM_COMPUTE_RETURN_ERROR_ON(output == input); @@ -170,6 +170,8 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo float hr = 0.f; std::tie(wr, hr) = calculate_scale_factors(*input->info(), *output->info()); + const bool call_quantized_kernel = is_data_type_quantized_asymmetric(input->info()->data_type()) && policy == InterpolationPolicy::BILINEAR; + DataLayout data_layout = input->info()->data_layout(); const int idx_width = get_data_layout_dimension_index(data_layout, DataLayoutDimension::WIDTH); const int idx_height = get_data_layout_dimension_index(data_layout, DataLayoutDimension::HEIGHT); @@ -200,11 +202,18 @@ void CLScaleKernel::configure(const ICLTensor *input, ICLTensor *output, Interpo build_opts.add_option("-DBORDER_SIZE=" + support::cpp11::to_string(border.right)); build_opts.add_option_if(border_mode == BorderMode::REPLICATE, "-DBORDER_MODE_REPLICATE"); build_opts.add_option_if_else(sampling_policy == SamplingPolicy::CENTER, "-DSAMPLING_POLICY_CENTER", "-DSAMPLING_POLICY_TOP_LEFT"); + if(call_quantized_kernel) + { + build_opts.add_option("-DSCALE=" + support::cpp11::to_string(input->info()->quantization_info().scale)); + build_opts.add_option("-DOFFSET=" + support::cpp11::to_string(input->info()->quantization_info().offset)); + } std::string interpolation_name = string_from_interpolation_policy(policy); std::transform(interpolation_name.begin(), interpolation_name.end(), interpolation_name.begin(), ::tolower); - std::string kernel_name = "scale_" + interpolation_name + "_" + lower_string(string_from_data_layout(data_layout)); - _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); + std::string kernel_name = "scale_" + interpolation_name; + kernel_name += call_quantized_kernel ? "_quantized_" : "_"; + kernel_name += lower_string(string_from_data_layout(data_layout)); + _kernel = static_cast(CLKernelLibrary::get().create_kernel(kernel_name, build_opts.options())); unsigned int idx = data_layout == DataLayout::NHWC ? 2 * num_arguments_per_3D_tensor() : 2 * num_arguments_per_2D_tensor(); //Skip the input and output parameters -- cgit v1.2.1