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

Merge pull request #1736 from alalek:ocl_fix_corner_memory_access

Browse Source
This commit is contained in:
Andrey Pavlenko 2013-11-01 18:37:34 +04:00 committed by OpenCV Buildbot
commit dd942df08b
3 changed files with 30 additions and 24 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

18
modules/ocl/src/opencl/imgproc_calcHarris.cl
View File

@ -119,18 +119,16 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g
__local float temp[6][THREADS]; __local float temp[6][THREADS];
#ifdef BORDER_CONSTANT #ifdef BORDER_CONSTANT
bool dx_con,dy_con;
float dx_s, dy_s;
for (int i=0; i < ksY+1; i++) for (int i=0; i < ksY+1; i++)
{ {
dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows; bool dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows;
dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+col)]; int indexDx = (dx_startY+i)*(dx_step>>2)+(dx_startX+col);
dx_data[i] = dx_con ? dx_s : 0.0f; float dx_s = dx_con ? Dx[indexDx] : 0.0f;
dx_data[i] = dx_s;
dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows; bool dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
dy_s = Dy[(dy_startY+i)*(dy_step>>2)+(dy_startX+col)]; int indexDy = (dy_startY+i)*(dy_step>>2)+(dy_startX+col);
dy_data[i] = dy_con ? dy_s : 0.0f; float dy_s = dx_con ? Dy[indexDy] : 0.0f;
dy_data[i] = dy_s;
data[0][i] = dx_data[i] * dx_data[i]; data[0][i] = dx_data[i] * dx_data[i];
data[1][i] = dx_data[i] * dy_data[i]; data[1][i] = dx_data[i] * dy_data[i];
data[2][i] = dy_data[i] * dy_data[i]; data[2][i] = dy_data[i] * dy_data[i];

16
modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl
View File

@ -118,16 +118,16 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
__local float temp[6][THREADS]; __local float temp[6][THREADS];
#ifdef BORDER_CONSTANT #ifdef BORDER_CONSTANT
bool dx_con, dy_con;
float dx_s, dy_s;
for (int i=0; i < ksY+1; i++) for (int i=0; i < ksY+1; i++)
{ {
dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows; bool dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows;
dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+col)]; int indexDx = (dx_startY+i)*(dx_step>>2)+(dx_startX+col);
dx_data[i] = dx_con ? dx_s : 0.0f; float dx_s = dx_con ? Dx[indexDx] : 0.0f;
dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows; dx_data[i] = dx_s;
dy_s = Dy[(dy_startY+i)*(dy_step>>2)+(dy_startX+col)]; bool dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
dy_data[i] = dy_con ? dy_s : 0.0f; int indexDy = (dy_startY+i)*(dy_step>>2)+(dy_startX+col);
float dy_s = dx_con ? Dy[indexDy] : 0.0f;
dy_data[i] = dy_s;
data[0][i] = dx_data[i] * dx_data[i]; data[0][i] = dx_data[i] * dx_data[i];
data[1][i] = dx_data[i] * dy_data[i]; data[1][i] = dx_data[i] * dy_data[i];
data[2][i] = dy_data[i] * dy_data[i]; data[2][i] = dy_data[i] * dy_data[i];

20
modules/ocl/test/test_imgproc.cpp
View File

@ -93,14 +93,22 @@ PARAM_TEST_CASE(ImgprocTestBase, MatType,
generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder); generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder);
} }
void Near(double threshold = 0.0) void Near(double threshold = 0.0, bool relative = false)
{ {
Mat whole, roi; Mat roi, whole;
gdst_whole.download(whole); gdst_whole.download(whole);
gdst_roi.download(roi); gdst_roi.download(roi);
EXPECT_MAT_NEAR(dst_whole, whole, threshold); if (relative)
EXPECT_MAT_NEAR(dst_roi, roi, threshold); {
EXPECT_MAT_NEAR_RELATIVE(dst_whole, whole, threshold);
EXPECT_MAT_NEAR_RELATIVE(dst_roi, roi, threshold);
}
else
{
EXPECT_MAT_NEAR(dst_whole, whole, threshold);
EXPECT_MAT_NEAR(dst_roi, roi, threshold);
}
} }
}; };
@ -228,7 +236,7 @@ OCL_TEST_P(CornerMinEigenVal, Mat)
cornerMinEigenVal(src_roi, dst_roi, blockSize, apertureSize, borderType); cornerMinEigenVal(src_roi, dst_roi, blockSize, apertureSize, borderType);
ocl::cornerMinEigenVal(gsrc_roi, gdst_roi, blockSize, apertureSize, borderType); ocl::cornerMinEigenVal(gsrc_roi, gdst_roi, blockSize, apertureSize, borderType);
Near(0.02); Near(1e-5, true);
} }
} }
@ -248,7 +256,7 @@ OCL_TEST_P(CornerHarris, Mat)
cornerHarris(src_roi, dst_roi, blockSize, apertureSize, k, borderType); cornerHarris(src_roi, dst_roi, blockSize, apertureSize, k, borderType);
ocl::cornerHarris(gsrc_roi, gdst_roi, blockSize, apertureSize, k, borderType); ocl::cornerHarris(gsrc_roi, gdst_roi, blockSize, apertureSize, k, borderType);
Near(0.02); Near(1e-5, true);
} }
} }