diff --git a/modules/calib3d/src/opencl/stereobm.cl b/modules/calib3d/src/opencl/stereobm.cl index e5f553b689..caca016eef 100644 --- a/modules/calib3d/src/opencl/stereobm.cl +++ b/modules/calib3d/src/opencl/stereobm.cl @@ -51,47 +51,27 @@ #define MAX_VAL 32767 void calcDisp(__local short * cost, __global short * disp, int uniquenessRatio/*, int textureTreshold, short textsum*/, - int mindisp, int ndisp, int w, __local short * dispbuf, int d, int x, int y, int cols, int rows, int wsz2) + int mindisp, int ndisp, int w, __local int * bestDisp, __local int * bestCost, int d, int x, int y, int cols, int rows, int wsz2) { short FILTERED = (mindisp - 1)<<4; - short best_disp, best_cost; - - dispbuf[d] = d; - barrier(CLK_LOCAL_MEM_FENCE); - - for(int lsize = tsize/2 >> 1; lsize > 0; lsize >>= 1) - { - short lid1 = dispbuf[d], lid2 = dispbuf[d+lsize], - cost1 = cost[lid1*w], cost2 = cost[lid2*w]; - if (d < lsize) - { - dispbuf[d] = (cost1 < cost2) ? lid1 : (cost1==cost2) ? max(lid1, lid2) : lid2; - } - barrier(CLK_LOCAL_MEM_FENCE); - } - best_disp = ndisp - dispbuf[0] - 1; - best_cost = cost[(ndisp-best_disp-1)*w]; - barrier(CLK_LOCAL_MEM_FENCE); + int best_disp = *bestDisp, best_cost = *bestCost, best_disp_back = ndisp - best_disp - 1; int thresh = best_cost + (best_cost * uniquenessRatio/100); - - bool notUniq = ( (cost[d*w] <= thresh) && (d < (ndisp - best_disp - 2) || d > (ndisp - best_disp) ) ); + bool notUniq = ( (cost[0] <= thresh) && (d < (best_disp_back - 1) || d > (best_disp_back + 1) ) ); if(notUniq) - dispbuf[0] = FILTERED; + *bestCost = FILTERED; barrier(CLK_LOCAL_MEM_FENCE); - // best_disp = (textsum < textureTreshold) ? FILTERED : best_disp; - if( dispbuf[0] != FILTERED && x < cols-wsz2-mindisp && y < rows-wsz2) + if( *bestCost != FILTERED && x < cols-wsz2-mindisp && y < rows-wsz2 && d == best_disp_back) { - cost = &cost[0] + (ndisp - best_disp - 1)*w; - int y3 = ((ndisp - best_disp - 1) > 0) ? cost[-w] : cost[w], + int y3 = (best_disp_back > 0) ? cost[-w] : cost[w], y2 = cost[0], - y1 = ((ndisp - best_disp - 1) < ndisp-1) ? cost[w] : cost[-w]; + y1 = (best_disp_back < ndisp-1) ? cost[w] : cost[-w]; int d = y3+y1-2*y2 + abs(y3-y1); - disp[0] = (short)(((ndisp - best_disp - 1 + mindisp)*256 + (d != 0 ? (y3-y1)*256/d : 0) + 15) >> 4); + disp[0] = (short)(((best_disp_back + mindisp)*256 + (d != 0 ? (y3-y1)*256/d : 0) + 15) >> 4); } } @@ -115,11 +95,11 @@ short calcCostBorder(__global const uchar * leftptr, __global const uchar * righ int idx = mad24(y+wsz2*(2*nthread-1), cols, x+wsz2*(1-2*nthread)); left = leftptr + idx; right = rightptr + (idx - d); + int shift = 1*nthread + cols*(1-nthread); short costdiff = 0; for(int i = 0; i < winsize; i++) { - int shift = 1*nthread + cols*(1-nthread); costdiff += abs( left[0] - right[0] ); left += shift; right += shift; @@ -131,7 +111,7 @@ short calcCostBorder(__global const uchar * leftptr, __global const uchar * righ } short calcCostInside(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y, - int wsz2, int cols, int d, short cost_up_left, short cost_up, short cost_left) + int wsz2, int cols, int d, short cost_up_left, short cost_up, short cost_left, int winsize) { __global const uchar * left, * right; int idx = mad24(y-wsz2-1, cols, x-wsz2-1); @@ -139,9 +119,9 @@ short calcCostInside(__global const uchar * leftptr, __global const uchar * righ right = rightptr + (idx - d); uchar corrner1 = abs(left[0] - right[0]), - corrner2 = abs(left[wsz] - right[wsz]), - corrner3 = abs(left[(wsz)*cols] - right[(wsz)*cols]), - corrner4 = abs(left[(wsz)*cols + wsz] - right[(wsz)*cols + wsz]); + corrner2 = abs(left[winsize] - right[winsize]), + corrner3 = abs(left[(winsize)*cols] - right[(winsize)*cols]), + corrner4 = abs(left[(winsize)*cols + winsize] - right[(winsize)*cols + winsize]); return cost_up + cost_left - cost_up_left + corrner1 - corrner2 - corrner3 + corrner4; @@ -162,9 +142,11 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar __global short * disp; __global const uchar * left, * right; - __local short dispbuf[tsize]; __local short costFunc[csize]; __local short * cost; + __local int best_disp[2]; + __local int best_cost[2]; + best_cost[nthread] = MAX_VAL; short costbuf[wsz]; int head = 0; @@ -180,18 +162,19 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar short tempcost = 0; if(x < cols-wsz2-mindisp && y < rows-wsz2) { - for(int i = 0; i < wsz; i++) + int shift = 1*nthread + cols*(1-nthread); + for(int i = 0; i < winsize; i++) { int idx = mad24(y-wsz2+i*nthread, cols, x-wsz2+i*(1-nthread)); left = leftptr + idx; right = rightptr + (idx - d); short costdiff = 0; - for(int j = 0; j < wsz; j++) + for(int j = 0; j < winsize; j++) { costdiff += abs( left[0] - right[0] ); - left += 1*nthread + cols*(1-nthread); - right += 1*nthread + cols*(1-nthread);// maybe use ? operator + left += shift; + right += shift; } if(nthread==1) { @@ -203,42 +186,55 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar } if(nthread==1) { - cost[0] = tempcost; + cost[0] = tempcost; + atomic_min(best_cost+nthread, tempcost); } barrier(CLK_LOCAL_MEM_FENCE); + if(best_cost[1] == tempcost) + best_disp[1] = ndisp - d - 1; + barrier(CLK_LOCAL_MEM_FENCE); + int dispIdx = mad24(gy, disp_step, disp_offset + gx*(int)sizeof(short)); disp = (__global short *)(dispptr + dispIdx); - calcDisp(&costFunc[sizeY - 1 + lx - ly], disp, uniquenessRatio, //textureTreshold, textsum, - mindisp, ndisp, 2*sizeY, &dispbuf[nthread*tsize/2], d, x, y, cols, rows, wsz2); + calcDisp(cost, disp, uniquenessRatio, //textureTreshold, textsum, + mindisp, ndisp, 2*sizeY, best_disp + 1, best_cost+1, d, x, y, cols, rows, wsz2); barrier(CLK_LOCAL_MEM_FENCE); lx = 1 - nthread; ly = nthread; - for(int i = 0; i < iters; i++) + for(int i = 0; i < sizeY*sizeX/2; i++) { x = (lx < sizeX) ? gx + shiftX + lx : cols; y = (ly < sizeY) ? gy + shiftY + ly : rows; + best_cost[nthread] = MAX_VAL; + barrier(CLK_LOCAL_MEM_FENCE); + costIdx = calcLocalIdx(lx, ly, d, sizeY); cost = costFunc + costIdx; + if(x < cols-wsz2-mindisp && y < rows-wsz2 ) { - cost[0] = ( ly*(1-nthread) + lx*nthread == 0 ) ? + tempcost = ( ly*(1-nthread) + lx*nthread == 0 ) ? calcCostBorder(leftptr, rightptr, x, y, nthread, wsz2, costbuf, &head, cols, d, - costFunc[calcLocalIdx(lx-1*(1-nthread), ly-1*nthread, d, sizeY)], winsize) : + cost[2*nthread-1], winsize) : calcCostInside(leftptr, rightptr, x, y, wsz2, cols, d, - costFunc[calcLocalIdx(lx-1, ly-1, d, sizeY)], - costFunc[calcLocalIdx(lx, ly-1, d, sizeY)], - costFunc[calcLocalIdx(lx-1, ly, d, sizeY)]); + cost[0], cost[1], cost[-1], winsize); + cost[0] = tempcost; + atomic_min(best_cost + nthread, tempcost); } barrier(CLK_LOCAL_MEM_FENCE); + if(best_cost[nthread] == tempcost) + best_disp[nthread] = ndisp - d - 1; + barrier(CLK_LOCAL_MEM_FENCE); + int dispIdx = mad24(gy+ly, disp_step, disp_offset + (gx+lx)*(int)sizeof(short)); disp = (__global short *)(dispptr + dispIdx); - calcDisp(&costFunc[sizeY - 1 - ly + lx], disp, uniquenessRatio, //textureTreshold, textsum, - mindisp, ndisp, 2*sizeY, &dispbuf[nthread*tsize/2], d, x, y, cols, rows, wsz2); + calcDisp(cost, disp, uniquenessRatio, //textureTreshold, textsum, + mindisp, ndisp, 2*sizeY, best_disp + nthread, best_cost + nthread, d, x, y, cols, rows, wsz2); barrier(CLK_LOCAL_MEM_FENCE); calcNewCoordinates(&lx, &ly, nthread); diff --git a/modules/calib3d/src/stereobm.cpp b/modules/calib3d/src/stereobm.cpp index 4d28b57d44..4d17b473da 100644 --- a/modules/calib3d/src/stereobm.cpp +++ b/modules/calib3d/src/stereobm.cpp @@ -743,9 +743,9 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right, int wsz = state->SADWindowSize; int wsz2 = wsz/2; - int sizeX = 13, sizeY = sizeX-1, N = ndisp*2; + int sizeX = std::max(11, 27 - ocl::Device::getDefault().maxComputeUnits() ), sizeY = sizeX-1, N = ndisp*2; - ocl::Kernel k("stereoBM_opt", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D tsize=%d -D wsz=%d -D iters=%d", (2*sizeY)*ndisp, 2*ndisp, wsz, sizeX*sizeY/2) ); + ocl::Kernel k("stereoBM_opt", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D wsz=%d", (2*sizeY)*ndisp, wsz) ); if(k.empty()) return false; diff --git a/modules/calib3d/test/opencl/test_stereobm.cpp b/modules/calib3d/test/opencl/test_stereobm.cpp index 3acb57e373..16b3a2a8de 100644 --- a/modules/calib3d/test/opencl/test_stereobm.cpp +++ b/modules/calib3d/test/opencl/test_stereobm.cpp @@ -94,8 +94,8 @@ OCL_TEST_P(StereoBMFixture, StereoBM) Mat t; absdiff(disp, udisp, t); for(int i = 0; i(i,j) > 0) - // if(i == 5 && j == 68) + // if(t.at(i,j) > 0) + if(i == 5 && j == 36) printf("%d %d cv: %d ocl: %d\n", i, j, disp.at(i,j), udisp.getMat(ACCESS_READ).at(i,j) ); /* imshow("diff.png", t*100); imshow("cv.png", disp*100);