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 // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. 7 // Third party copyrights are property of their respective owners. 8 // 9 // @Authors 10 // Niko Li, newlife20080214 (a] gmail.com 11 // Jia Haipeng, jiahaipeng95 (a] gmail.com 12 // Xu Pang, pangxu010 (a] 163.com 13 // Wenju He, wenju (a] multicorewareinc.com 14 // Redistribution and use in source and binary forms, with or without modification, 15 // are permitted provided that the following conditions are met: 16 // 17 // * Redistribution's of source code must retain the above copyright notice, 18 // this list of conditions and the following disclaimer. 19 // 20 // * Redistribution's in binary form must reproduce the above copyright notice, 21 // this list of conditions and the following disclaimer in the documentation 22 // and/or other materials provided with the distribution. 23 // 24 // * The name of the copyright holders may not be used to endorse or promote products 25 // derived from this software without specific prior written permission. 26 // 27 // This software is provided by the copyright holders and contributors as is and 28 // any express or implied warranties, including, but not limited to, the implied 29 // warranties of merchantability and fitness for a particular purpose are disclaimed. 30 // In no event shall the Intel Corporation or contributors be liable for any direct, 31 // indirect, incidental, special, exemplary, or consequential damages 32 // (including, but not limited to, procurement of substitute goods or services; 33 // loss of use, data, or profits; or business interruption) however caused 34 // and on any theory of liability, whether in contract, strict liability, 35 // or tort (including negligence or otherwise) arising in any way out of 36 // the use of this software, even if advised of the possibility of such damage. 37 // 38 // 39 40 #define OUT_OF_RANGE -1 41 42 // for identical rounding after dividing on different platforms 43 #define ROUNDING_EPS 0.000001f 44 45 #if histdims == 1 46 47 __kernel void calcLUT(__global const uchar * histptr, int hist_step, int hist_offset, int hist_bins, 48 __global int * lut, float scale, __constant float * ranges) 49 { 50 int x = get_global_id(0); 51 float value = convert_float(x); 52 53 if (value > ranges[1] || value < ranges[0]) 54 lut[x] = OUT_OF_RANGE; 55 else 56 { 57 float lb = ranges[0], ub = ranges[1], gap = (ub - lb) / hist_bins; 58 value -= lb; 59 int bin = convert_int_sat_rtn(value / gap + ROUNDING_EPS); 60 61 if (bin >= hist_bins) 62 lut[x] = OUT_OF_RANGE; 63 else 64 { 65 int hist_index = mad24(hist_step, bin, hist_offset); 66 __global const float * hist = (__global const float *)(histptr + hist_index); 67 68 lut[x] = (int)convert_uchar_sat_rte(hist[0] * scale); 69 } 70 } 71 } 72 73 __kernel void LUT(__global const uchar * src, int src_step, int src_offset, 74 __constant int * lut, 75 __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols) 76 { 77 int x = get_global_id(0); 78 int y = get_global_id(1); 79 80 if (x < dst_cols && y < dst_rows) 81 { 82 int src_index = mad24(y, src_step, src_offset + x * scn); 83 int dst_index = mad24(y, dst_step, dst_offset + x); 84 85 int value = lut[src[src_index]]; 86 dst[dst_index] = value == OUT_OF_RANGE ? 0 : convert_uchar(value); 87 } 88 } 89 90 #elif histdims == 2 91 92 __kernel void calcLUT(int hist_bins, __global int * lut, int lut_offset, 93 __constant float * ranges, int roffset) 94 { 95 int x = get_global_id(0); 96 float value = convert_float(x); 97 98 ranges += roffset; 99 lut += lut_offset; 100 101 if (value > ranges[1] || value < ranges[0]) 102 lut[x] = OUT_OF_RANGE; 103 else 104 { 105 float lb = ranges[0], ub = ranges[1], gap = (ub - lb) / hist_bins; 106 value -= lb; 107 int bin = convert_int_sat_rtn(value / gap + ROUNDING_EPS); 108 109 lut[x] = bin >= hist_bins ? OUT_OF_RANGE : bin; 110 } 111 } 112 113 __kernel void LUT(__global const uchar * src1, int src1_step, int src1_offset, 114 __global const uchar * src2, int src2_step, int src2_offset, 115 __global const uchar * histptr, int hist_step, int hist_offset, 116 __constant int * lut, float scale, 117 __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols) 118 { 119 int x = get_global_id(0); 120 int y = get_global_id(1); 121 122 if (x < dst_cols && y < dst_rows) 123 { 124 int src1_index = mad24(y, src1_step, src1_offset + x * scn1); 125 int src2_index = mad24(y, src2_step, src2_offset + x * scn2); 126 int dst_index = mad24(y, dst_step, dst_offset + x); 127 128 int bin1 = lut[src1[src1_index]]; 129 int bin2 = lut[src2[src2_index] + 256]; 130 dst[dst_index] = bin1 == OUT_OF_RANGE || bin2 == OUT_OF_RANGE ? 0 : 131 convert_uchar_sat_rte(*(__global const float *)(histptr + 132 mad24(hist_step, bin1, hist_offset + bin2 * (int)sizeof(float))) * scale); 133 } 134 } 135 136 #else 137 #error "(nimages <= 2) should be true" 138 #endif 139
