src/core/CL/cl_kernels/fill_border.cl

   1 /*
   2  * Copyright (c) 2016-2018 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 #include "fixed_point.h"
  28 #endif /* FIXED_POINT_POSITION */
  29
  30 /** Fill N pixel of the padding edge of a single channel image by replicating the closest valid pixel.
  31  *
  32  * @attention  The DATA_TYPE needs to be passed at the compile time.
  33  * e.g. -DDATA_TYPE=int
  34  *
  35  * @attention  The border size for top, bottom, left, right needs to be passed at the compile time.
  36  * e.g. --DBORDER_SIZE_TOP=0 -DBORDER_SIZE_BOTTOM=2 -DBORDER_SIZE_LEFT=0 -DBORDER_SIZE_RIGHT=2
  37  *
  38  * @param[in,out] buf_ptr                           Pointer to the source image. Supported data types: U8/U16/S16/U32/S32/F16/F32
  39  * @param[in]     buf_stride_x                      Stride of the source image in X dimension (in bytes)
  40  * @param[in]     buf_step_x                        buf_stride_x * number of elements along X processed per workitem(in bytes)
  41  * @param[in]     buf_stride_y                      Stride of the source image in Y dimension (in bytes)
  42  * @param[in]     buf_step_y                        buf_stride_y * number of elements along Y processed per workitem(in bytes)
  43  * @param[in]     buf_stride_z                      Stride between images if batching images (in bytes)
  44  * @param[in]     buf_step_z                        buf_stride_z * number of elements along Z processed per workitem(in bytes)
  45  * @param[in]     buf_offset_first_element_in_bytes The offset of the first element in the source image
  46  * @param[in]     width                             Width of the valid region of the image
  47  * @param[in]     height                            Height of the valid region of the image
  48  * @param[in]     start_pos                         XY coordinate indicating the start point of the valid region
  49  */
  50 __kernel void fill_image_borders_replicate(
  51     TENSOR3D_DECLARATION(buf),
  52     uint width,
  53     uint height,
  54     int2 start_pos)
  55 {
  56     Image buf = CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(buf);
  57
  58     // Update pointer to point to the starting point of the valid region
  59     buf.ptr += start_pos.y * buf.stride_y + start_pos.x * buf.stride_x;
  60
  61     const int total_width = BORDER_SIZE_LEFT + width + BORDER_SIZE_RIGHT;
  62     const int gid0        = get_global_id(0);
  63     const int gidH        = gid0 - total_width;
  64     const int gidW        = gid0 - BORDER_SIZE_LEFT;
  65
  66     if(gidH >= 0)
  67     {
  68         // Handle left border
  69         DATA_TYPE left_val = *(__global DATA_TYPE *)offset(&buf, 0, gidH);
  70         for(int i = -BORDER_SIZE_LEFT; i < 0; ++i)
  71         {
  72             *(__global DATA_TYPE *)offset(&buf, i, gidH) = left_val;
  73         }
  74         // Handle right border
  75         DATA_TYPE right_val = *(__global DATA_TYPE *)offset(&buf, width - 1, gidH);
  76         for(int i = 0; i < BORDER_SIZE_RIGHT; ++i)
  77         {
  78             *(__global DATA_TYPE *)offset(&buf, width + i, gidH) = right_val;
  79         }
  80     }
  81     else
  82     {
  83         // Get value for corners
  84         int val_idx = gidW;
  85         if(gidW < 0 || gidW > (width - 1))
  86         {
  87             val_idx = gidW < 0 ? 0 : width - 1;
  88         }
  89
  90         // Handle top border
  91         DATA_TYPE top_val = *(__global DATA_TYPE *)offset(&buf, val_idx, 0);
  92         for(int i = -BORDER_SIZE_TOP; i < 0; ++i)
  93         {
  94             *(__global DATA_TYPE *)offset(&buf, gidW, i) = top_val;
  95         }
  96         // Handle bottom border
  97         DATA_TYPE bottom_val = *(__global DATA_TYPE *)offset(&buf, val_idx, height - 1);
  98         for(int i = 0; i < BORDER_SIZE_BOTTOM; ++i)
  99         {
 100             *(__global DATA_TYPE *)offset(&buf, gidW, height + i) = bottom_val;
 101         }
 102     }
 103 }
 104
 105 /** Fill N pixels of the padding edge of a single channel image with a constant value.
 106  *
 107  * @attention  The DATA_TYPE needs to be passed at the compile time.
 108  * e.g. -DDATA_TYPE=int
 109  *
 110  * @attention  The border size for top, bottom, left, right needs to be passed at the compile time.
 111  * e.g. --DBORDER_SIZE_TOP=0 -DBORDER_SIZE_BOTTOM=2 -DBORDER_SIZE_LEFT=0 -DBORDER_SIZE_RIGHT=2
 112  *
 113  * @param[out] buf_ptr                           Pointer to the source image. Supported data types: U8/U16/S16/U32/S32/F16/F32
 114  * @param[in]  buf_stride_x                      Stride of the source image in X dimension (in bytes)
 115  * @param[in]  buf_step_x                        buf_stride_x * number of elements along X processed per workitem(in bytes)
 116  * @param[in]  buf_stride_y                      Stride of the source image in Y dimension (in bytes)
 117  * @param[in]  buf_step_y                        buf_stride_y * number of elements along Y processed per workitem(in bytes)
 118  * @param[in]  buf_stride_z                      Stride between images if batching images (in bytes)
 119  * @param[in]  buf_step_z                        buf_stride_z * number of elements along Z processed per workitem(in bytes)
 120  * @param[in]  buf_offset_first_element_in_bytes The offset of the first element in the source image
 121  * @param[in]  width                             Width of the valid region of the image
 122  * @param[in]  height                            Height of the valid region of the image
 123  * @param[in]  start_pos                         XY coordinate indicating the start point of the valid region
 124  * @param[in]  constant_value                    Constant value to use to fill the edges
 125  */
 126 __kernel void fill_image_borders_constant(
 127     TENSOR3D_DECLARATION(buf),
 128     uint      width,
 129     uint      height,
 130     int2      start_pos,
 131     DATA_TYPE constant_value)
 132 {
 133     Image buf = CONVERT_TENSOR3D_TO_IMAGE_STRUCT_NO_STEP(buf);
 134
 135     // Update pointer to point to the starting point of the valid region
 136     buf.ptr += start_pos.y * buf.stride_y + start_pos.x * buf.stride_x;
 137
 138     const int total_width = BORDER_SIZE_LEFT + width + BORDER_SIZE_RIGHT;
 139     const int gid0        = get_global_id(0);
 140     const int gidH        = gid0 - total_width;
 141     const int gidW        = gid0 - BORDER_SIZE_LEFT;
 142
 143     if(gidH >= 0)
 144     {
 145         // Handle left border
 146         for(int i = -BORDER_SIZE_LEFT; i < 0; ++i)
 147         {
 148             *(__global DATA_TYPE *)offset(&buf, i, gidH) = constant_value;
 149         }
 150         // Handle right border
 151         for(int i = 0; i < BORDER_SIZE_RIGHT; ++i)
 152         {
 153             *(__global DATA_TYPE *)offset(&buf, width + i, gidH) = constant_value;
 154         }
 155     }
 156     else
 157     {
 158         // Handle top border
 159         for(int i = -BORDER_SIZE_TOP; i < 0; ++i)
 160         {
 161             *(__global DATA_TYPE *)offset(&buf, gidW, i) = constant_value;
 162         }
 163         // Handle bottom border
 164         for(int i = 0; i < BORDER_SIZE_BOTTOM; ++i)
 165         {
 166             *(__global DATA_TYPE *)offset(&buf, gidW, height + i) = constant_value;
 167         }
 168     }
 169 }