optimizations

diff --git a/modules/calib3d/perf/opencl/perf_stereobm.cpp b/modules/calib3d/perf/opencl/perf_stereobm.cpp
index dd2bc9e..b795a35 100644 (file)
--- 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<int, int> StereoBMFixture_t;
 typedef TestBaseWithParam<StereoBMFixture_t> 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<StereoBM> 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 7ab58df..2e74f59 100644 (file)
--- 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<cols && y<rows && d<ndisp)
+    __local short costFunc[csize];
+    short textsum[tsize];
+    __local short * cost = &costFunc[0] + d + ndisp*ndisp/nthreads*z;
+    __global uchar * left, * right;
+    int dispIdx = mad24(total_y, disp_step, disp_offset + x*(int)sizeof(short) );
+    __global short * disp = (__global short*)(dispptr + dispIdx);
+    if( x < cols && total_y < rows)
     {
-        int dispIdx = mad24(y, disp_step, disp_offset + total_x*(int)sizeof(short) );
-        __global short * disp = (__global short*)(dispptr + dispIdx);
         disp[0] = FILTERED;
+    }
+
+    short costbuf[wsz];
+    short textbuf[wsz];
+    int head = 0;
 
-        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<wsz2) ? wsz2 : y;
+        cost[0] = 0;
+        textsum[y-(gy*ndisp)] = 0;
+        for(int i = -wsz2; (i < wsz2+1) && (y < rows-wsz2); i++)
         {
-            for(; (x <= ndisp+mindisp+wsz2-2); x++)
-            {
-                cost[(d-mindisp)+ndisp*(x%(gx*ndisp))] = INT_MAX;
-                textsum[x%(gx*ndisp)] = INT_MAX;
-            }
-            cost[(d-mindisp)+ndisp*(x%(gx*ndisp))] = 0;
-            textsum[x%(gx*ndisp)] = 0;
-            for(int i = -wsz2; i < wsz2+1; i++)
-                for(int j = -wsz2; j < wsz2+1; j++)
-                {
-                    cost[(d-mindisp)+ndisp*(x%(gx*ndisp))] += 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[x%(gx*ndisp)] += abs( left[min( y+i, rows-1 ) * cols + min( x+j, cols-1 )] - preFilterCap );
-                }
-            x++;
-            for(; (x < gx*ndisp + ndisp) && (x < cols-wsz2-mindisp); x++)
-            {
-                cost[(d-mindisp)+ndisp*(x%(gx*ndisp))] = cost[(d-mindisp)+ndisp*((x-1)%(gx*ndisp))];
-                textsum[x%(gx*ndisp)] = textsum[(x-1)%(gx*ndisp)];
-                for(int i = -wsz2; i < wsz2+1; i++)
-                {
-                    cost[(d-mindisp)+ndisp*(x%(gx*ndisp))] += -abs( left[min( y+i, rows-1 ) * cols + min( x-wsz2-1, cols-1 )]
-                            - right[min( y+i, rows-1 ) * cols + min( x-wsz2-1-d, cols-1 )] ) +
-                        abs( left[min( y+i, rows-1 ) * cols + min( x+wsz2, cols-1 )]
-                            - right[min( y+i, rows-1 ) * cols + min( x+wsz2-d, cols-1 )] );
-                    textsum[x%(gx*ndisp)] += -abs( left[min( y+i, rows-1 ) * cols + min( x-wsz2-1, cols-1 )] - preFilterCap ) +
-                        abs( left[min( y+i, rows-1 ) * cols + min( x+wsz2, cols-1 )] - preFilterCap );
-                }
-            }
+            left = leftptr + mad24(y+i, cols, x-wsz2);
+            right = rightptr + mad24(y+i, cols, x-wsz2-d-mindisp);
 
-            for(; (x > 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<cols-1) ? x+1 : x-1)] * (1) +
-                  input[              (y)     * cols + (x-1)] * (-2) + input[       (y)            * cols + ((x<cols-1) ? x+1 : x-1)] * (2) +
-                  input[             (y+1)    * cols + (x-1)] * (-1) + input[      (y+1)           * cols + ((x<cols-1) ? x+1 : x-1)] * (1);
+        int cov = input[ ((y > 0) ? y-1 : y+1)  * cols + (x-1)] * (-1) + input[ ((y > 0) ? y-1 : y+1)  * cols + ((x<cols-1) ? x+1 : x-1)] * (1) +
+                  input[              (y)       * cols + (x-1)] * (-2) + input[        (y)             * cols + ((x<cols-1) ? x+1 : x-1)] * (2) +
+                  input[((y<rows-1)?(y+1):(y-1))* cols + (x-1)] * (-1) + input[((y<rows-1)?(y+1):(y-1))* cols + ((x<cols-1) ? x+1 : x-1)] * (1);
 
-        cov = min(min(max(-prefilterCap, cov), prefilterCap) + prefilterCap, 255);
+        cov = min(clamp(cov, -prefilterCap, prefilterCap) + prefilterCap, 255);
         output[y * cols + x] = cov & 0xFF;
     }
 }

diff --git a/modules/calib3d/src/stereobm.cpp b/modules/calib3d/src/stereobm.cpp
index 05652f0..876215b 100644 (file)
--- a/modules/calib3d/src/stereobm.cpp
+++ b/modules/calib3d/src/stereobm.cpp
@@ -180,13 +180,11 @@ static bool ocl_prefilter_xsobel(InputArray _input, OutputArray _output, int pre
     _output.create(input.size(), input.type());
     output = _output.getUMat();
 
-    size_t blockSize = 1;
     size_t globalThreads[3] = { input.cols, input.rows, 1 };
-    size_t localThreads[3]  = { blockSize, blockSize, 1 };
 
     k.args(ocl::KernelArg::PtrReadOnly(input), ocl::KernelArg::PtrWriteOnly(output), input.rows, input.cols, prefilterCap);
 
-    return k.run(2, globalThreads, localThreads, false);
+    return k.run(2, globalThreads, NULL, false);
 }
 
 static void
@@ -655,6 +653,7 @@ findStereoCorrespondenceBM( const Mat& left, const Mat& right,
                     mind = d;
                 }
             }
+
             tsum += htext[y + wsz2] - htext[y - wsz2 - 1];
             if( tsum < textureThreshold )
             {
@@ -738,9 +737,9 @@ struct PrefilterInvoker : public ParallelLoopBody
 
 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", 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_<short>) + 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 15fa93a..a683e69 100644 (file)
--- 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<StereoBM> 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<t.rows; i++)
+        for(int j = 0; j< t.cols; j++)
+       //    if(t.at<short>(i,j) > 0)
+            if(i>=5 && i <=16 && j == 36+15)
+                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);
 }
 
 OCL_INSTANTIATE_TEST_CASE_P(StereoMatcher, StereoBMFixture, testing::Combine(testing::Values(32, 64, 128),