1 // License Agreement 2 // For Open Source Computer Vision Library 3 // 4 // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved. 5 // Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. 6 // Third party copyrights are property of their respective owners. 7 // 8 // @Authors 9 // Niko Li, newlife20080214 (a] gmail.com 10 // Zero Lin zero.lin (a] amd.com 11 // Redistribution and use in source and binary forms, with or without modification, 12 // are permitted provided that the following conditions are met: 13 // 14 // * Redistribution's of source code must retain the above copyright notice, 15 // this list of conditions and the following disclaimer. 16 // 17 // * Redistribution's in binary form must reproduce the above copyright notice, 18 // this list of conditions and the following disclaimer in the documentation 19 // and/or other materials provided with the distribution. 20 // 21 // * The name of the copyright holders may not be used to endorse or promote products 22 // derived from this software without specific prior written permission. 23 // 24 // This software is provided by the copyright holders and contributors as is and 25 // any express or implied warranties, including, but not limited to, the implied 26 // warranties of merchantability and fitness for a particular purpose are disclaimed. 27 // In no event shall the Intel Corporation or contributors be liable for any direct, 28 // indirect, incidental, special, exemplary, or consequential damages 29 // (including, but not limited to, procurement of substitute goods or services; 30 // loss of use, data, or profits; or business interruption) however caused 31 // and on any theory of liability, whether in contract, strict liability, 32 // or tort (including negligence or otherwise) arising in any way out of 33 // the use of this software, even if advised of the possibility of such damage. 34 // 35 // 36 37 #ifdef DOUBLE_SUPPORT 38 #ifdef cl_amd_fp64 39 #pragma OPENCL EXTENSION cl_amd_fp64:enable 40 #elif defined (cl_khr_fp64) 41 #pragma OPENCL EXTENSION cl_khr_fp64:enable 42 #endif 43 #endif 44 45 #if cn != 3 46 #define loadpix(addr) *(__global const T*)(addr) 47 #define storepix(val, addr) *(__global T*)(addr) = val 48 #define TSIZE ((int)sizeof(T)) 49 #define convertScalar(a) (a) 50 #else 51 #define loadpix(addr) vload3(0, (__global const T1*)(addr)) 52 #define storepix(val, addr) vstore3(val, 0, (__global T1*)(addr)) 53 #define TSIZE ((int)sizeof(T1)*3) 54 #define convertScalar(a) (T)(a.x, a.y, a.z) 55 #endif 56 57 #ifdef BORDER_CONSTANT 58 #define EXTRAPOLATE(x, cols) \ 59 ; 60 #elif defined BORDER_REPLICATE 61 #define EXTRAPOLATE(x, cols) \ 62 x = clamp(x, 0, cols - 1); 63 #elif defined BORDER_WRAP 64 #define EXTRAPOLATE(x, cols) \ 65 { \ 66 if (x < 0) \ 67 x -= ((x - cols + 1) / cols) * cols; \ 68 if (x >= cols) \ 69 x %= cols; \ 70 } 71 #elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) 72 #ifdef BORDER_REFLECT 73 #define DELTA int delta = 0 74 #else 75 #define DELTA int delta = 1 76 #endif 77 #define EXTRAPOLATE(x, cols) \ 78 { \ 79 DELTA; \ 80 if (cols == 1) \ 81 x = 0; \ 82 else \ 83 do \ 84 { \ 85 if( x < 0 ) \ 86 x = -x - 1 + delta; \ 87 else \ 88 x = cols - 1 - (x - cols) - delta; \ 89 } \ 90 while (x >= cols || x < 0); \ 91 } 92 #else 93 #error "No extrapolation method" 94 #endif 95 96 #define NEED_EXTRAPOLATION(x, cols) (x >= cols || x < 0) 97 98 __kernel void copyMakeBorder(__global const uchar * srcptr, int src_step, int src_offset, int src_rows, int src_cols, 99 __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, 100 int top, int left, ST nVal) 101 { 102 int x = get_global_id(0); 103 int y0 = get_global_id(1) * rowsPerWI; 104 105 #ifdef BORDER_CONSTANT 106 T scalar = convertScalar(nVal); 107 #endif 108 109 if (x < dst_cols) 110 { 111 int src_x = x - left, src_y; 112 int dst_index = mad24(y0, dst_step, mad24(x, (int)TSIZE, dst_offset)); 113 114 if (NEED_EXTRAPOLATION(src_x, src_cols)) 115 { 116 #ifdef BORDER_CONSTANT 117 for (int y = y0, y1 = min(y0 + rowsPerWI, dst_rows); y < y1; ++y, dst_index += dst_step) 118 storepix(scalar, dstptr + dst_index); 119 return; 120 #endif 121 EXTRAPOLATE(src_x, src_cols) 122 } 123 src_x = mad24(src_x, TSIZE, src_offset); 124 125 for (int y = y0, y1 = min(y0 + rowsPerWI, dst_rows); y < y1; ++y, dst_index += dst_step) 126 { 127 src_y = y - top; 128 if (NEED_EXTRAPOLATION(src_y, src_rows)) 129 { 130 EXTRAPOLATE(src_y, src_rows) 131 #ifdef BORDER_CONSTANT 132 storepix(scalar, dstptr + dst_index); 133 continue; 134 #endif 135 } 136 int src_index = mad24(src_y, src_step, src_x); 137 storepix(loadpix(srcptr + src_index), dstptr + dst_index); 138 } 139 } 140 } 141
