src/core/CL/cl_kernels/normalization_layer.cl

   1 /*
   2  * Copyright (c) 2017 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(FIXED_POINT_POSITION)
  27
  28 #include "fixed_point.h"
  29 #define MUL_OP(x, y) MUL_SAT_OP_EXPAND((x), (y), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
  30 #define ADD_OP(x, y) ADD_SAT_OP_EXPAND((x), (y), DATA_TYPE, VEC_SIZE)
  31 #define DIV_OP(x, y) DIV_SAT_OP_VEC_EXPAND((x), (y), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
  32 #define EXP_OP(x) EXP_OP_EXPAND((x), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
  33 #define LOG_OP(x) LOG_OP_EXPAND((x), DATA_TYPE, VEC_SIZE, FIXED_POINT_POSITION)
  34 #define POW_OP(x, y) EXP_OP(MUL_OP(LOG_OP((x)), (y)))
  35 #define SQCVT_SAT(a) SQCVT_SAT_OP_EXPAND((a), DATA_TYPE, FIXED_POINT_POSITION)
  36
  37 #define LOAD_OP(offset, ptr) vload16(offset, ptr)
  38 #define STORE_OP(data, offset, ptr) vstore16(data, offset, ptr)
  39
  40 #else // FIXED_POINT_POSITION
  41
  42 #define MUL_OP(x, y) ((x) * (y))
  43 #define ADD_OP(x, y) ((x) + (y))
  44 #define DIV_OP(x, y) ((x) / (y))
  45 #define POW_OP(x, y) pow((x), (y))
  46 #define SQCVT_SAT(a) (a)
  47
  48 #define LOAD_OP(offset, ptr) vload4(offset, ptr)
  49 #define STORE_OP(data, offset, ptr) vstore4(data, offset, ptr)
  50
  51 #endif // FIXED_POINT_POSITION
  52
  53 /** Apply cross-map normalization.
  54  *
  55  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
  56  * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE=size, e.g. -DVEC_SIZE=16
  57  * @note The radius should be given as a preprocessor argument using -DRADIUS=size. e.g. -DRADIUS=5
  58  * @note The number of slices should be given as a preprocessor argument using -DNUM_SLICES=size. e.g. -DNUM_SLICES=192
  59  * @note In case of fixed-point operation -DFIXED_POINT_POSITION=fixed_point_position must be provided: e.g. -DFIXED_POINT_POSITION=3
  60  * @note Scaling coefficient (= alpha/norm_size), beta and kappa need to be passed at compile time using -DCOEFF, -DALPHA and -DKAPPA
  61  *
  62  * @param[in]  input_ptr                            Pointer to the first source tensor. Supported data types: QS8/QS16/F16/F32
  63  * @param[in]  input_stride_x                       Stride of the first source tensor in X dimension (in bytes)
  64  * @param[in]  input_step_x                         input_stride_x * number of elements along X processed per workitem(in bytes)
  65  * @param[in]  input_stride_y                       Stride of the first source tensor in Y dimension (in bytes)
  66  * @param[in]  input_step_y                         input_stride_y * number of elements along Y processed per workitem(in bytes)
  67  * @param[in]  input_stride_z                       Stride of the first source tensor in Z dimension (in bytes)
  68  * @param[in]  input_step_z                         input_stride_z * number of elements along Z processed per workitem(in bytes)
  69  * @param[in]  input_offset_first_element_in_bytes  The offset of the first element in the first source tensor
  70  * @param[out] output_ptr                           Pointer to the destination tensor. Supported data types: same as @p input_ptr
  71  * @param[in]  output_stride_x                      Stride of the destination tensor in X dimension (in bytes)
  72  * @param[in]  output_step_x                        output_stride_x * number of elements along X processed per workitem(in bytes)
  73  * @param[in]  output_stride_y                      Stride of the destination tensor in Y dimension (in bytes)
  74  * @param[in]  output_step_y                        output_stride_y * number of elements along Y processed per workitem(in bytes)
  75  * @param[in]  output_stride_z                      Stride of the destination tensor in Z dimension (in bytes)
  76  * @param[in]  output_step_z                        output_stride_z * number of elements along Z processed per workitem(in bytes)
  77  * @param[in]  output_offset_first_element_in_bytes The offset of the first element in the destination tensor
  78  */
  79 __kernel void normalization_layer_cross_map(TENSOR3D_DECLARATION(input),
  80                                             TENSOR3D_DECLARATION(output))
  81 {
  82     Tensor3D in  = CONVERT_TO_TENSOR3D_STRUCT(input);
  83     Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output);
  84
  85     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
  86     acc = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0;
  87     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
  88     coeff_v = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))SQCVT_SAT(COEFF);
  89     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
  90     beta_v = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))SQCVT_SAT(BETA);
  91     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
  92     kappa_v = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))SQCVT_SAT(KAPPA);
  93
  94     const int current_slice = get_global_id(2);
  95     const int left_slice    = max(-(int)RADIUS, -current_slice);
  96     const int right_slice   = min((int)RADIUS, (int)NUM_SLICES - 1 - current_slice);
  97
  98     for(int i = left_slice; i <= right_slice; i++)
  99     {
 100         VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 101         values = LOAD_OP(0, (__global DATA_TYPE *)tensor3D_offset(&in, 0, 0, i));
 102         acc    = ADD_OP(acc, MUL_OP(values, values));
 103     }
 104
 105     acc = ADD_OP(MUL_OP(acc, coeff_v), kappa_v);
 106     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 107     normalized = POW_OP(acc, beta_v);
 108     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 109     normalized_pixel = DIV_OP(LOAD_OP(0, (__global DATA_TYPE *)in.ptr), normalized);
 110
 111     STORE_OP(normalized_pixel, 0, (__global DATA_TYPE *)out.ptr);
 112 }
 113
 114 /** Apply in-map normalization.
 115  *
 116  * @note Datatype should be given as a preprocessor argument using -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
 117  * @note Vector size should be given as a preprocessor argument using -DVEC_SIZE=size, e.g. -DVEC_SIZE=16
 118  * @note The radius should be given as a preprocessor argument using -DRADIUS=size. e.g. -DRADIUS=5
 119  * @note In case of fixed-point operation -DFIXED_POINT_POSITION=fixed_point_position must be provided: e.g. -DFIXED_POINT_POSITION=3
 120  * @note Scaling coefficient (= alpha/norm_size), beta and kappa need to be passed at compile time using -DCOEFF, -DALPHA and -DKAPPA
 121  *
 122  * @param[in]  input_ptr                            Pointer to the first source tensor. Supported data types: QS8/F16/F32
 123  * @param[in]  input_stride_x                       Stride of the first source tensor in X dimension (in bytes)
 124  * @param[in]  input_step_x                         input_stride_x * number of elements along X processed per workitem(in bytes)
 125  * @param[in]  input_stride_y                       Stride of the first source tensor in Y dimension (in bytes)
 126  * @param[in]  input_step_y                         input_stride_y * number of elements along Y processed per workitem(in bytes)
 127  * @param[in]  input_stride_z                       Stride of the first source tensor in Z dimension (in bytes)
 128  * @param[in]  input_step_z                         input_stride_z * number of elements along Z processed per workitem(in bytes)
 129  * @param[in]  input_offset_first_element_in_bytes  The offset of the first element in the first source tensor
 130  * @param[out] output_ptr                           Pointer to the destination tensor. Supported data types: same as @p input_ptr
 131  * @param[in]  output_stride_x                      Stride of the destination tensor in X dimension (in bytes)
 132  * @param[in]  output_step_x                        output_stride_x * number of elements along X processed per workitem(in bytes)
 133  * @param[in]  output_stride_y                      Stride of the first destination tensor in Y dimension (in bytes)
 134  * @param[in]  output_step_y                        output_stride_y * number of elements along Y processed per workitem(in bytes)
 135  * @param[in]  output_stride_z                      Stride of the first source tensor in Z dimension (in bytes)
 136  * @param[in]  output_step_z                        output_stride_z * number of elements along Z processed per workitem(in bytes)
 137  * @param[in]  output_offset_first_element_in_bytes The offset of the first element in the destination tensor
 138  */
 139 __kernel void normalization_layer_in_map(TENSOR3D_DECLARATION(input),
 140                                          TENSOR3D_DECLARATION(output))
 141 {
 142     Tensor3D in  = CONVERT_TO_TENSOR3D_STRUCT(input);
 143     Tensor3D out = CONVERT_TO_TENSOR3D_STRUCT(output);
 144
 145     VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 146     acc = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))0;
 147     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 148     coeff_v = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))SQCVT_SAT(COEFF);
 149     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 150     beta_v = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))SQCVT_SAT(BETA);
 151     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 152     kappa_v = (VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))SQCVT_SAT(KAPPA);
 153
 154     const int current_col = get_global_id(0) << 2;
 155     const int left_pos    = max(-(int)RADIUS, -3 - current_col);
 156     const int right_pos   = min((int)RADIUS, (int)((get_global_size(0) << 2) + 3 - 1 - current_col));
 157
 158 #if defined(IN_MAP_2D)
 159     const int current_row = get_global_id(1);
 160     const int first_row   = max(-(int)RADIUS, -current_row);
 161     const int last_row    = min((int)RADIUS, (int)get_global_size(1) - 1 - current_row);
 162 #endif /* defined(IN_MAP_2D) */
 163
 164 #if defined(IN_MAP_2D)
 165     for(int j = first_row; j <= last_row; ++j)
 166     {
 167 #endif /* defined(IN_MAP_2D) */
 168         for(int i = left_pos; i <= right_pos; ++i)
 169         {
 170 #if defined(IN_MAP_2D)
 171             VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 172             values = LOAD_OP(0, (__global DATA_TYPE *)tensor3D_offset(&in, i, j, 0));
 173 #else  /* defined(IN_MAP_2D) */
 174             VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 175             values = LOAD_OP(0, (__global DATA_TYPE *)tensor3D_offset(&in, i, 0, 0));
 176 #endif /* defined(IN_MAP_2D) */
 177             acc = ADD_OP(acc, MUL_OP(values, values));
 178         }
 179 #if defined(IN_MAP_2D)
 180     }
 181 #endif /* defined(IN_MAP_2D) */
 182
 183     acc = ADD_OP(MUL_OP(acc, coeff_v), kappa_v);
 184     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 185     normalized = POW_OP(acc, beta_v);
 186     const VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)
 187     normalized_pixel = DIV_OP(LOAD_OP(0, (__global DATA_TYPE *)in.ptr), normalized);
 188
 189     STORE_OP(normalized_pixel, 0, (__global DATA_TYPE *)out.ptr);
 190 }