From 429f84e59e45a66f56540e67f5ab3a9c8dc34249 Mon Sep 17 00:00:00 2001 From: peng xiao Date: Thu, 6 Jun 2013 11:44:35 +0800 Subject: [PATCH] Fix a bug of cornerHarris and cornerMinEigenVal. The bug is a buffer overrun when border type is reflect101. It is found that gfft crashed with input of size 100x100 on Intel CPU. --- modules/ocl/src/opencl/imgproc_calcHarris.cl | 15 ++++++++------- modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl | 16 +++++++++------- modules/ocl/test/test_optflow.cpp | 2 +- 3 files changed, 18 insertions(+), 15 deletions(-) diff --git a/modules/ocl/src/opencl/imgproc_calcHarris.cl b/modules/ocl/src/opencl/imgproc_calcHarris.cl index 15742d6..1911a72 100644 --- a/modules/ocl/src/opencl/imgproc_calcHarris.cl +++ b/modules/ocl/src/opencl/imgproc_calcHarris.cl @@ -130,28 +130,29 @@ __kernel void calcHarris(__global const float *Dx,__global const float *Dy, __gl data[2][i] = dy_data[i] * dy_data[i]; } #else - for(int i=0; i < ksY+1; i++) - { + int clamped_col = min(dst_cols, col); + for(int i=0; i < ksY+1; i++) + { int dx_selected_row; int dx_selected_col; dx_selected_row = ADDR_H(dx_startY+i, 0, dx_whole_rows); dx_selected_row = ADDR_B(dx_startY+i, dx_whole_rows, dx_selected_row); - dx_selected_col = ADDR_L(dx_startX+col, 0, dx_whole_cols); - dx_selected_col = ADDR_R(dx_startX+col, dx_whole_cols, dx_selected_col); + dx_selected_col = ADDR_L(dx_startX+clamped_col, 0, dx_whole_cols); + dx_selected_col = ADDR_R(dx_startX+clamped_col, dx_whole_cols, dx_selected_col); dx_data[i] = Dx[dx_selected_row * (dx_step>>2) + dx_selected_col]; int dy_selected_row; int dy_selected_col; dy_selected_row = ADDR_H(dy_startY+i, 0, dy_whole_rows); dy_selected_row = ADDR_B(dy_startY+i, dy_whole_rows, dy_selected_row); - dy_selected_col = ADDR_L(dy_startX+col, 0, dy_whole_cols); - dy_selected_col = ADDR_R(dy_startX+col, dy_whole_cols, dy_selected_col); + dy_selected_col = ADDR_L(dy_startX+clamped_col, 0, dy_whole_cols); + dy_selected_col = ADDR_R(dy_startX+clamped_col, dy_whole_cols, dy_selected_col); dy_data[i] = Dy[dy_selected_row * (dy_step>>2) + dy_selected_col]; 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]; - } + } #endif float sum0 = 0.0, sum1 = 0.0, sum2 = 0.0; for(int i=1; i < ksY; i++) diff --git a/modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl b/modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl index 662fbb0..462ec77 100644 --- a/modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl +++ b/modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl @@ -130,28 +130,30 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy, data[2][i] = dy_data[i] * dy_data[i]; } #else - for(int i=0; i < ksY+1; i++) - { + int clamped_col = min(dst_cols, col); + + for(int i=0; i < ksY+1; i++) + { int dx_selected_row; int dx_selected_col; dx_selected_row = ADDR_H(dx_startY+i, 0, dx_whole_rows); dx_selected_row = ADDR_B(dx_startY+i, dx_whole_rows, dx_selected_row); - dx_selected_col = ADDR_L(dx_startX+col, 0, dx_whole_cols); - dx_selected_col = ADDR_R(dx_startX+col, dx_whole_cols, dx_selected_col); + dx_selected_col = ADDR_L(dx_startX+clamped_col, 0, dx_whole_cols); + dx_selected_col = ADDR_R(dx_startX+clamped_col, dx_whole_cols, dx_selected_col); dx_data[i] = Dx[dx_selected_row * (dx_step>>2) + dx_selected_col]; int dy_selected_row; int dy_selected_col; dy_selected_row = ADDR_H(dy_startY+i, 0, dy_whole_rows); dy_selected_row = ADDR_B(dy_startY+i, dy_whole_rows, dy_selected_row); - dy_selected_col = ADDR_L(dy_startX+col, 0, dy_whole_cols); - dy_selected_col = ADDR_R(dy_startX+col, dy_whole_cols, dy_selected_col); + dy_selected_col = ADDR_L(dy_startX+clamped_col, 0, dy_whole_cols); + dy_selected_col = ADDR_R(dy_startX+clamped_col, dy_whole_cols, dy_selected_col); dy_data[i] = Dy[dy_selected_row * (dy_step>>2) + dy_selected_col]; 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]; - } + } #endif float sum0 = 0.0, sum1 = 0.0, sum2 = 0.0; for(int i=1; i < ksY; i++) diff --git a/modules/ocl/test/test_optflow.cpp b/modules/ocl/test/test_optflow.cpp index 72689f3..0121be8 100644 --- a/modules/ocl/test/test_optflow.cpp +++ b/modules/ocl/test/test_optflow.cpp @@ -121,7 +121,7 @@ TEST_P(GoodFeaturesToTrack, EmptyCorners) cv::ocl::GoodFeaturesToTrackDetector_OCL detector(maxCorners, qualityLevel, minDistance); - cv::ocl::oclMat src(100, 128, CV_8UC1, cv::Scalar::all(0)); + cv::ocl::oclMat src(100, 100, CV_8UC1, cv::Scalar::all(0)); cv::ocl::oclMat corners(1, maxCorners, CV_32FC2); detector(src, corners); -- 2.7.4