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 // Rock Li, Rock.li (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 #if lcn == 1 38 #if dcn == 4 39 #define LUT_OP \ 40 int idx = *(__global const int *)(srcptr + src_index); \ 41 dst = (__global dstT *)(dstptr + dst_index); \ 42 dst[0] = lut_l[idx & 0xff]; \ 43 dst[1] = lut_l[(idx >> 8) & 0xff]; \ 44 dst[2] = lut_l[(idx >> 16) & 0xff]; \ 45 dst[3] = lut_l[(idx >> 24) & 0xff]; 46 #elif dcn == 3 47 #define LUT_OP \ 48 uchar3 idx = vload3(0, srcptr + src_index); \ 49 dst = (__global dstT *)(dstptr + dst_index); \ 50 dst[0] = lut_l[idx.x]; \ 51 dst[1] = lut_l[idx.y]; \ 52 dst[2] = lut_l[idx.z]; 53 #elif dcn == 2 54 #define LUT_OP \ 55 short idx = *(__global const short *)(srcptr + src_index); \ 56 dst = (__global dstT *)(dstptr + dst_index); \ 57 dst[0] = lut_l[idx & 0xff]; \ 58 dst[1] = lut_l[(idx >> 8) & 0xff]; 59 #elif dcn == 1 60 #define LUT_OP \ 61 uchar idx = (srcptr + src_index)[0]; \ 62 dst = (__global dstT *)(dstptr + dst_index); \ 63 dst[0] = lut_l[idx]; 64 #else 65 #define LUT_OP \ 66 __global const srcT * src = (__global const srcT *)(srcptr + src_index); \ 67 dst = (__global dstT *)(dstptr + dst_index); \ 68 for (int cn = 0; cn < dcn; ++cn) \ 69 dst[cn] = lut_l[src[cn]]; 70 #endif 71 #else 72 #if dcn == 4 73 #define LUT_OP \ 74 __global const uchar4 * src_pixel = (__global const uchar4 *)(srcptr + src_index); \ 75 int4 idx = mad24(convert_int4(src_pixel[0]), (int4)(lcn), (int4)(0, 1, 2, 3)); \ 76 dst = (__global dstT *)(dstptr + dst_index); \ 77 dst[0] = lut_l[idx.x]; \ 78 dst[1] = lut_l[idx.y]; \ 79 dst[2] = lut_l[idx.z]; \ 80 dst[3] = lut_l[idx.w]; 81 #elif dcn == 3 82 #define LUT_OP \ 83 uchar3 src_pixel = vload3(0, srcptr + src_index); \ 84 int3 idx = mad24(convert_int3(src_pixel), (int3)(lcn), (int3)(0, 1, 2)); \ 85 dst = (__global dstT *)(dstptr + dst_index); \ 86 dst[0] = lut_l[idx.x]; \ 87 dst[1] = lut_l[idx.y]; \ 88 dst[2] = lut_l[idx.z]; 89 #elif dcn == 2 90 #define LUT_OP \ 91 __global const uchar2 * src_pixel = (__global const uchar2 *)(srcptr + src_index); \ 92 int2 idx = mad24(convert_int2(src_pixel[0]), lcn, (int2)(0, 1)); \ 93 dst = (__global dstT *)(dstptr + dst_index); \ 94 dst[0] = lut_l[idx.x]; \ 95 dst[1] = lut_l[idx.y]; 96 #elif dcn == 1 //error case (1 < lcn) ==> lcn == scn == dcn 97 #define LUT_OP \ 98 uchar idx = (srcptr + src_index)[0]; \ 99 dst = (__global dstT *)(dstptr + dst_index); \ 100 dst[0] = lut_l[idx]; 101 #else 102 #define LUT_OP \ 103 __global const srcT * src = (__global const srcT *)(srcptr + src_index); \ 104 dst = (__global dstT *)(dstptr + dst_index); \ 105 for (int cn = 0; cn < dcn; ++cn) \ 106 dst[cn] = lut_l[mad24(src[cn], lcn, cn)]; 107 #endif 108 #endif 109 110 __kernel void LUT(__global const uchar * srcptr, int src_step, int src_offset, 111 __global const uchar * lutptr, int lut_step, int lut_offset, 112 __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols) 113 { 114 int x = get_global_id(0); 115 int y = get_global_id(1) << 2; 116 117 __local dstT lut_l[256 * lcn]; 118 __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset); 119 120 for (int i = mad24((int)get_local_id(1), (int)get_local_size(0), (int)get_local_id(0)), 121 step = get_local_size(0) * get_local_size(1); i < 256 * lcn; i += step) 122 lut_l[i] = lut[i]; 123 barrier(CLK_LOCAL_MEM_FENCE); 124 125 if (x < cols && y < rows) 126 { 127 int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset)); 128 int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset)); 129 130 __global dstT * dst; 131 132 LUT_OP; 133 134 if (y < rows - 1) 135 { 136 src_index += src_step; 137 dst_index += dst_step; 138 LUT_OP; 139 140 if (y < rows - 2) 141 { 142 src_index += src_step; 143 dst_index += dst_step; 144 LUT_OP; 145 146 if (y < rows - 3) 147 { 148 src_index += src_step; 149 dst_index += dst_step; 150 LUT_OP; 151 } 152 } 153 } 154 155 } 156 } 157
