1 /*M///////////////////////////////////////////////////////////////////////////////////////
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install,
7 // copy or use the software.
11 // For Open Source Computer Vision Library
13 // Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
14 // Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
15 // Third party copyrights are property of their respective owners.
17 // Redistribution and use in source and binary forms, with or without modification,
18 // are permitted provided that the following conditions are met:
20 // * Redistribution's of source code must retain the above copyright notice,
21 // this list of conditions and the following disclaimer.
23 // * Redistribution's in binary form must reproduce the above copyright notice,
24 // this list of conditions and the following disclaimer in the documentation
25 // and/or other materials provided with the distribution.
27 // * The name of the copyright holders may not be used to endorse or promote products
28 // derived from this software without specific prior written permission.
30 // This software is provided by the copyright holders and contributors as is and
31 // any express or implied warranties, including, but not limited to, the implied
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
34 // indirect, incidental, special, exemplary, or consequential damages
35 // (including, but not limited to, procurement of substitute goods or services;
36 // loss of use, data, or profits; or business interruption) however caused
37 // and on any theory of liability, whether in contract, strict liability,
38 // or tort (including negligence or otherwise) arising in any way out of
39 // the use of this software, even if advised of the possibility of such damage.
43 //////////////////////////////////////////////////////////////////////////////////////////////////
44 ////////////////////////////////////////// stereoBM //////////////////////////////////////////////
45 //////////////////////////////////////////////////////////////////////////////////////////////////
53 #define WSZ2 (WSZ / 2)
55 #ifdef DEFINE_KERNEL_STEREOBM
57 #define DISPARITY_SHIFT 4
58 #define FILTERED ((MIN_DISP - 1) << DISPARITY_SHIFT)
60 void calcDisp(__local short * cost, __global short * disp, int uniquenessRatio,
61 __local int * bestDisp, __local int * bestCost, int d, int x, int y, int cols, int rows)
63 int best_disp = *bestDisp, best_cost = *bestCost;
64 barrier(CLK_LOCAL_MEM_FENCE);
67 int thresh = best_cost + (best_cost * uniquenessRatio / 100);
68 bool notUniq = ( (c <= thresh) && (d < (best_disp - 1) || d > (best_disp + 1) ) );
72 barrier(CLK_LOCAL_MEM_FENCE);
74 if( *bestCost != FILTERED && x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2 && d == best_disp)
78 if ((0 < best_disp) && (best_disp < NUM_DISP - 1))
80 yp = cost[-2 * BLOCK_SIZE_Y];
81 yn = cost[2 * BLOCK_SIZE_Y];
82 d_aprox = yp + yn - 2 * c + abs(yp - yn);
84 disp[0] = (short)(((best_disp + MIN_DISP)*256 + (d_aprox != 0 ? (yp - yn) * 256 / d_aprox : 0) + 15) >> 4);
88 short calcCostBorder(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y, int nthread,
89 short * costbuf, int *h, int cols, int d, short cost)
91 int head = (*h) % WSZ;
92 __global const uchar * left, * right;
93 int idx = mad24(y + WSZ2 * (2 * nthread - 1), cols, x + WSZ2 * (1 - 2 * nthread));
95 right = rightptr + (idx - d);
101 for (int i = 0; i < WSZ; i++)
103 costdiff += abs( left[0] - right[0] );
108 else // (1 == nthread)
111 for (int i = 0; i < WSZ; i++)
113 costdiff += abs(left[i] - right[i]);
116 cost += costdiff - costbuf[head];
117 costbuf[head] = costdiff;
122 short calcCostInside(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y,
123 int cols, int d, short cost_up_left, short cost_up, short cost_left)
125 __global const uchar * left, * right;
126 int idx = mad24(y - WSZ2 - 1, cols, x - WSZ2 - 1);
127 left = leftptr + idx;
128 right = rightptr + (idx - d);
131 uchar corrner1 = abs(left[0] - right[0]),
132 corrner2 = abs(left[WSZ] - right[WSZ]),
133 corrner3 = abs(left[idx2] - right[idx2]),
134 corrner4 = abs(left[idx2 + WSZ] - right[idx2 + WSZ]);
136 return cost_up + cost_left - cost_up_left + corrner1 -
137 corrner2 - corrner3 + corrner4;
140 __kernel void stereoBM(__global const uchar * leftptr,
141 __global const uchar * rightptr,
142 __global uchar * dispptr, int disp_step, int disp_offset,
143 int rows, int cols, // rows, cols of left and right images, not disp
144 int textureTreshold, int uniquenessRatio)
146 int lz = get_local_id(0);
147 int gx = get_global_id(1) * BLOCK_SIZE_X;
148 int gy = get_global_id(2) * BLOCK_SIZE_Y;
150 int nthread = lz / NUM_DISP;
151 int disp_idx = lz % NUM_DISP;
153 __global short * disp;
154 __global const uchar * left, * right;
156 __local short costFunc[2 * BLOCK_SIZE_Y * NUM_DISP];
158 __local short * cost;
159 __local int best_disp[2];
160 __local int best_cost[2];
161 best_cost[nthread] = MAX_VAL;
162 best_disp[nthread] = -1;
163 barrier(CLK_LOCAL_MEM_FENCE);
168 int shiftX = WSZ2 + NUM_DISP + MIN_DISP - 1;
171 int x = gx + shiftX, y = gy + shiftY, lx = 0, ly = 0;
173 int costIdx = disp_idx * 2 * BLOCK_SIZE_Y + (BLOCK_SIZE_Y - 1);
174 cost = costFunc + costIdx;
177 if (x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2)
182 for (int i = 0; i < WSZ; i++)
184 int idx = mad24(y - WSZ2, cols, x - WSZ2 + i);
185 left = leftptr + idx;
186 right = rightptr + (idx - disp_idx);
188 for(int j = 0; j < WSZ; j++)
190 costdiff += abs( left[0] - right[0] );
194 costbuf[i] = costdiff;
197 else // (1 == nthread)
200 for (int i = 0; i < WSZ; i++)
202 int idx = mad24(y - WSZ2 + i, cols, x - WSZ2);
203 left = leftptr + idx;
204 right = rightptr + (idx - disp_idx);
206 for (int j = 0; j < WSZ; j++)
208 costdiff += abs( left[j] - right[j]);
210 tempcost += costdiff;
211 costbuf[i] = costdiff;
218 atomic_min(best_cost + 1, tempcost);
220 barrier(CLK_LOCAL_MEM_FENCE);
222 if (best_cost[1] == tempcost)
223 atomic_max(best_disp + 1, disp_idx);
224 barrier(CLK_LOCAL_MEM_FENCE);
226 int dispIdx = mad24(gy, disp_step, mad24((int)sizeof(short), gx, disp_offset));
227 disp = (__global short *)(dispptr + dispIdx);
228 calcDisp(cost, disp, uniquenessRatio, best_disp + 1, best_cost + 1, disp_idx, x, y, cols, rows);
229 barrier(CLK_LOCAL_MEM_FENCE);
234 for (int i = 0; i < BLOCK_SIZE_Y * BLOCK_SIZE_X / 2; i++)
236 x = (lx < BLOCK_SIZE_X) ? gx + shiftX + lx : cols;
237 y = (ly < BLOCK_SIZE_Y) ? gy + shiftY + ly : rows;
239 best_cost[nthread] = MAX_VAL;
240 best_disp[nthread] = -1;
241 barrier(CLK_LOCAL_MEM_FENCE);
243 costIdx = mad24(2 * BLOCK_SIZE_Y, disp_idx, (BLOCK_SIZE_Y - 1 - ly + lx));
245 costIdx = BLOCK_SIZE_Y - 1;
246 cost = costFunc + costIdx;
247 if (x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2)
249 tempcost = (ly * (1 - nthread) + lx * nthread == 0) ?
250 calcCostBorder(leftptr, rightptr, x, y, nthread, costbuf, &head, cols, disp_idx, cost[2*nthread-1]) :
251 calcCostInside(leftptr, rightptr, x, y, cols, disp_idx, cost[0], cost[1], cost[-1]);
254 atomic_min(best_cost + nthread, tempcost);
255 barrier(CLK_LOCAL_MEM_FENCE);
257 if (best_cost[nthread] == tempcost)
258 atomic_max(best_disp + nthread, disp_idx);
259 barrier(CLK_LOCAL_MEM_FENCE);
261 dispIdx = mad24(gy + ly, disp_step, mad24((int)sizeof(short), (gx + lx), disp_offset));
262 disp = (__global short *)(dispptr + dispIdx);
263 calcDisp(cost, disp, uniquenessRatio, best_disp + nthread, best_cost + nthread, disp_idx, x, y, cols, rows);
265 barrier(CLK_LOCAL_MEM_FENCE);
267 if (lx + nthread - 1 == ly)
269 lx = (lx + nthread + 1) * (1 - nthread);
270 ly = (ly + 1) * nthread;
275 ly = ly - nthread + 1;
279 #endif //DEFINE_KERNEL_STEREOBM
281 //////////////////////////////////////////////////////////////////////////////////////////////////
282 /////////////////////////////////////// Norm Prefiler ////////////////////////////////////////////
283 //////////////////////////////////////////////////////////////////////////////////////////////////
285 __kernel void prefilter_norm(__global unsigned char *input, __global unsigned char *output,
286 int rows, int cols, int prefilterCap, int scale_g, int scale_s)
288 // prefilterCap in range 1..63, checked in StereoBMImpl::compute
290 int x = get_global_id(0);
291 int y = get_global_id(1);
293 if(x < cols && y < rows)
295 int cov1 = input[ max(y-1, 0) * cols + x] * 1 +
296 input[y * cols + max(x-1,0)] * 1 + input[ y * cols + x] * 4 + input[y * cols + min(x+1, cols-1)] * 1 +
297 input[min(y+1, rows-1) * cols + x] * 1;
299 for(int i = -WSZ2; i < WSZ2+1; i++)
300 for(int j = -WSZ2; j < WSZ2+1; j++)
301 cov2 += input[clamp(y+i, 0, rows-1) * cols + clamp(x+j, 0, cols-1)];
303 int res = (cov1*scale_g - cov2*scale_s)>>10;
304 res = clamp(res, -prefilterCap, prefilterCap) + prefilterCap;
305 output[y * cols + x] = res;
310 //////////////////////////////////////////////////////////////////////////////////////////////////
311 ////////////////////////////////////// Sobel Prefiler ////////////////////////////////////////////
312 //////////////////////////////////////////////////////////////////////////////////////////////////
314 __kernel void prefilter_xsobel(__global unsigned char *input, __global unsigned char *output,
315 int rows, int cols, int prefilterCap)
317 // prefilterCap in range 1..63, checked in StereoBMImpl::compute
318 int x = get_global_id(0);
319 int y = get_global_id(1);
320 if(x < cols && y < rows)
322 if (0 < x && !((y == rows-1) & (rows%2==1) ) )
324 int cov = input[ ((y > 0) ? y-1 : y+1) * cols + (x-1)] * (-1) + input[ ((y > 0) ? y-1 : y+1) * cols + ((x<cols-1) ? x+1 : x-1)] * (1) +
325 input[ (y) * cols + (x-1)] * (-2) + input[ (y) * cols + ((x<cols-1) ? x+1 : x-1)] * (2) +
326 input[((y<rows-1)?(y+1):(y-1))* cols + (x-1)] * (-1) + input[((y<rows-1)?(y+1):(y-1))* cols + ((x<cols-1) ? x+1 : x-1)] * (1);
328 cov = clamp(cov, -prefilterCap, prefilterCap) + prefilterCap;
329 output[y * cols + x] = cov;
332 output[y * cols + x] = prefilterCap;