|
|
@ -186,21 +186,13 @@ __kernel void calcLut(__global __const uchar * src, const int srcStep, |
|
|
|
#else |
|
|
|
#else |
|
|
|
clipped = smem[0]; |
|
|
|
clipped = smem[0]; |
|
|
|
#endif |
|
|
|
#endif |
|
|
|
|
|
|
|
|
|
|
|
// broadcast evaluated value |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
__local int totalClipped; |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
if (tid == 0) |
|
|
|
|
|
|
|
totalClipped = clipped; |
|
|
|
|
|
|
|
barrier(CLK_LOCAL_MEM_FENCE); |
|
|
|
barrier(CLK_LOCAL_MEM_FENCE); |
|
|
|
|
|
|
|
|
|
|
|
// redistribute clipped samples evenly |
|
|
|
// redistribute clipped samples evenly |
|
|
|
|
|
|
|
int redistBatch = clipped / 256; |
|
|
|
int redistBatch = totalClipped / 256; |
|
|
|
|
|
|
|
tHistVal += redistBatch; |
|
|
|
tHistVal += redistBatch; |
|
|
|
|
|
|
|
|
|
|
|
int residual = totalClipped - redistBatch * 256; |
|
|
|
int residual = clipped - redistBatch * 256; |
|
|
|
int rStep = 256 / residual; |
|
|
|
int rStep = 256 / residual; |
|
|
|
if (rStep < 1) |
|
|
|
if (rStep < 1) |
|
|
|
rStep = 1; |
|
|
|
rStep = 1; |
|
|
|