generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

bad experiment =(

Browse Source
This commit is contained in:
Konstantin Matskevich 2014-02-28 09:28:07 +04:00
parent 0904f10ab5
commit 799d7e7a50
3 changed files with 99 additions and 72 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

149
modules/calib3d/src/opencl/stereobm.cl
View File

@ -49,38 +49,85 @@
#define MAX_VAL 32767 #define MAX_VAL 32767
__kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar * rightptr, __global uchar * dispptr, void calcDisp(__local short * costFunc, __global short * disp, int uniquenessRatio, int textureTreshold, short textsum, int mindisp, int ndisp)
int disp_step, int disp_offset, int rows, int cols, int mindisp, int ndisp,
int preFilterCap, int nthreads, int textureTreshold, int uniquenessRatio)
{ {
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; short FILTERED = (mindisp - 1)<<4;
__local short costFunc[csize]; short best_disp = FILTERED, best_cost = MAX_VAL-1;
short textsum[tsize]; __local short * cost;
__local short * cost = &costFunc[0] + d + ndisp*ndisp/nthreads*z; cost = &costFunc[0];
__global uchar * left, * right; #pragma unroll
int dispIdx = mad24(total_y, disp_step, disp_offset + x*(int)sizeof(short) ); for(int i = 0; i < tsize/2; i++)
__global short * disp = (__global short*)(dispptr + dispIdx);
if( x < cols && total_y < rows)
{ {
disp[0] = FILTERED; short c = cost[0];
best_cost = (c < best_cost) ? c : best_cost;
best_disp = (best_cost == c) ? ndisp - i - 1 : best_disp;
cost++;
} }
cost = &costFunc[0];
int thresh = best_cost + (best_cost * uniquenessRatio/100);
#pragma unroll
for(int i = 0; (i < tsize/2) && (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;
if( best_disp != FILTERED )
{
cost = &costFunc[0] + (ndisp - best_disp - 1);
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];
int d = y3+y1-2*y2 + abs(y3-y1);
disp[0] = (short)best_disp;//(((ndisp - best_disp - 1 + mindisp)*256 + (d != 0 ? (y3-y1)*256/d : 0) + 15) >> 4);
}
}
__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 textureTreshold, int uniquenessRatio)
{
int x = get_global_id(0);
int ly = get_local_id(1);
int y = get_global_id(1)*32;
int d = get_local_id(2);
int wsz2 = wsz/2;
short FILTERED = (mindisp - 1)<<4;
__local short costFunc[tsize];
__local short bestdisp[tsize];
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) );
__global short * disp = (__global short*)(dispptr + dispIdx);
short best_cost = MAX_VAL-1, best_disp = FILTERED;
short costbuf[wsz]; short costbuf[wsz];
short textbuf[wsz]; short textbuf[wsz];
int head = 0; 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)
{
disp[0] = FILTERED;
}
if( (x > ndisp+mindisp+wsz2-2) && (x < cols - wsz2 - mindisp) ) if( (x > ndisp+mindisp+wsz2-2) && (x < cols - wsz2 - mindisp) )
{ {
cost += (y < wsz2) ? ndisp*wsz2 : 0;
y = (y<wsz2) ? wsz2 : y;
cost[0] = 0;
textsum[y-(gy*ndisp)] = 0;
for(int i = -wsz2; (i < wsz2+1) && (y < rows-wsz2); i++) for(int i = -wsz2; (i < wsz2+1) && (y < rows-wsz2); i++)
{ {
left = leftptr + mad24(y+i, cols, x-wsz2); left = leftptr + mad24(y+i, cols, x-wsz2);
@ -95,18 +142,30 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar
left++; right++; left++; right++;
} }
cost[0] += costdiff; cost[0] += costdiff;
textsum[y-(gy*ndisp)] += textdiff; textsum += textdiff;
costbuf[head] = costdiff; costbuf[head] = costdiff;
textbuf[head] = textdiff; textbuf[head] = textdiff;
head++; head++;
} }
barrier(CLK_LOCAL_MEM_FENCE);
if( (y < rows-wsz2) && (y >= wsz2) )
{
calcDisp(&costFunc[ly*ndisp], &disp[0], uniquenessRatio, textureTreshold, textsum, mindisp, ndisp);
}
barrier(CLK_LOCAL_MEM_FENCE);
}
y++; y++;
for(; (y < gy*ndisp + ndisp/nthreads + z*ndisp/nthreads) && (y < rows-wsz2); y++) cost = &costFunc[0] + d+ly*ndisp;
for(; (y < endy) && (y<rows); y++)
{
disp+=cols;
disp[0] = FILTERED;
if(y < rows - wsz2 && y > wsz2 && (x > ndisp+mindisp+wsz2-2) && (x < cols - wsz2 - mindisp) )
{ {
head = head%wsz; 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); left = leftptr + mad24(y-wsz2-1, cols, x - wsz2);
right = rightptr + mad24(y-wsz2-1, cols, x - wsz2 - d - mindisp); right = rightptr + mad24(y-wsz2-1, cols, x - wsz2 - d - mindisp);
@ -120,47 +179,17 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar
left++; right++; left++; right++;
} }
cost[0] += costdiff - costbuf[head]; cost[0] += costdiff - costbuf[head];
textsum[y-(gy*ndisp)] += textdiff - textbuf[head]; textsum += textdiff - textbuf[head];
costbuf[head] = costdiff; costbuf[head] = costdiff;
textbuf[head] = textdiff; textbuf[head] = textdiff;
head++; head++;
}
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
cost = &costFunc[0] + d*ndisp; if(d == 0)
short best_disp = FILTERED, best_cost = MAX_VAL-1;
#pragma unroll
for(int i = 0; i < tsize; i++)
{ {
short c = cost[0]; calcDisp(&costFunc[ly*ndisp], &disp[0], uniquenessRatio, textureTreshold, textsum, mindisp, ndisp);
best_cost = (c < best_cost) ? c : best_cost;
best_disp = (best_cost == c) ? ndisp - i - 1 : best_disp;
cost++;
}
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++;
}
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)
{
disp[0] = (short)(((ndisp - best_disp - 1 + mindisp)*256 + (d != 0 ? (y3-y1)*256/d : 0) + 15) >> 4);
} }
barrier(CLK_LOCAL_MEM_FENCE);
} }
} }
} }

