Compute Library
 22.08
dequantization_layer.cl
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2017-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 "helpers.h"
25 
26 #if defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST)
27 /** This performs per channel dequantization of 8-bit signed integers to floating point. (NCHW)
28  *
29  * @note Source datatype should be given as a preprocessor argument using -DDATA_TYPE_SRC=type. e.g. -DDATA_TYPE_SRC=char
30  * @note Destination datatype should be given as a preprocessor argument using -DDATA_TYPE_DST=type. e.g. -DDATA_TYPE_DST=float
31  * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE=size. e.g. -DVEC_SIZE=16
32  *
33  * @param[in] input_ptr Pointer to the source tensor. Supported data types: QSYMM8_PER_CHANNEL
34  * @param[in] input_stride_x Stride of the source tensor in X dimension (in bytes)
35  * @param[in] input_step_x input_stride_x * number of elements along X processed per workitem(in bytes)
36  * @param[in] input_stride_y Stride of the source tensor in Y dimension (in bytes)
37  * @param[in] input_step_y input_stride_y * number of elements along Y processed per workitem(in bytes)
38  * @param[in] input_stride_z Stride of the source tensor in Z dimension (in bytes)
39  * @param[in] input_step_z input_stride_z * number of elements along Z processed per workitem(in bytes)
40  * @param[in] input_offset_first_element_in_bytes The offset of the first element in the source tensor
41  * @param[out] output_ptr Pointer to the destination tensor. Supported data types: F16/F32
42  * @param[in] output_stride_x Stride of the destination tensor in X dimension (in bytes)
43  * @param[in] output_step_x output_stride_x * number of elements along X processed per workitem(in bytes)
44  * @param[in] output_stride_y Stride of the destination tensor in Y dimension (in bytes)
45  * @param[in] output_step_y output_stride_y * number of elements along Y processed per workitem(in bytes)
46  * @param[in] output_stride_z Stride of the source tensor in Z dimension (in bytes)
47  * @param[in] output_step_z output_stride_z * number of elements along Z processed per workitem(in bytes)
48  * @param[in] output_offset_first_element_in_bytes The offset of the first element in the destination tensor
49  * @param[in] scale Pointer to buffer with the per channel quantized scales
50  */
51 __kernel void dequantization_layer_per_channel_nchw(
53  TENSOR3D_DECLARATION(output),
54  __global float *scale)
55 {
56  // Get pixels pointer
58  Tensor3D output = CONVERT_TO_TENSOR3D_STRUCT(output);
59 
60 #if defined(LAST_ACCESSED_X)
61  // Check if access on width gets out of bounds
62  // If it does shift access vector to access elements within bounds
63  const int xi = (int)(get_global_id(0) * VEC_SIZE);
64  input.ptr -= max(xi - (int)LAST_ACCESSED_X, 0) * input_stride_x;
65  output.ptr -= max(xi - (int)LAST_ACCESSED_X, 0) * output_stride_x;
66 
67  // Load data
69  val = CONVERT(VLOAD(VEC_SIZE)(0, (__global DATA_TYPE_SRC *)input.ptr), VEC_DATA_TYPE(int, VEC_SIZE));
70 
71  // Create scale vectors
72  const VEC_DATA_TYPE(float, VEC_SIZE)
73  vscale = scale[get_global_id(2)];
74 
75  // Dequantize
76  VEC_DATA_TYPE(float, VEC_SIZE)
77  res = vscale * CONVERT((val), VEC_DATA_TYPE(float, VEC_SIZE));
78 
79  // Store result
81  (CONVERT(res, VEC_DATA_TYPE(DATA_TYPE_DST, VEC_SIZE)), 0, (__global DATA_TYPE_DST *)output.ptr);
82 #else // !defined(LAST_ACCESSED_X)
83  *((__global DATA_TYPE_DST *)(output.ptr)) = (DATA_TYPE_DST)((float)((int)(*((__global DATA_TYPE_SRC *)(input.ptr)))) * scale[get_global_id(2)]);
84 #endif // defined(LAST_ACCESSED_X)
85 }
86 #endif // defined(VEC_SIZE) && defined(DATA_TYPE_SRC) && defined(DATA_TYPE_DST)
#define VEC_SIZE
#define CONVERT(x, type)
Definition: helpers.h:731
Structure to hold 3D tensor information.
Definition: helpers.h:906
#define CONVERT_TO_TENSOR3D_STRUCT(name)
Definition: helpers.h:870
#define VSTORE(size)
Definition: helpers.h:458
__global uchar * ptr
Pointer to the starting postion of the buffer.
Definition: helpers.h:908
#define VLOAD(size)
Definition: helpers.h:204
#define TENSOR3D_DECLARATION(name)
Definition: helpers.h:813
#define VEC_DATA_TYPE(type, size)
Definition: helpers.h:728