fixes

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,6 +180,8 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar
     cost = costFunc + costIdx;
 
     short tempcost = 0;
+    if(x < cols-wsz2-mindisp && y < rows-wsz2)
+    {
         for(int i = 0; i < wsz; i++)
         {
             int idx = mad24(y-wsz2+i*nthread, cols, x-wsz2+i*(1-nthread));
@@ -199,28 +202,28 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar
             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);
+        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;
+        x = (lx < sizeX) ? gx + shiftX + lx : cols;
-        y = gy + shiftY + ly;
+        y = (ly < sizeY) ? gy + shiftY + ly : rows;
 
         costIdx = calcLocalIdx(lx, ly, d, sizeY);
         cost = costFunc + costIdx;
+        if(x < cols-wsz2-mindisp && y < rows-wsz2 )
+        {
             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)]) :
@@ -228,19 +231,20 @@ __kernel void stereoBM_opt(__global const uchar * leftptr, __global const uchar
                     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)
-        {
         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);
+            mindisp, ndisp, 2*sizeY, &dispbuf[nthread*tsize/2], d, x, y, cols, rows, wsz2);
-        }
         barrier(CLK_LOCAL_MEM_FENCE);
 
         calcNewCoordinates(&lx, &ly, nthread);
     }
+
+
+
 }
 
 #endif

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<t.rows; i++)
         for(int j = 0; j< t.cols; j++)
            if(t.at<short>(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<short>(i,j), udisp.getMat(ACCESS_READ).at<short>(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),