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

   1 // Copyright (c) 2018 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 #include "include/include_all.cl"
  16
  17
  18 KERNEL(border_gpu_ref)(
  19     const __global UNIT_TYPE* input,
  20     __global UNIT_TYPE* output)
  21 {
  22     // [CONSTEXPR]
  23     // Border sizes (left-top set and right-bottom set):
  24     const uint blt_sx = LT_SIZES_SIZE_X;
  25     const uint blt_sy = LT_SIZES_SIZE_Y;
  26     const uint blt_sf = LT_SIZES_FEATURE_NUM;
  27     const uint blt_sb = LT_SIZES_BATCH_NUM;
  28
  29     const uint brb_sx = RB_SIZES_SIZE_X;
  30     const uint brb_sy = RB_SIZES_SIZE_Y;
  31     const uint brb_sf = RB_SIZES_FEATURE_NUM;
  32     const uint brb_sb = RB_SIZES_BATCH_NUM;
  33
  34     // Input sizes:
  35     const uint in_sx = INPUT0_SIZE_X;
  36     const uint in_sy = INPUT0_SIZE_Y;
  37     const uint in_sf = INPUT0_FEATURE_NUM;
  38     const uint in_sb = INPUT0_BATCH_NUM;
  39
  40     // Input limits (exclusive; tested on output position):
  41     const uint in_lx = in_sx + blt_sx;
  42     const uint in_ly = in_sy + blt_sy;
  43     const uint in_lf = in_sf + blt_sf;
  44     const uint in_lb = in_sb + blt_sb;
  45
  46
  47     const uint out_x  = (uint) get_global_id(0);
  48     const uint out_y  = (uint) get_global_id(1);
  49     const uint out_fb = (uint) get_global_id(2);
  50
  51     const uint out_f  = out_fb % OUTPUT_FEATURE_NUM;
  52     const uint out_b  = out_fb / OUTPUT_FEATURE_NUM;
  53
  54 #ifdef BORDER_TYPE_CONSTANT
  55     UNIT_TYPE in_val = TO_UNIT_TYPE(BORDER_VALUE);
  56     if (out_x >= blt_sx & out_x < in_lx &
  57         out_y >= blt_sy & out_y < in_ly &
  58         out_f >= blt_sf & out_f < in_lf &
  59         out_b >= blt_sb & out_b < in_lb)
  60     {
  61         const uint in_x = out_x - blt_sx;
  62         const uint in_y = out_y - blt_sy;
  63         const uint in_f = out_f - blt_sf;
  64         const uint in_b = out_b - blt_sb;
  65
  66         const uint in_pos = GET_DATA_INDEX(INPUT0, in_b, in_f, in_y, in_x);
  67         in_val = input[in_pos];
  68     }
  69 #elif defined BORDER_TYPE_EDGE
  70     const uint in_x = (out_x >= blt_sx & out_x < in_lx) ? out_x - blt_sx : (out_x < blt_sx ? 0 : in_sx - 1);
  71     const uint in_y = (out_y >= blt_sy & out_y < in_ly) ? out_y - blt_sy : (out_y < blt_sy ? 0 : in_sy - 1);
  72     const uint in_f = (out_f >= blt_sf & out_f < in_lf) ? out_f - blt_sf : (out_f < blt_sf ? 0 : in_sf - 1);
  73     const uint in_b = (out_b >= blt_sb & out_b < in_lb) ? out_b - blt_sb : (out_b < blt_sb ? 0 : in_sb - 1);
  74
  75     const uint in_pos = GET_DATA_INDEX(INPUT0, in_b, in_f, in_y, in_x);
  76     UNIT_TYPE in_val = input[in_pos];
  77 #elif defined BORDER_TYPE_MIRROR
  78     const uint in_x = (out_x >= blt_sx & out_x < in_lx) ? out_x - blt_sx : (out_x < blt_sx ? blt_sx - 1 - out_x : in_sx + in_lx - 1 - out_x);
  79     const uint in_y = (out_y >= blt_sy & out_y < in_ly) ? out_y - blt_sy : (out_y < blt_sy ? blt_sy - 1 - out_y : in_sy + in_ly - 1 - out_y);
  80     const uint in_f = (out_f >= blt_sf & out_f < in_lf) ? out_f - blt_sf : (out_f < blt_sf ? blt_sf - 1 - out_f : in_sf + in_lf - 1 - out_f);
  81     const uint in_b = (out_b >= blt_sb & out_b < in_lb) ? out_b - blt_sb : (out_b < blt_sb ? blt_sb - 1 - out_b : in_sb + in_lb - 1 - out_b);
  82
  83     const uint in_pos = GET_DATA_INDEX(INPUT0, in_b, in_f, in_y, in_x);
  84     UNIT_TYPE in_val = input[in_pos];
  85 #elif defined BORDER_TYPE_MIRROR_101
  86     const uint in_x = (out_x >= blt_sx & out_x < in_lx) ? out_x - blt_sx : (out_x < blt_sx ? blt_sx - out_x : in_sx + in_lx - 2 - out_x);
  87     const uint in_y = (out_y >= blt_sy & out_y < in_ly) ? out_y - blt_sy : (out_y < blt_sy ? blt_sy - out_y : in_sy + in_ly - 2 - out_y);
  88     const uint in_f = (out_f >= blt_sf & out_f < in_lf) ? out_f - blt_sf : (out_f < blt_sf ? blt_sf - out_f : in_sf + in_lf - 2 - out_f);
  89     const uint in_b = (out_b >= blt_sb & out_b < in_lb) ? out_b - blt_sb : (out_b < blt_sb ? blt_sb - out_b : in_sb + in_lb - 2 - out_b);
  90
  91     const uint in_pos = GET_DATA_INDEX(INPUT0, in_b, in_f, in_y, in_x);
  92     UNIT_TYPE in_val = input[in_pos];
  93 #else
  94     #error Unsupported border type.
  95 #endif
  96
  97     const uint out_pos = GET_DATA_INDEX(OUTPUT, out_b, out_f, out_y, out_x);
  98     output[out_pos] = in_val;
  99 }