inference-engine/thirdparty/clDNN/kernel_selector/core/cl_kernels/reorder_data.cl

   1 // Copyright (c) 2016-2017 Intel Corporation
   2 //
   3 // Licensed under the Apache License, Version 2.0 (the "License");
   4 // you may not use this file except in compliance with the License.
   5 // You may obtain a copy of the License at
   6 //
   7 //      http://www.apache.org/licenses/LICENSE-2.0
   8 //
   9 // Unless required by applicable law or agreed to in writing, software
  10 // distributed under the License is distributed on an "AS IS" BASIS,
  11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  12 // See the License for the specific language governing permissions and
  13 // limitations under the License.
  14
  15
  16 #include "include/reshape_dims.cl"
  17 #include "include/fetch.cl"
  18
  19 #include "include/data_types.cl"
  20
  21 ///////////////////////// Input Index /////////////////////////
  22 inline uint FUNC(get_input_index)(uint b, uint f, uint y, uint x)
  23 {
  24 #if   INPUT0_SIMPLE
  25     return GET_DATA_INDEX(INPUT0, b, f, y, x);
  26 #elif defined INPUT0_LAYOUT_BS_F_BSV8__AF8  || \
  27       defined INPUT0_LAYOUT_BS_F_BSV16__AF8
  28     return GET_DATA_BS_FYX_BSV8_INDEX(INPUT0, b, f, y, x, SUB_GROUP_SIZE);
  29 #elif defined INPUT0_LAYOUT_BF8_XY16
  30     return GET_DATA_BF8_XY16_INDEX(INPUT0, b, f, y, x);
  31 #elif defined INPUT0_LAYOUT_BYXF_AF32
  32         return GET_DATA_BYXF_AF32_INDEX(INPUT0, b, f, y, x);
  33 #elif defined INPUT0_LAYOUT_BYX8_F4
  34         return GET_DATA_BYX8_F4_INDEX(INPUT0, b, f, y, x);
  35 #elif defined INPUT0_LAYOUT_FS_BS_YX_BSV4_FSV32
  36     return GET_DATA_FS_BS_YX_BSV4_FSV32_INDEX(INPUT0, b, f, y, x);
  37 #elif defined INPUT0_LAYOUT_B_FS_YX_FSV4
  38     return GET_DATA_B_FS_YX_FSV4_INDEX(INPUT0, b, f, y, x);
  39 #else
  40 #error reorder_data.cl: input format - not supported
  41 #endif
  42 }
  43
  44 ///////////////////////// Output Index /////////////////////////
  45
  46 inline uint FUNC(get_output_index)(uint b, uint f, uint y, uint x)
  47 {
  48 #if   OUTPUT_SIMPLE
  49     return GET_DATA_INDEX(OUTPUT, b, f, y, x);
  50 #elif defined OUTPUT_LAYOUT_BS_F_BSV8__AF8  || \
  51       defined OUTPUT_LAYOUT_BS_F_BSV16__AF8
  52     return GET_DATA_BS_FYX_BSV8_INDEX(OUTPUT, b, f, y, x, SUB_GROUP_SIZE);
  53 #elif defined OUTPUT_LAYOUT_BF8_XY16
  54     return GET_DATA_BF8_XY16_INDEX(OUTPUT, b, f, y, x);
  55 #elif defined OUTPUT_LAYOUT_BYXF_AF32
  56         return GET_DATA_BYXF_AF32_INDEX(OUTPUT, b, f, y, x);
  57 #elif defined OUTPUT_LAYOUT_BYX8_F4
  58         return GET_DATA_BYX8_F4_INDEX(OUTPUT, b, f, y, x);
  59 #elif defined OUTPUT_LAYOUT_FS_BS_YX_BSV4_FSV32
  60     return GET_DATA_FS_BS_YX_BSV4_FSV32_INDEX(OUTPUT, b, f, y, x);
  61 #elif defined OUTPUT_LAYOUT_B_FS_YX_FSV4
  62     return GET_DATA_B_FS_YX_FSV4_INDEX(OUTPUT, b, f, y, x);
  63 #else
  64 #error reorder_data.cl: output format - not supported
  65 #endif
  66 }
  67
  68 KERNEL (reorder_data)(
  69     const __global INPUT_REORDER_TYPE* input,
  70     __global OUTPUT_REORDER_TYPE* output
  71 #ifdef MEAN_SUBTRACT_IN_BUFFER
  72     , __global MEAN_SUBTRACT_TYPE* mean_subtract
  73 #endif
  74     )
  75 {
  76     const uint b = get_global_id(GWS_BATCH);
  77     const uint f = get_global_id(GWS_FEATURE);
  78 #if   INPUT0_DIMS == 2
  79     const uint y = 0;
  80     const uint x = 0;
  81 #elif INPUT0_DIMS == 4
  82     const uint y = ((uint)(get_global_id(GWS_YX))) / INPUT0_SIZE_X;
  83     const uint x = ((uint)(get_global_id(GWS_YX))) % INPUT0_SIZE_X;
  84 #endif
  85
  86     uint4 ov = FUNC_CALL(reshape_dims)(b,f,y,x, INPUT0_SIZE_Y, INPUT0_SIZE_X, OUTPUT_SIZE_Y, OUTPUT_SIZE_X, INPUT0_DIMS, OUTPUT_DIMS);
  87     const uint input_idx  = FUNC_CALL(get_input_index)(b, f, y, x);
  88     const uint output_idx = FUNC_CALL(get_output_index)(ov[0],ov[1],ov[2],ov[3]);
  89
  90 #if defined MEAN_SUBTRACT_INSIDE_PARAMS
  91     float res = TO_MEAN_TYPE(input[input_idx]);
  92     res = MEAN_OP(res, VALUE_TO_SUBTRACT[f % VALUE_TO_SUBTRACT_SIZE]);
  93 #elif defined MEAN_SUBTRACT_IN_BUFFER
  94 #if defined MEAN_PER_FEATURE
  95     MEAN_SUBTRACT_TYPE res = TO_MEAN_TYPE(input[input_idx]);
  96     res = MEAN_OP(res, mean_subtract[f]);
  97 #else
  98     MEAN_SUBTRACT_TYPE res = TO_MEAN_TYPE(input[input_idx]);
  99     uint4 msv = FUNC_CALL(reshape_dims)(b,f,y,x, INPUT0_SIZE_Y, INPUT0_SIZE_X, MEAN_SUBTRACT_SIZE_Y, MEAN_SUBTRACT_SIZE_X, INPUT0_DIMS, MEAN_SUBTRACT_DIMS);
 100     res = MEAN_OP(res, mean_subtract[GET_DATA_INDEX_SAFE(MEAN_SUBTRACT, msv[0], msv[1], msv[2], msv[3])]);
 101 #endif
 102 #else
 103     CALC_TYPE res = TO_CALC_TYPE(input[input_idx]);
 104 #endif
 105
 106     output[output_idx] = ACTIVATION(TO_OUTPUT_REORDER_TYPE(res), NL_M ,NL_N);
 107 }