diff --git a/modules/calib3d/src/opencl/stereobm.cl b/modules/calib3d/src/opencl/stereobm.cl index 63765357fe73f0e7811a276d7f217fe2a3c17efa..bca6dd48b675db51f5fd18c4c59d99b2847dc1f1 100644 --- a/modules/calib3d/src/opencl/stereobm.cl +++ b/modules/calib3d/src/opencl/stereobm.cl @@ -56,7 +56,7 @@ void calcDisp(__local short * costFunc, __global short * disp, int uniquenessRat __local short * cost; cost = &costFunc[0]; #pragma unroll - for(int i = 0; i < tsize/2; i++) + for(int i = 0; i < tsize; i++) { short c = cost[0]; best_cost = (c < best_cost) ? c : best_cost; @@ -67,14 +67,14 @@ void calcDisp(__local short * costFunc, __global short * disp, int uniquenessRat cost = &costFunc[0]; int thresh = best_cost + (best_cost * uniquenessRatio/100); #pragma unroll - for(int i = 0; (i < tsize/2) && (uniquenessRatio > 0); i++) + 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++; } - best_disp = (textsum < textureTreshold) ? FILTERED : best_disp; +// best_disp = (textsum < textureTreshold) ? FILTERED : best_disp; if( best_disp != FILTERED ) { @@ -92,104 +92,73 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar int preFilterCap, int textureTreshold, int uniquenessRatio) { int x = get_global_id(0); - int ly = get_local_id(1); - int y = get_global_id(1)*32; + int total_y = get_global_id(1); int d = get_local_id(2); + int ly = get_local_id(1); + int gy = get_group_id(1), y = gy*wsz; int wsz2 = wsz/2; short FILTERED = (mindisp - 1)<<4; - __local short costFunc[tsize]; - __local short bestdisp[tsize]; + __local short costFunc[csize]; short textsum; - __local short * cost = &costFunc[0] + d +ly*ndisp; - __global uchar * left, * right; - int dispIdx = mad24(y, disp_step, disp_offset + x*(int)sizeof(short) ); + __local short * cost = costFunc + d; + + __global const uchar * left, * right; + int dispIdx = mad24(total_y, disp_step, disp_offset + x*(int)sizeof(short) ); __global short * disp = (__global short*)(dispptr + dispIdx); short best_cost = MAX_VAL-1, best_disp = FILTERED; short costbuf[wsz]; - short textbuf[wsz]; - int head = 0; - - int endy = y+32; - cost[0] = 0; - bestdisp[d + ly*ndisp] = d; - textsum = 0; - - for(; y < wsz2; y++) - { - disp[0] = FILTERED; - disp += cols; - } - if( x < cols && y < rows) + if( x < cols && total_y < rows) { disp[0] = FILTERED; } if( (x > ndisp+mindisp+wsz2-2) && (x < cols - wsz2 - mindisp) ) { - for(int i = -wsz2; (i < wsz2+1) && (y < rows-wsz2); i++) - { - left = leftptr + mad24(y+i, cols, x-wsz2); - right = rightptr + mad24(y+i, cols, x-wsz2-d-mindisp); + cost[ly*ndisp] = 0; + cost += (y < wsz2) ? ndisp*wsz2 : 0; + y = (y= wsz2) ) + for( int i = 0; i < wsz; i++) { - calcDisp(&costFunc[ly*ndisp], &disp[0], uniquenessRatio, textureTreshold, textsum, mindisp, ndisp); + if(ly == i) + cost[0] += costdiff; } barrier(CLK_LOCAL_MEM_FENCE); - } - - y++; - cost = &costFunc[0] + d+ly*ndisp; - - for(; (y < endy) && (y wsz2 && (x > ndisp+mindisp+wsz2-2) && (x < cols - wsz2 - mindisp) ) + y++; + for(; (y < gy*wsz + wsz) && (y < rows-wsz2); y++) { - head = head%wsz; - left = leftptr + mad24(y-wsz2-1, cols, x - wsz2); - right = rightptr + mad24(y-wsz2-1, cols, x - wsz2 - d - mindisp); - - int costdiff = 0, textdiff = 0; - #pragma unroll - for(int i = 0; i < wsz; i++) + cost += ndisp; + left += cols; + right += cols; + costdiff += abs(left[0] - right[0]) - abs(left[(-wsz2-1)*cols] - right[(-wsz2-1)*cols]);//costbuf[(y-1)%wsz]; + for( int i = 0; i < wsz; i++) { - costdiff += - abs( left[wsz*cols] - right[wsz*cols] ); - textdiff += abs( left[wsz*cols] - preFilterCap ); - left++; right++; + if(ly == i) + cost[0] += costdiff; } - cost[0] += costdiff - costbuf[head]; - textsum += textdiff - textbuf[head]; - costbuf[head] = costdiff; - textbuf[head] = textdiff; - head++; barrier(CLK_LOCAL_MEM_FENCE); + } + barrier(CLK_LOCAL_MEM_FENCE); +/* + if(total_y >= wsz2 && total_y < rows - wsz2 && d == 0) + { + cost = costFunc + ly*ndisp; + disp[0] = cost[wsz-1]; + }*/ - if(d == 0) - { - calcDisp(&costFunc[ly*ndisp], &disp[0], uniquenessRatio, textureTreshold, textsum, mindisp, ndisp); - } - barrier(CLK_LOCAL_MEM_FENCE); + if(total_y >= wsz2 && total_y < rows - wsz2 && d == 0) + { + calcDisp(&(costFunc + ly*ndisp)[0], disp, uniquenessRatio, textureTreshold, textsum, mindisp, ndisp); } } } diff --git a/modules/calib3d/src/stereobm.cpp b/modules/calib3d/src/stereobm.cpp index 67864b07dd6ccb7d0c7efadc3010f033abcd82eb..972499e27b160a7562f9de993b1cc37b0c197899 100644 --- a/modules/calib3d/src/stereobm.cpp +++ b/modules/calib3d/src/stereobm.cpp @@ -739,7 +739,8 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right, OutputArray _disp, StereoBMParams* state) {//printf("opt\n"); int ndisp = state->numDisparities; - ocl::Kernel k("stereoBM_opt", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D tsize=%d -D wsz=%d", ndisp*ndisp, 2*ndisp, state->SADWindowSize) ); + int wsz = state->SADWindowSize; + ocl::Kernel k("stereoBM_opt", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D tsize=%d -D wsz=%d", wsz*ndisp, ndisp, wsz) ); if(k.empty()) return false; @@ -747,8 +748,8 @@ 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-left.rows%32 + 32)/32, ndisp}; - size_t localThreads[3] = {1, 2, ndisp}; + size_t globalThreads[3] = { left.cols, (left.rows-left.rows%wsz + wsz), ndisp}; + size_t localThreads[3] = {1, wsz, ndisp}; int idx = 0; idx = k.set(idx, ocl::KernelArg::PtrReadOnly(left)); @@ -797,7 +798,7 @@ static bool ocl_stereo(InputArray _left, InputArray _right, if(ocl::Device::getDefault().localMemSize() > state->numDisparities * state->numDisparities * sizeof(short) ) return ocl_stereobm_opt(_left, _right, _disp, state); else - return false;//ocl_stereobm_bf(_left, _right, _disp, state); + return ocl_stereobm_bf(_left, _right, _disp, state); } struct FindStereoCorrespInvoker : public ParallelLoopBody diff --git a/modules/calib3d/test/opencl/test_stereobm.cpp b/modules/calib3d/test/opencl/test_stereobm.cpp index 1852e0dfde8751cd9d020c613aebffc26e14ccad..b6f777653d1af50560191d777f0bc74a94932564 100644 --- a/modules/calib3d/test/opencl/test_stereobm.cpp +++ b/modules/calib3d/test/opencl/test_stereobm.cpp @@ -92,11 +92,11 @@ OCL_TEST_P(StereoBMFixture, StereoBM) 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 == 125 && j == 174) - printf("%d %d cv: %d ocl: %d\n", i, j, disp.at(i,j), udisp.getMat(ACCESS_READ).at(i,j) );*/ + // if(t.at(i,j) > 0) + if(i == 5 && j == 38) + 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);