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

temp

Browse Source
This commit is contained in:
Konstantin Matskevich 2014-03-12 13:09:57 +04:00
parent 18a59b48ba
commit 91e4334227
4 changed files with 39 additions and 40 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
modules/calib3d/perf/opencl/perf_stereobm.cpp
View File

@ -51,7 +51,7 @@ namespace ocl {
typedef std::tr1::tuple<int, int> StereoBMFixture_t; typedef std::tr1::tuple<int, int> StereoBMFixture_t;
typedef TestBaseWithParam<StereoBMFixture_t> StereoBMFixture; typedef TestBaseWithParam<StereoBMFixture_t> StereoBMFixture;
OCL_PERF_TEST_P(StereoBMFixture, StereoBM, ::testing::Combine(OCL_PERF_ENUM(32, 64), OCL_PERF_ENUM(11,21) ) ) OCL_PERF_TEST_P(StereoBMFixture, StereoBM, ::testing::Combine(OCL_PERF_ENUM(32, 64, 128), OCL_PERF_ENUM(11,21) ) )
{ {
const int n_disp = get<0>(GetParam()), winSize = get<1>(GetParam()); const int n_disp = get<0>(GetParam()), winSize = get<1>(GetParam());
UMat left, right, disp; UMat left, right, disp;

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

@ -50,14 +50,12 @@
#define MAX_VAL 32767 #define MAX_VAL 32767
void calcDisp(__local short * costFunc, __global short * disp, int uniquenessRatio/*, int textureTreshold, short textsum*/, 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 short * dispbuf, int d, int x, int y, int cols, int rows, int wsz2)
{ {
short FILTERED = (mindisp - 1)<<4; short FILTERED = (mindisp - 1)<<4;
short best_disp = FILTERED, best_cost = MAX_VAL-1; short best_disp, best_cost;
__local short * cost;
cost = &costFunc[0];
dispbuf[d] = d; dispbuf[d] = d;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
@ -72,28 +70,23 @@ void calcDisp(__local short * costFunc, __global short * disp, int uniquenessRat
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
best_disp = ndisp - dispbuf[0] - 1; best_disp = ndisp - dispbuf[0] - 1;
best_cost = costFunc[(ndisp-best_disp-1)*w]; best_cost = cost[(ndisp-best_disp-1)*w];
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
int thresh = best_cost + (best_cost * uniquenessRatio/100); 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;
bool notUniq = ( (cost[d*w] <= thresh) && (d < (ndisp - best_disp - 2) || d > (ndisp - best_disp) ) );
if(notUniq)
dispbuf[0] = FILTERED;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
for(int lsize = tsize/2 >> 1; lsize > 0; lsize >>= 1)
{
short val1 = dispbuf[d], val2 = dispbuf[d+lsize];
if (d < lsize)
{
dispbuf[d] = min(val1, val2);
}
barrier(CLK_LOCAL_MEM_FENCE);
}
// best_disp = (textsum < textureTreshold) ? FILTERED : best_disp; // best_disp = (textsum < textureTreshold) ? FILTERED : best_disp;
if( dispbuf[0] != FILTERED && x < cols-wsz2-mindisp && y < rows-wsz2) if( dispbuf[0] != FILTERED && x < cols-wsz2-mindisp && y < rows-wsz2)
{ {
cost = &costFunc[0] + (ndisp - best_disp - 1)*w; cost = &cost[0] + (ndisp - best_disp - 1)*w;
int y3 = ((ndisp - best_disp - 1) > 0) ? cost[-w] : cost[w], int y3 = ((ndisp - best_disp - 1) > 0) ? cost[-w] : cost[w],
y2 = cost[0], y2 = cost[0],
y1 = ((ndisp - best_disp - 1) < ndisp-1) ? cost[w] : cost[-w]; y1 = ((ndisp - best_disp - 1) < ndisp-1) ? cost[w] : cost[-w];
@ -115,7 +108,7 @@ void calcNewCoordinates(int * x, int * y, int nthread)
} }
short calcCostBorder(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y, int nthread, short calcCostBorder(__global const uchar * leftptr, __global const uchar * rightptr, int x, int y, int nthread,
int wsz2, short * costbuf, int * h, int cols, int d, short cost) int wsz2, short * costbuf, int * h, int cols, int d, short cost, int winsize)
{ {
int head = (*h)%wsz; int head = (*h)%wsz;
__global const uchar * left, * right; __global const uchar * left, * right;
@ -124,11 +117,12 @@ short calcCostBorder(__global const uchar * leftptr, __global const uchar * righ
right = rightptr + (idx - d); right = rightptr + (idx - d);
short costdiff = 0; short costdiff = 0;
for(int i = 0; i < wsz; i++) for(int i = 0; i < winsize; i++)
{ {
costdiff += abs( left[0] - right[0] ); int shift = 1*nthread + cols*(1-nthread);
left += 1*nthread + cols*(1-nthread); costdiff += abs( left[0] - right[0] );
right += 1*nthread + cols*(1-nthread);// maybe use ? operator left += shift;
right += shift;
} }
cost += costdiff - costbuf[head]; cost += costdiff - costbuf[head];
costbuf[head] = costdiff; costbuf[head] = costdiff;
@ -144,21 +138,25 @@ short calcCostInside(__global const uchar * leftptr, __global const uchar * righ
left = leftptr + idx; left = leftptr + idx;
right = rightptr + (idx - d); right = rightptr + (idx - d);
return cost_up + cost_left - cost_up_left + abs(left[0] - right[0]) - uchar corrner1 = abs(left[0] - right[0]),
abs(left[wsz] - right[wsz]) - abs(left[(wsz)*cols] - right[(wsz)*cols]) + corrner2 = abs(left[wsz] - right[wsz]),
abs(left[(wsz)*cols + wsz] - right[(wsz)*cols + wsz]); corrner3 = abs(left[(wsz)*cols] - right[(wsz)*cols]),
corrner4 = abs(left[(wsz)*cols + wsz] - right[(wsz)*cols + wsz]);
return cost_up + cost_left - cost_up_left + corrner1 -
corrner2 - corrner3 + corrner4;
} }
__kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar * rightptr, __global uchar * dispptr, __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 disp_step, int disp_offset, int rows, int cols, int mindisp, int ndisp,
int preFilterCap, int textureTreshold, int uniquenessRatio, int sizeX, int sizeY) int preFilterCap, int textureTreshold, int uniquenessRatio, int sizeX, int sizeY, int winsize)
{ {
int gx = get_global_id(0)*sizeX; int gx = get_global_id(0)*sizeX;
int gy = get_global_id(1)*sizeY; int gy = get_global_id(1)*sizeY;
int lz = get_local_id(2); int lz = get_local_id(2);
int nthread = lz/32;// only 0 or 1 int nthread = lz/ndisp;// only 0 or 1
int d = lz%32;// 1 .. 32 int d = lz%ndisp;// 1 .. ndisp
int wsz2 = wsz/2; int wsz2 = wsz/2;
__global short * disp; __global short * disp;
@ -203,19 +201,22 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar
head++; head++;
} }
} }
barrier(CLK_LOCAL_MEM_FENCE); if(nthread==1)
{
cost[0] = tempcost; cost[0] = tempcost;
}
barrier(CLK_LOCAL_MEM_FENCE);
int dispIdx = mad24(gy, disp_step, disp_offset + gx*(int)sizeof(short)); int dispIdx = mad24(gy, disp_step, disp_offset + gx*(int)sizeof(short));
disp = (__global short *)(dispptr + dispIdx); disp = (__global short *)(dispptr + dispIdx);
calcDisp(&costFunc[sizeY - 1 + lx - ly], disp, uniquenessRatio, /*textureTreshold, textsum,*/ calcDisp(&costFunc[sizeY - 1 + lx - ly], disp, uniquenessRatio, //textureTreshold, textsum,
mindisp, ndisp, 2*sizeY, &dispbuf[nthread*tsize/2], d, x, y, cols, rows, wsz2); mindisp, ndisp, 2*sizeY, &dispbuf[nthread*tsize/2], d, x, y, cols, rows, wsz2);
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
lx = 1 - nthread; lx = 1 - nthread;
ly = nthread; ly = nthread;
while(lx < sizeX || ly < sizeY ) for(int i = 0; i < iters; i++)
{ {
x = (lx < sizeX) ? gx + shiftX + lx : cols; x = (lx < sizeX) ? gx + shiftX + lx : cols;
y = (ly < sizeY) ? gy + shiftY + ly : rows; y = (ly < sizeY) ? gy + shiftY + ly : rows;
@ -226,7 +227,7 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar
{ {
cost[0] = ( ly*(1-nthread) + lx*nthread == 0 ) ? cost[0] = ( ly*(1-nthread) + lx*nthread == 0 ) ?
calcCostBorder(leftptr, rightptr, x, y, nthread, wsz2, costbuf, &head, cols, d, calcCostBorder(leftptr, rightptr, x, y, nthread, wsz2, costbuf, &head, cols, d,
costFunc[calcLocalIdx(lx-1*(1-nthread), ly-1*nthread, d, sizeY)]) : costFunc[calcLocalIdx(lx-1*(1-nthread), ly-1*nthread, d, sizeY)], winsize) :
calcCostInside(leftptr, rightptr, x, y, wsz2, cols, d, calcCostInside(leftptr, rightptr, x, y, wsz2, cols, d,
costFunc[calcLocalIdx(lx-1, ly-1, d, sizeY)], costFunc[calcLocalIdx(lx-1, ly-1, d, sizeY)],
costFunc[calcLocalIdx(lx, ly-1, d, sizeY)], costFunc[calcLocalIdx(lx, ly-1, d, sizeY)],
@ -242,9 +243,6 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar
calcNewCoordinates(&lx, &ly, nthread); calcNewCoordinates(&lx, &ly, nthread);
} }
} }
#endif #endif

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

@ -743,9 +743,9 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right,
int wsz = state->SADWindowSize; int wsz = state->SADWindowSize;
int wsz2 = wsz/2; int wsz2 = wsz/2;
int sizeX = 9, sizeY = sizeX-1, N = ndisp*2; int sizeX = 13, 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", (2*sizeY)*ndisp, 2*ndisp, wsz) ); 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) );
if(k.empty()) if(k.empty())
return false; return false;
@ -754,7 +754,7 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right,
_disp.create(_left.size(), CV_16S); _disp.create(_left.size(), CV_16S);
_disp.setTo((mindisp - 1)<<4); _disp.setTo((mindisp - 1)<<4);
Rect roi = Rect(Point(wsz2 + mindisp + ndisp - 1, wsz2), Point(cols-1-wsz2-mindisp, rows-1-wsz2) ); Rect roi = Rect(Point(wsz2 + mindisp + ndisp - 1, wsz2), Point(cols-wsz2-mindisp, rows-wsz2) );
UMat disp = (_disp.getUMat())(roi); UMat disp = (_disp.getUMat())(roi);
int globalX = disp.cols/sizeX, globalY = disp.rows/sizeY; int globalX = disp.cols/sizeX, globalY = disp.rows/sizeY;
@ -776,6 +776,7 @@ static bool ocl_stereobm_opt( InputArray _left, InputArray _right,
idx = k.set(idx, state->uniquenessRatio); idx = k.set(idx, state->uniquenessRatio);
idx = k.set(idx, sizeX); idx = k.set(idx, sizeX);
idx = k.set(idx, sizeY); idx = k.set(idx, sizeY);
idx = k.set(idx, wsz);
return k.run(3, globalThreads, localThreads, false); return k.run(3, globalThreads, localThreads, false);
} }

4
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== 12 && j == 44) // if(i == 5 && j == 68)
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);