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@gmail.com 11// Jia Haipeng, jiahaipeng95@gmail.com 12// Xu Pang, pangxu010@163.com 13// Wenju He, wenju@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