12
modules/calib3d/src/stereobm.cpp
View File

@ -681,7 +681,7 @@ findStereoCorrespondenceBM( const Mat& left, const Mat& right,
sad[ndisp] = sad[ndisp-2]; sad[ndisp] = sad[ndisp-2];
int p = sad[mind+1], n = sad[mind-1]; int p = sad[mind+1], n = sad[mind-1];
d = p + n - 2*sad[mind] + std::abs(p - n); d = p + n - 2*sad[mind] + std::abs(p - n);
dptr[y*dstep] = (short)(((ndisp - mind - 1 + mindisp)*256 + (d != 0 ? (p-n)*256/d : 0) + 15) >> 4); dptr[y*dstep] = (short)mind;//(((ndisp - mind - 1 + mindisp)*256 + (d != 0 ? (p-n)*256/d : 0) + 15) >> 4);
costptr[y*coststep] = sad[mind]; costptr[y*coststep] = sad[mind];
} }
} }
@ -739,7 +739,7 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right,
OutputArray _disp, StereoBMParams* state) OutputArray _disp, StereoBMParams* state)
{//printf("opt\n"); {//printf("opt\n");
int ndisp = state->numDisparities; 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, ndisp, state->SADWindowSize) ); 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) );
if(k.empty()) if(k.empty())
return false; return false;
@ -747,9 +747,8 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right,
_disp.create(_left.size(), CV_16S); _disp.create(_left.size(), CV_16S);
UMat disp = _disp.getUMat(); UMat disp = _disp.getUMat();
int nthreads = (ndisp <= 64) ? 2 : 4; size_t globalThreads[3] = { left.cols, (left.rows-left.rows%32 + 32)/32, ndisp};
size_t globalThreads[3] = { left.cols, (left.rows - left.rows%ndisp + ndisp), nthreads}; size_t localThreads[3] = {1, 2, ndisp};
size_t localThreads[3] = {1, ndisp, nthreads};
int idx = 0; int idx = 0;
idx = k.set(idx, ocl::KernelArg::PtrReadOnly(left)); idx = k.set(idx, ocl::KernelArg::PtrReadOnly(left));
@ -758,7 +757,6 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right,
idx = k.set(idx, state->minDisparity); idx = k.set(idx, state->minDisparity);
idx = k.set(idx, ndisp); idx = k.set(idx, ndisp);
idx = k.set(idx, state->preFilterCap); idx = k.set(idx, state->preFilterCap);
idx = k.set(idx, nthreads);
idx = k.set(idx, state->textureThreshold); idx = k.set(idx, state->textureThreshold);
idx = k.set(idx, state->uniquenessRatio); idx = k.set(idx, state->uniquenessRatio);
@ -993,7 +991,7 @@ public:
bufSize2 = width*height*(sizeof(Point_<short>) + sizeof(int) + sizeof(uchar)); bufSize2 = width*height*(sizeof(Point_<short>) + sizeof(int) + sizeof(uchar));
#if CV_SSE2 #if CV_SSE2
bool useShorts = params.preFilterCap <= 31 && params.SADWindowSize <= 21 && checkHardwareSupport(CV_CPU_SSE2); bool useShorts = false;//params.preFilterCap <= 31 && params.SADWindowSize <= 21 && checkHardwareSupport(CV_CPU_SSE2);
#else #else
const bool useShorts = false; const bool useShorts = false;
#endif #endif

6
modules/calib3d/test/opencl/test_stereobm.cpp
View File

@ -90,12 +90,12 @@ OCL_TEST_P(StereoBMFixture, StereoBM)
cv::ocl::finish(); cv::ocl::finish();
long t3 = clock(); long t3 = clock();
std::cout << (double)(t2-t1)/CLOCKS_PER_SEC << " " << (double)(t3-t2)/CLOCKS_PER_SEC << std::endl; std::cout << (double)(t2-t1)/CLOCKS_PER_SEC << " " << (double)(t3-t2)/CLOCKS_PER_SEC << std::endl;
/*
Mat t; absdiff(disp, udisp, t); Mat t; absdiff(disp, udisp, t);
/* for(int i = 0; i<t.rows; i++) /* for(int i = 0; i<t.rows; i++)
for(int j = 0; j< t.cols; j++) for(int j = 0; j< t.cols; j++)
// if(t.at<short>(i,j) > 0) if(t.at<short>(i,j) > 0)
if(i>=5 && i <=16 && j == 36+15) // if(i == 125 && j == 174)
printf("%d %d cv: %d ocl: %d\n", i, j, disp.at<short>(i,j), udisp.getMat(ACCESS_READ).at<short>(i,j) );*/ printf("%d %d cv: %d ocl: %d\n", i, j, disp.at<short>(i,j), udisp.getMat(ACCESS_READ).at<short>(i,j) );*/
/* imshow("diff.png", t*100); /* imshow("diff.png", t*100);
imshow("cv.png", disp*100); imshow("cv.png", disp*100);