diff --git a/modules/calib3d/src/opencl/stereobm.cl b/modules/calib3d/src/opencl/stereobm.cl index 7036cdcf36..307b668b33 100644 --- a/modules/calib3d/src/opencl/stereobm.cl +++ b/modules/calib3d/src/opencl/stereobm.cl @@ -51,7 +51,7 @@ #define MAX_VAL 32767 void calcDisp(__local short * costFunc, __global short * disp, int uniquenessRatio/*, int textureTreshold, short textsum*/, - int mindisp, int ndisp, int w, __local short * dispbuf, int d) + int mindisp, int ndisp, int w, __local short * dispbuf, int d, int x, int y, int cols, int rows, int wsz2) { short FILTERED = (mindisp - 1)<<4; short best_disp = FILTERED, best_cost = MAX_VAL-1; @@ -73,6 +73,7 @@ void calcDisp(__local short * costFunc, __global short * disp, int uniquenessRat } best_disp = ndisp - dispbuf[0] - 1; best_cost = costFunc[(ndisp-best_disp-1)*w]; + barrier(CLK_LOCAL_MEM_FENCE); int thresh = best_cost + (best_cost * uniquenessRatio/100); dispbuf[d] = ( (cost[d*w] <= thresh) && (d < (ndisp - best_disp - 2) || d > (ndisp - best_disp) ) ) ? FILTERED : best_disp; @@ -90,7 +91,7 @@ void calcDisp(__local short * costFunc, __global short * disp, int uniquenessRat // best_disp = (textsum < textureTreshold) ? FILTERED : best_disp; - if( dispbuf[0] != FILTERED ) + if( dispbuf[0] != FILTERED && x < cols-wsz2-mindisp && y < rows-wsz2) { cost = &costFunc[0] + (ndisp - best_disp - 1)*w; int y3 = ((ndisp - best_disp - 1) > 0) ? cost[-w] : cost[w], @@ -179,68 +180,71 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar cost = costFunc + costIdx; short tempcost = 0; - for(int i = 0; i < wsz; i++) + if(x < cols-wsz2-mindisp && y < rows-wsz2) { - 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++) - { - costdiff += abs( left[0] - right[0] ); - left += 1*nthread + cols*(1-nthread); - right += 1*nthread + cols*(1-nthread);// maybe use ? operator - } - if(nthread==1) + for(int i = 0; i < wsz; i++) { - tempcost += costdiff; + 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++) + { + costdiff += abs( left[0] - right[0] ); + left += 1*nthread + cols*(1-nthread); + right += 1*nthread + cols*(1-nthread);// maybe use ? operator + } + if(nthread==1) + { + tempcost += costdiff; + } + costbuf[head] = costdiff; + head++; } - costbuf[head] = costdiff; - head++; } barrier(CLK_LOCAL_MEM_FENCE); cost[0] = tempcost; - if(x < cols-wsz2-mindisp && y < rows-wsz2 && nthread == 1) - { - 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); - } + 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); barrier(CLK_LOCAL_MEM_FENCE); lx = 1 - nthread; ly = nthread; - while(lx < sizeX && ly < sizeY ) + while(lx < sizeX || ly < sizeY ) { - x = gx + shiftX + lx; - y = gy + shiftY + ly; + x = (lx < sizeX) ? gx + shiftX + lx : cols; + y = (ly < sizeY) ? gy + shiftY + ly : rows; costIdx = calcLocalIdx(lx, ly, d, sizeY); cost = costFunc + costIdx; - cost[0] = ( 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)]) : - 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)]); - barrier(CLK_LOCAL_MEM_FENCE); - - if(x < cols-mindisp-wsz2 && y < rows-wsz2) + if(x < cols-wsz2-mindisp && y < rows-wsz2 ) { - 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); + cost[0] = ( 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)]) : + 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)]); } 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); + barrier(CLK_LOCAL_MEM_FENCE); + calcNewCoordinates(&lx, &ly, nthread); } + + + } #endif diff --git a/modules/calib3d/test/opencl/test_stereobm.cpp b/modules/calib3d/test/opencl/test_stereobm.cpp index 7b28483d3e..3b034c6200 100644 --- a/modules/calib3d/test/opencl/test_stereobm.cpp +++ b/modules/calib3d/test/opencl/test_stereobm.cpp @@ -90,18 +90,18 @@ OCL_TEST_P(StereoBMFixture, StereoBM) cv::ocl::finish(); long t3 = clock(); std::cout << (double)(t2-t1)/CLOCKS_PER_SEC << " " << (double)(t3-t2)/CLOCKS_PER_SEC << std::endl; -/* + Mat t; absdiff(disp, udisp, t); for(int i = 0; i(i,j) > 0) - // if(i== 255 && j == 375) + // if(i== 12 && j == 44) 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); imshow("ocl.png", udisp.getMat(ACCESS_READ)*100); waitKey(0);*/ -// Near(1e-3); + Near(1e-3); } OCL_INSTANTIATE_TEST_CASE_P(StereoMatcher, StereoBMFixture, testing::Combine(testing::Values(32, 64, 128),