Path: blob/master/modules/calib3d/src/opencl/stereobm.cl
16348 views
/*M///////////////////////////////////////////////////////////////////////////////////////1//2// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.3//4// By downloading, copying, installing or using the software you agree to this license.5// If you do not agree to this license, do not download, install,6// copy or use the software.7//8//9// License Agreement10// For Open Source Computer Vision Library11//12// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.13// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.14// Third party copyrights are property of their respective owners.15//16// Redistribution and use in source and binary forms, with or without modification,17// are permitted provided that the following conditions are met:18//19// * Redistribution's of source code must retain the above copyright notice,20// this list of conditions and the following disclaimer.21//22// * Redistribution's in binary form must reproduce the above copyright notice,23// this list of conditions and the following disclaimer in the documentation24// and/or other materials provided with the distribution.25//26// * The name of the copyright holders may not be used to endorse or promote products27// derived from this software without specific prior written permission.28//29// This software is provided by the copyright holders and contributors as is and30// any express or implied warranties, including, but not limited to, the implied31// warranties of merchantability and fitness for a particular purpose are disclaimed.32// In no event shall the Intel Corporation or contributors be liable for any direct,33// indirect, incidental, special, exemplary, or consequential damages34// (including, but not limited to, procurement of substitute goods or services;35// loss of use, data, or profits; or business interruption) however caused36// and on any theory of liability, whether in contract, strict liability,37// or tort (including negligence or otherwise) arising in any way out of38// the use of this software, even if advised of the possibility of such damage.39//40//M*/4142//////////////////////////////////////////////////////////////////////////////////////////////////43////////////////////////////////////////// stereoBM //////////////////////////////////////////////44//////////////////////////////////////////////////////////////////////////////////////////////////4546#define MAX_VAL 327674748#ifndef WSZ49#define WSZ 250#endif5152#define WSZ2 (WSZ / 2)5354#ifdef DEFINE_KERNEL_STEREOBM5556#define DISPARITY_SHIFT 457#define FILTERED ((MIN_DISP - 1) << DISPARITY_SHIFT)5859void calcDisp(__local short * cost, __global short * disp, int uniquenessRatio,60__local int * bestDisp, __local int * bestCost, int d, int x, int y, int cols, int rows)61{62int best_disp = *bestDisp, best_cost = *bestCost;63barrier(CLK_LOCAL_MEM_FENCE);6465short c = cost[0];66int thresh = best_cost + (best_cost * uniquenessRatio / 100);67bool notUniq = ( (c <= thresh) && (d < (best_disp - 1) || d > (best_disp + 1) ) );6869if (notUniq)70*bestCost = FILTERED;71barrier(CLK_LOCAL_MEM_FENCE);7273if( *bestCost != FILTERED && x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2 && d == best_disp)74{75int d_aprox = 0;76int yp =0, yn = 0;77if ((0 < best_disp) && (best_disp < NUM_DISP - 1))78{79yp = cost[-2 * BLOCK_SIZE_Y];80yn = cost[2 * BLOCK_SIZE_Y];81d_aprox = yp + yn - 2 * c + abs(yp - yn);82}83disp[0] = (short)(((best_disp + MIN_DISP)*256 + (d_aprox != 0 ? (yp - yn) * 256 / d_aprox : 0) + 15) >> 4);84}85}8687short calcCostBorder(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y, int nthread,88short * costbuf, int *h, int cols, int d, short cost)89{90int head = (*h) % WSZ;91__global const uchar * left, * right;92int idx = mad24(y + WSZ2 * (2 * nthread - 1), cols, x + WSZ2 * (1 - 2 * nthread));93left = leftptr + idx;94right = rightptr + (idx - d);9596short costdiff = 0;97if (0 == nthread)98{99#pragma unroll100for (int i = 0; i < WSZ; i++)101{102costdiff += abs( left[0] - right[0] );103left += cols;104right += cols;105}106}107else // (1 == nthread)108{109#pragma unroll110for (int i = 0; i < WSZ; i++)111{112costdiff += abs(left[i] - right[i]);113}114}115cost += costdiff - costbuf[head];116costbuf[head] = costdiff;117*h = head + 1;118return cost;119}120121short calcCostInside(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y,122int cols, int d, short cost_up_left, short cost_up, short cost_left)123{124__global const uchar * left, * right;125int idx = mad24(y - WSZ2 - 1, cols, x - WSZ2 - 1);126left = leftptr + idx;127right = rightptr + (idx - d);128int idx2 = WSZ*cols;129130uchar corrner1 = abs(left[0] - right[0]),131corrner2 = abs(left[WSZ] - right[WSZ]),132corrner3 = abs(left[idx2] - right[idx2]),133corrner4 = abs(left[idx2 + WSZ] - right[idx2 + WSZ]);134135return cost_up + cost_left - cost_up_left + corrner1 -136corrner2 - corrner3 + corrner4;137}138139__kernel void stereoBM(__global const uchar * leftptr,140__global const uchar * rightptr,141__global uchar * dispptr, int disp_step, int disp_offset,142int rows, int cols, // rows, cols of left and right images, not disp143int textureThreshold, int uniquenessRatio)144{145int lz = get_local_id(0);146int gx = get_global_id(1) * BLOCK_SIZE_X;147int gy = get_global_id(2) * BLOCK_SIZE_Y;148149int nthread = lz / NUM_DISP;150int disp_idx = lz % NUM_DISP;151152__global short * disp;153__global const uchar * left, * right;154155__local short costFunc[2 * BLOCK_SIZE_Y * NUM_DISP];156157__local short * cost;158__local int best_disp[2];159__local int best_cost[2];160best_cost[nthread] = MAX_VAL;161best_disp[nthread] = -1;162barrier(CLK_LOCAL_MEM_FENCE);163164short costbuf[WSZ];165int head = 0;166167int shiftX = WSZ2 + NUM_DISP + MIN_DISP - 1;168int shiftY = WSZ2;169170int x = gx + shiftX, y = gy + shiftY, lx = 0, ly = 0;171172int costIdx = disp_idx * 2 * BLOCK_SIZE_Y + (BLOCK_SIZE_Y - 1);173cost = costFunc + costIdx;174175int tempcost = 0;176if (x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2)177{178if (0 == nthread)179{180#pragma unroll181for (int i = 0; i < WSZ; i++)182{183int idx = mad24(y - WSZ2, cols, x - WSZ2 + i);184left = leftptr + idx;185right = rightptr + (idx - disp_idx);186short costdiff = 0;187for(int j = 0; j < WSZ; j++)188{189costdiff += abs( left[0] - right[0] );190left += cols;191right += cols;192}193costbuf[i] = costdiff;194}195}196else // (1 == nthread)197{198#pragma unroll199for (int i = 0; i < WSZ; i++)200{201int idx = mad24(y - WSZ2 + i, cols, x - WSZ2);202left = leftptr + idx;203right = rightptr + (idx - disp_idx);204short costdiff = 0;205for (int j = 0; j < WSZ; j++)206{207costdiff += abs( left[j] - right[j]);208}209tempcost += costdiff;210costbuf[i] = costdiff;211}212}213}214if (nthread == 1)215{216cost[0] = tempcost;217atomic_min(best_cost + 1, tempcost);218}219barrier(CLK_LOCAL_MEM_FENCE);220221if (best_cost[1] == tempcost)222atomic_max(best_disp + 1, disp_idx);223barrier(CLK_LOCAL_MEM_FENCE);224225int dispIdx = mad24(gy, disp_step, mad24((int)sizeof(short), gx, disp_offset));226disp = (__global short *)(dispptr + dispIdx);227calcDisp(cost, disp, uniquenessRatio, best_disp + 1, best_cost + 1, disp_idx, x, y, cols, rows);228barrier(CLK_LOCAL_MEM_FENCE);229230lx = 1 - nthread;231ly = nthread;232233for (int i = 0; i < BLOCK_SIZE_Y * BLOCK_SIZE_X / 2; i++)234{235x = (lx < BLOCK_SIZE_X) ? gx + shiftX + lx : cols;236y = (ly < BLOCK_SIZE_Y) ? gy + shiftY + ly : rows;237238best_cost[nthread] = MAX_VAL;239best_disp[nthread] = -1;240barrier(CLK_LOCAL_MEM_FENCE);241242costIdx = mad24(2 * BLOCK_SIZE_Y, disp_idx, (BLOCK_SIZE_Y - 1 - ly + lx));243if (0 > costIdx)244costIdx = BLOCK_SIZE_Y - 1;245cost = costFunc + costIdx;246if (x < cols - WSZ2 - MIN_DISP && y < rows - WSZ2)247{248tempcost = (ly * (1 - nthread) + lx * nthread == 0) ?249calcCostBorder(leftptr, rightptr, x, y, nthread, costbuf, &head, cols, disp_idx, cost[2*nthread-1]) :250calcCostInside(leftptr, rightptr, x, y, cols, disp_idx, cost[0], cost[1], cost[-1]);251}252cost[0] = tempcost;253atomic_min(best_cost + nthread, tempcost);254barrier(CLK_LOCAL_MEM_FENCE);255256if (best_cost[nthread] == tempcost)257atomic_max(best_disp + nthread, disp_idx);258barrier(CLK_LOCAL_MEM_FENCE);259260dispIdx = mad24(gy + ly, disp_step, mad24((int)sizeof(short), (gx + lx), disp_offset));261disp = (__global short *)(dispptr + dispIdx);262calcDisp(cost, disp, uniquenessRatio, best_disp + nthread, best_cost + nthread, disp_idx, x, y, cols, rows);263264barrier(CLK_LOCAL_MEM_FENCE);265266if (lx + nthread - 1 == ly)267{268lx = (lx + nthread + 1) * (1 - nthread);269ly = (ly + 1) * nthread;270}271else272{273lx += nthread;274ly = ly - nthread + 1;275}276}277}278#endif //DEFINE_KERNEL_STEREOBM279280//////////////////////////////////////////////////////////////////////////////////////////////////281/////////////////////////////////////// Norm Prefiler ////////////////////////////////////////////282//////////////////////////////////////////////////////////////////////////////////////////////////283284__kernel void prefilter_norm(__global unsigned char *input, __global unsigned char *output,285int rows, int cols, int prefilterCap, int scale_g, int scale_s)286{287// prefilterCap in range 1..63, checked in StereoBMImpl::compute288289int x = get_global_id(0);290int y = get_global_id(1);291292if(x < cols && y < rows)293{294int cov1 = input[ max(y-1, 0) * cols + x] * 1 +295input[y * cols + max(x-1,0)] * 1 + input[ y * cols + x] * 4 + input[y * cols + min(x+1, cols-1)] * 1 +296input[min(y+1, rows-1) * cols + x] * 1;297int cov2 = 0;298for(int i = -WSZ2; i < WSZ2+1; i++)299for(int j = -WSZ2; j < WSZ2+1; j++)300cov2 += input[clamp(y+i, 0, rows-1) * cols + clamp(x+j, 0, cols-1)];301302int res = (cov1*scale_g - cov2*scale_s)>>10;303res = clamp(res, -prefilterCap, prefilterCap) + prefilterCap;304output[y * cols + x] = res;305}306}307308309//////////////////////////////////////////////////////////////////////////////////////////////////310////////////////////////////////////// Sobel Prefiler ////////////////////////////////////////////311//////////////////////////////////////////////////////////////////////////////////////////////////312313__kernel void prefilter_xsobel(__global unsigned char *input, __global unsigned char *output,314int rows, int cols, int prefilterCap)315{316// prefilterCap in range 1..63, checked in StereoBMImpl::compute317int x = get_global_id(0);318int y = get_global_id(1);319if(x < cols && y < rows)320{321if (0 < x && !((y == rows-1) & (rows%2==1) ) )322{323int 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) +324input[ (y) * cols + (x-1)] * (-2) + input[ (y) * cols + ((x<cols-1) ? x+1 : x-1)] * (2) +325input[((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);326327cov = clamp(cov, -prefilterCap, prefilterCap) + prefilterCap;328output[y * cols + x] = cov;329}330else331output[y * cols + x] = prefilterCap;332}333}334335