diff --git a/modules/calib3d/perf/opencl/perf_stereobm.cpp b/modules/calib3d/perf/opencl/perf_stereobm.cpp index dd2bc9e0a9..b795a35265 100644 --- a/modules/calib3d/perf/opencl/perf_stereobm.cpp +++ b/modules/calib3d/perf/opencl/perf_stereobm.cpp @@ -51,7 +51,7 @@ namespace ocl { typedef std::tr1::tuple StereoBMFixture_t; typedef TestBaseWithParam StereoBMFixture; -OCL_PERF_TEST_P(StereoBMFixture, StereoBM, ::testing::Combine(OCL_PERF_ENUM(32, 64, 128), OCL_PERF_ENUM(11,21) ) ) +OCL_PERF_TEST_P(StereoBMFixture, StereoBM, ::testing::Combine(OCL_PERF_ENUM(32, 64), OCL_PERF_ENUM(11,21) ) ) { const int n_disp = get<0>(GetParam()), winSize = get<1>(GetParam()); UMat left, right, disp; @@ -64,11 +64,11 @@ OCL_PERF_TEST_P(StereoBMFixture, StereoBM, ::testing::Combine(OCL_PERF_ENUM(32, declare.in(left, right); Ptr bm = createStereoBM( n_disp, winSize ); - bm->setPreFilterType(bm->PREFILTER_NORMALIZED_RESPONSE); + bm->setPreFilterType(bm->PREFILTER_XSOBEL); OCL_TEST_CYCLE() bm->compute(left, right, disp); - SANITY_CHECK(disp, 0.05, ERROR_RELATIVE); + SANITY_CHECK(disp, 1e-3, ERROR_RELATIVE); } }//ocl diff --git a/modules/calib3d/src/opencl/stereobm.cl b/modules/calib3d/src/opencl/stereobm.cl index 7ab58dfec1..2e74f591cf 100644 --- a/modules/calib3d/src/opencl/stereobm.cl +++ b/modules/calib3d/src/opencl/stereobm.cl @@ -47,90 +47,119 @@ #ifdef csize -__kernel void stereoBM_opt(__global const uchar * left, __global const uchar * right, __global uchar * dispptr, +#define MAX_VAL 32767 + +__kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar * rightptr, __global uchar * dispptr, int disp_step, int disp_offset, int rows, int cols, int mindisp, int ndisp, - int preFilterCap, int winsize, int textureTreshold, int uniquenessRatio) + int preFilterCap, int nthreads, int textureTreshold, int uniquenessRatio) { - int total_x = get_global_id(0); - int gx = get_group_id(0), x = gx*ndisp; - int y = get_global_id(1); - int d = get_local_id(0) + mindisp; - int wsz2 = winsize/2; + int x = get_global_id(0); + int total_y = get_global_id(1); + int z = get_local_id(2); + int d = get_local_id(1); + int gy = get_group_id(1), y = gy*ndisp + z*ndisp/nthreads; + int wsz2 = wsz/2; short FILTERED = (mindisp - 1)<<4; - __local int cost[csize]; - int textsum[tsize]; - if( total_x ndisp-1) && (y > wsz2-1) && (total_x < cols + ndisp - cols%ndisp) && (y < rows - wsz2)) + if( (x > ndisp+mindisp+wsz2-2) && (x < cols - wsz2 - mindisp) ) + { + cost += (y < wsz2) ? ndisp*wsz2 : 0; + y = (y cols - (cols-1)%ndisp - 1) && (x < cols + ndisp - (cols-1)%ndisp - 1); x++) + int costdiff = 0, textdiff = 0; + #pragma unroll + for(int j = 0; j < wsz; j++) { - cost[(d-mindisp)+ndisp*(x%(gx*ndisp))] = INT_MAX; - textsum[x%(gx*ndisp)] = INT_MAX; + costdiff += abs( left[0] - right[0] ); + textdiff += abs( left[0] - preFilterCap ); + left++; right++; } - barrier(CLK_LOCAL_MEM_FENCE); + cost[0] += costdiff; + textsum[y-(gy*ndisp)] += textdiff; + costbuf[head] = costdiff; + textbuf[head] = textdiff; + head++; + } + y++; + for(; (y < gy*ndisp + ndisp/nthreads + z*ndisp/nthreads) && (y < rows-wsz2); y++) + { + head = head%wsz; + cost += ndisp; + cost[0] = cost[-ndisp]; + textsum[y-(gy*ndisp)] = textsum[(y-1)-(gy*ndisp)]; + left = leftptr + mad24(y-wsz2-1, cols, x - wsz2); + right = rightptr + mad24(y-wsz2-1, cols, x - wsz2 - d - mindisp); - int best_disp = FILTERED, best_cost = INT_MAX-1; - for(int i = 0; (i < ndisp); i++) + int costdiff = 0, textdiff = 0; + #pragma unroll + for(int i = 0; i < wsz; i++) { - best_cost = (cost[i + ndisp*(d-mindisp)] < best_cost) ? cost[i + ndisp*(d-mindisp)] : best_cost; - best_disp = (best_cost == cost[i + ndisp*(d-mindisp)]) ? i+mindisp : best_disp; + costdiff += + abs( left[wsz*cols] - right[wsz*cols] ); + textdiff += abs( left[wsz*cols] - preFilterCap ); + left++; right++; } + cost[0] += costdiff - costbuf[head]; + textsum[y-(gy*ndisp)] += textdiff - textbuf[head]; + costbuf[head] = costdiff; + textbuf[head] = textdiff; + head++; + } + barrier(CLK_LOCAL_MEM_FENCE); - int thresh = best_cost + (best_cost * uniquenessRatio/100); - for(int i = 0; (i < ndisp) && (uniquenessRatio > 0); i++) - { - best_disp = ( (cost[i + ndisp*(d-mindisp)] <= thresh) && (i < best_disp - mindisp - 1 || i > best_disp - mindisp + 1) ) ? - FILTERED : best_disp; - } + cost = &costFunc[0] + d*ndisp; + short best_disp = FILTERED, best_cost = MAX_VAL-1; + #pragma unroll + for(int i = 0; i < tsize; i++) + { + short c = cost[0]; + best_cost = (c < best_cost) ? c : best_cost; + best_disp = (best_cost == c) ? ndisp - i - 1 : best_disp; + cost++; + } - disp[0] = textsum[d-mindisp] < textureTreshold ? (FILTERED) : (best_disp == FILTERED) ? (short)(best_disp) : (short)(best_disp); + cost = &costFunc[0] + d*ndisp; + int thresh = best_cost + (best_cost * uniquenessRatio/100); + #pragma unroll + for(int i = 0; (i < tsize) && (uniquenessRatio > 0); i++) + { + best_disp = ( (cost[0] <= thresh) && (i < (ndisp - best_disp - 2) || i > (ndisp - best_disp) ) ) ? + FILTERED : best_disp; + cost++; + } - if( best_disp != FILTERED ) + best_disp = (total_y >= rows-wsz2) || (total_y < wsz2) || (textsum[d] < textureTreshold) ? FILTERED : best_disp; + + if( best_disp != FILTERED ) + { + cost = &costFunc[0] + (ndisp - best_disp - 1) + ndisp*d; + int y3 = ((ndisp - best_disp - 1) > 0) ? cost[-1] : cost[1], + y2 = cost[0], + y1 = ((ndisp - best_disp - 1) < ndisp-1) ? cost[1] : cost[-1]; + d = y3+y1-2*y2 + abs(y3-y1); + if( x < cols && total_y < rows) { - int y1 = (best_disp > mindisp) ? cost[(best_disp-mindisp-1) + ndisp*(d-mindisp)] : - cost[(best_disp-mindisp+1) + ndisp*(d-mindisp)], - y2 = cost[(best_disp-mindisp) + ndisp*(d-mindisp)], - y3 = (best_disp < mindisp+ndisp-1) ? cost[(best_disp-mindisp+1) + ndisp*(d-mindisp)] : - cost[(best_disp-mindisp-1) + ndisp*(d-mindisp)]; - float a = (y3 - ((best_disp+1)*(y2-y1) + best_disp*y1 - (best_disp-1)*y2)/(best_disp - (best_disp-1)) )/ - ((best_disp+1)*((best_disp+1) - (best_disp-1) - best_disp) + (best_disp-1)*best_disp); - float b = (y2 - y1)/(best_disp - (best_disp-1)) - a*((best_disp-1)+best_disp); - disp[0] = (y1 == y2 || y3 == y2) ? (short)(best_disp*16) :(short)(-b/(2*a)*16); + disp[0] = (short)(((ndisp - best_disp - 1 + mindisp)*256 + (d != 0 ? (y3-y1)*256/d : 0) + 15) >> 4); } } } @@ -148,7 +177,7 @@ __kernel void stereoBM_BF(__global const uchar * left, __global const uchar * ri int y = get_global_id(1); int wsz2 = winsize/2; short FILTERED = (mindisp - 1)<<4; - + if(x < cols && y < rows ) { int dispIdx = mad24(y, disp_step, disp_offset + x*(int)sizeof(short) ); @@ -161,21 +190,20 @@ __kernel void stereoBM_BF(__global const uchar * left, __global const uchar * ri for(int d = mindisp; d < ndisp+mindisp; d++) { - cost[d-mindisp] = 0; + cost[(ndisp-1) - (d - mindisp)] = 0; for(int i = -wsz2; i < wsz2+1; i++) for(int j = -wsz2; j < wsz2+1; j++) { - textsum += abs( left[min( y+i, rows-1 ) * cols + min( x+j, cols-1 )] - preFilterCap ); - cost[d-mindisp] += abs( left[min( y+i, rows-1 ) * cols + min( x+j, cols-1 )] - - right[min( y+i, rows-1 ) * cols + min( x+j-d, cols-1 )] ); + textsum += (d == mindisp) ? abs( left[ (y+i) * cols + x + j] - preFilterCap ) : 0; + cost[(ndisp-1) - (d - mindisp)] += abs(left[(y+i) * cols + x+j] - right[(y+i) * cols + x+j-d] ); } } - int best_disp = mindisp, best_cost = cost[0]; - for(int d = mindisp; d < ndisp+mindisp; d++) + int best_disp = -1, best_cost = INT_MAX; + for(int d = ndisp + mindisp - 1; d > mindisp-1; d--) { best_cost = (cost[d-mindisp] < best_cost) ? cost[d-mindisp] : best_cost; - best_disp = (best_cost == cost[d-mindisp]) ? d : best_disp; + best_disp = (best_cost == cost[d-mindisp]) ? (d) : best_disp; } int thresh = best_cost + (best_cost * uniquenessRatio/100); @@ -191,10 +219,8 @@ __kernel void stereoBM_BF(__global const uchar * left, __global const uchar * ri int y1 = (best_disp > mindisp) ? cost[best_disp-mindisp-1] : cost[best_disp-mindisp+1], y2 = cost[best_disp-mindisp], y3 = (best_disp < mindisp+ndisp-1) ? cost[best_disp-mindisp+1] : cost[best_disp-mindisp-1]; - float a = (y3 - ((best_disp+1)*(y2-y1) + best_disp*y1 - (best_disp-1)*y2)/(best_disp - (best_disp-1)) )/ - ((best_disp+1)*((best_disp+1) - (best_disp-1) - best_disp) + (best_disp-1)*best_disp); - float b = (y2 - y1)/(best_disp - (best_disp-1)) - a*((best_disp-1)+best_disp); - disp[0] = (y1 == y2 || y2 == y3) ? (short)(best_disp*16) : (short)(-b/(2*a)*16); + int _d = y3+y1-2*y2 + abs(y3-y1); + disp[0] = (short)(((ndisp - (best_disp-mindisp) - 1 + mindisp)*256 + (_d != 0 ? (y3-y1)*256/_d : 0) + 15) >> 4); } } } @@ -221,10 +247,10 @@ __kernel void prefilter_norm(__global unsigned char *input, __global unsigned ch int cov2 = 0; for(int i = -wsz2; i < wsz2+1; i++) for(int j = -wsz2; j < wsz2+1; j++) - cov2 += input[min( max( (y+i),0 ),rows-1 ) * cols + min( max( (x+j),0 ),cols-1 )]; + cov2 += input[clamp(y+i, 0, rows-1) * cols + clamp(x+j, 0, cols-1)]; int res = (cov1*scale_g - cov2*scale_s)>>10; - res = min(min(max(-prefilterCap, res), prefilterCap) + prefilterCap, 255); + res = min(clamp(res, -prefilterCap, prefilterCap) + prefilterCap, 255); output[y * cols + x] = res & 0xFF; } } @@ -240,13 +266,13 @@ __kernel void prefilter_xsobel(__global unsigned char *input, __global unsigned int x = get_global_id(0); int y = get_global_id(1); output[y * cols + x] = min(prefilterCap, 255) & 0xFF; - if(x < cols && y < rows-1 && x > 0) + if(x < cols && y < rows && x > 0 && !((y == rows-1)&(rows%2==1) ) ) { - int cov = input[((y > 0) ? y-1 : y+1) * cols + (x-1)] * (-1) + input[((y > 0) ? y-1 : y+1) * cols + ((x 0) ? y-1 : y+1) * cols + (x-1)] * (-1) + input[ ((y > 0) ? y-1 : y+1) * cols + ((xnumDisparities; - ocl::Kernel k("stereoBM_opt", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D tsize=%d", ndisp*ndisp, ndisp) ); + ocl::Kernel k("stereoBM_opt", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D tsize=%d -D wsz=%d", ndisp*ndisp, ndisp, state->SADWindowSize) ); if(k.empty()) return false; @@ -748,8 +747,9 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right, _disp.create(_left.size(), CV_16S); UMat disp = _disp.getUMat(); - size_t globalThreads[3] = { left.cols, left.rows, 1 }; - size_t localThreads[3] = {ndisp, 1, 1}; + int nthreads = (ndisp <= 64) ? 2 : 4; + size_t globalThreads[3] = { left.cols, (left.rows - left.rows%ndisp + ndisp), nthreads}; + size_t localThreads[3] = {1, ndisp, nthreads}; int idx = 0; idx = k.set(idx, ocl::KernelArg::PtrReadOnly(left)); @@ -758,11 +758,11 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right, idx = k.set(idx, state->minDisparity); idx = k.set(idx, ndisp); idx = k.set(idx, state->preFilterCap); - idx = k.set(idx, state->SADWindowSize); + idx = k.set(idx, nthreads); idx = k.set(idx, state->textureThreshold); idx = k.set(idx, state->uniquenessRatio); - return k.run(2, globalThreads, localThreads, false); + return k.run(3, globalThreads, localThreads, false); } static bool ocl_stereobm_bf(InputArray _left, InputArray _right, @@ -790,15 +790,16 @@ static bool ocl_stereobm_bf(InputArray _left, InputArray _right, idx = k.set(idx, state->uniquenessRatio); return k.run(2, globalThreads, NULL, false); + return false; } static bool ocl_stereo(InputArray _left, InputArray _right, OutputArray _disp, StereoBMParams* state) { - if(ocl::Device::getDefault().localMemSize() > state->numDisparities * state->numDisparities * sizeof(int) ) + if(ocl::Device::getDefault().localMemSize() > state->numDisparities * state->numDisparities * sizeof(short) ) return ocl_stereobm_opt(_left, _right, _disp, state); else - return ocl_stereobm_bf(_left, _right, _disp, state); + return false;//ocl_stereobm_bf(_left, _right, _disp, state); } struct FindStereoCorrespInvoker : public ParallelLoopBody @@ -992,7 +993,7 @@ public: bufSize2 = width*height*(sizeof(Point_) + sizeof(int) + sizeof(uchar)); #if CV_SSE2 - bool useShorts = false;//params.preFilterCap <= 31 && params.SADWindowSize <= 21 && checkHardwareSupport(CV_CPU_SSE2); + bool useShorts = params.preFilterCap <= 31 && params.SADWindowSize <= 21 && checkHardwareSupport(CV_CPU_SSE2); #else const bool useShorts = false; #endif diff --git a/modules/calib3d/test/opencl/test_stereobm.cpp b/modules/calib3d/test/opencl/test_stereobm.cpp index 15fa93aa4f..a683e6938b 100644 --- a/modules/calib3d/test/opencl/test_stereobm.cpp +++ b/modules/calib3d/test/opencl/test_stereobm.cpp @@ -81,11 +81,27 @@ OCL_TEST_P(StereoBMFixture, StereoBM) { Ptr bm = createStereoBM( n_disp, winSize); bm->setPreFilterType(bm->PREFILTER_XSOBEL); +// bm->setMinDisparity(15); + long t1 = clock(); OCL_OFF(bm->compute(left, right, disp)); + long t2 = clock(); OCL_ON(bm->compute(uleft, uright, udisp)); - - Near(0.05); + 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>=5 && i <=16 && j == 36+15) + 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); } OCL_INSTANTIATE_TEST_CASE_P(StereoMatcher, StereoBMFixture, testing::Combine(testing::Values(32, 64, 128),