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