From 4ceaf44fa0645308af3b39eee7cb0d4acedcd830 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Thu, 27 Mar 2014 11:18:30 +0400 Subject: [PATCH] Fixed incorrect calculation of best_disp --- modules/calib3d/src/opencl/stereobm.cl | 15 ++++----------- modules/calib3d/src/stereobm.cpp | 3 +-- 2 files changed, 5 insertions(+), 13 deletions(-) diff --git a/modules/calib3d/src/opencl/stereobm.cl b/modules/calib3d/src/opencl/stereobm.cl index e23cfdd..73402a6 100644 --- a/modules/calib3d/src/opencl/stereobm.cl +++ b/modules/calib3d/src/opencl/stereobm.cl @@ -147,6 +147,7 @@ __kernel void stereoBM(__global const uchar * leftptr, __global const uchar * ri __local int best_disp[2]; __local int best_cost[2]; best_cost[nthread] = MAX_VAL; + best_disp[nthread] = MAX_VAL; barrier(CLK_LOCAL_MEM_FENCE); short costbuf[wsz]; @@ -187,16 +188,12 @@ __kernel void stereoBM(__global const uchar * leftptr, __global const uchar * ri if(nthread==1) { cost[0] = tempcost; -#ifndef CPU atomic_min(best_cost+nthread, tempcost); -#else - *(best_cost+nthread) = min(*(best_cost+nthread), tempcost); -#endif } barrier(CLK_LOCAL_MEM_FENCE); if(best_cost[1] == tempcost) - best_disp[1] = ndisp - d - 1; + atomic_min(best_disp + 1, ndisp - d - 1); barrier(CLK_LOCAL_MEM_FENCE); int dispIdx = mad24(gy, disp_step, disp_offset + gx*(int)sizeof(short)); @@ -214,6 +211,7 @@ __kernel void stereoBM(__global const uchar * leftptr, __global const uchar * ri y = (ly < sizeY) ? gy + shiftY + ly : rows; best_cost[nthread] = MAX_VAL; + best_disp[nthread] = MAX_VAL; barrier(CLK_LOCAL_MEM_FENCE); costIdx = calcLocalIdx(lx, ly, d, sizeY); @@ -228,20 +226,15 @@ __kernel void stereoBM(__global const uchar * leftptr, __global const uchar * ri cost[0], cost[1], cost[-1], winsize); } cost[0] = tempcost; -#ifndef CPU atomic_min(best_cost + nthread, tempcost); -#else - *(best_cost + nthread) = min(*(best_cost + nthread), tempcost); -#endif barrier(CLK_LOCAL_MEM_FENCE); if(best_cost[nthread] == tempcost) - best_disp[nthread] = ndisp - d - 1; + atomic_min(best_disp + nthread, ndisp - d - 1); barrier(CLK_LOCAL_MEM_FENCE); int dispIdx = mad24(gy+ly, disp_step, disp_offset + (gx+lx)*(int)sizeof(short)); disp = (__global short *)(dispptr + dispIdx); - calcDisp(cost, disp, uniquenessRatio, mindisp, ndisp, 2*sizeY, best_disp + nthread, best_cost + nthread, d, x, y, cols, rows, wsz2); barrier(CLK_LOCAL_MEM_FENCE); diff --git a/modules/calib3d/src/stereobm.cpp b/modules/calib3d/src/stereobm.cpp index bd32b4f..7c06deb 100644 --- a/modules/calib3d/src/stereobm.cpp +++ b/modules/calib3d/src/stereobm.cpp @@ -744,9 +744,8 @@ static bool ocl_stereobm( InputArray _left, InputArray _right, int wsz2 = wsz/2; int sizeX = std::max(11, 27 - ocl::Device::getDefault().maxComputeUnits() ), sizeY = sizeX-1, N = ndisp*2; - bool is_cpu = cv::ocl::Device::getDefault().type() == cv::ocl::Device::TYPE_CPU; - ocl::Kernel k("stereoBM", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D wsz=%d%s", (2*sizeY)*ndisp, wsz, is_cpu ? " -D CPU" : "")); + ocl::Kernel k("stereoBM", ocl::calib3d::stereobm_oclsrc, cv::format("-D csize=%d -D wsz=%d", (2*sizeY)*ndisp, wsz) ); if(k.empty()) return false; -- 2.7.4