From 99ae9d9cc1db4d4906ad061d63ad438aa2bbe5e5 Mon Sep 17 00:00:00 2001 From: Alexander Alekhin Date: Fri, 1 Nov 2013 16:38:04 +0400 Subject: [PATCH] ocl: corner*: fix memory access in kernels; change error check to relative --- modules/ocl/src/opencl/imgproc_calcHarris.cl | 18 ++++++++--------- .../ocl/src/opencl/imgproc_calcMinEigenVal.cl | 16 +++++++-------- modules/ocl/test/test_imgproc.cpp | 20 +++++++++++++------ 3 files changed, 30 insertions(+), 24 deletions(-) diff --git a/modules/ocl/src/opencl/imgproc_calcHarris.cl b/modules/ocl/src/opencl/imgproc_calcHarris.cl index 0a981e12e8..3f53ddf9a5 100644 --- a/modules/ocl/src/opencl/imgproc_calcHarris.cl +++ b/modules/ocl/src/opencl/imgproc_calcHarris.cl @@ -119,18 +119,16 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g __local float temp[6][THREADS]; #ifdef BORDER_CONSTANT - bool dx_con,dy_con; - float dx_s, dy_s; 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; - dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+col)]; - dx_data[i] = dx_con ? dx_s : 0.0f; - - 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)]; - dy_data[i] = dy_con ? dy_s : 0.0f; - + bool dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows; + int indexDx = (dx_startY+i)*(dx_step>>2)+(dx_startX+col); + float dx_s = dx_con ? Dx[indexDx] : 0.0f; + dx_data[i] = dx_s; + bool dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows; + 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[1][i] = dx_data[i] * dy_data[i]; data[2][i] = dy_data[i] * dy_data[i]; diff --git a/modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl b/modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl index 110d204a59..c598246aec 100644 --- a/modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl +++ b/modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl @@ -118,16 +118,16 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy, __local float temp[6][THREADS]; #ifdef BORDER_CONSTANT - bool dx_con, dy_con; - float dx_s, dy_s; 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; - dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+col)]; - dx_data[i] = dx_con ? dx_s : 0.0f; - 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)]; - dy_data[i] = dy_con ? dy_s : 0.0f; + bool dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows; + int indexDx = (dx_startY+i)*(dx_step>>2)+(dx_startX+col); + float dx_s = dx_con ? Dx[indexDx] : 0.0f; + dx_data[i] = dx_s; + bool dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows; + 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[1][i] = dx_data[i] * dy_data[i]; data[2][i] = dy_data[i] * dy_data[i]; diff --git a/modules/ocl/test/test_imgproc.cpp b/modules/ocl/test/test_imgproc.cpp index e981d437e8..7e4b14ecae 100644 --- a/modules/ocl/test/test_imgproc.cpp +++ b/modules/ocl/test/test_imgproc.cpp @@ -93,14 +93,22 @@ PARAM_TEST_CASE(ImgprocTestBase, MatType, 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_roi.download(roi); - EXPECT_MAT_NEAR(dst_whole, whole, threshold); - EXPECT_MAT_NEAR(dst_roi, roi, threshold); + if (relative) + { + 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); 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); ocl::cornerHarris(gsrc_roi, gdst_roi, blockSize, apertureSize, k, borderType); - Near(0.02); + Near(1e-5, true); } }