Skip to content

O
opencv
Commit d9489bbe authored by Alexey Spizhevoy's avatar Alexey Spizhevoy
Browse files
Options

fixed bug in gpu::cornerHarris and gpu::cornerMinEigenVal

parent db5cce88
Hide whitespace changes
Inline Side-by-side
Showing with 41 additions and 23 deletions
modules/gpu/src/cuda/imgproc.cu
View file @ d9489bbe
...@@ -478,10 +478,13 @@ namespace cv { namespace gpu { namespace imgproc ...@@ -478,10 +478,13 @@ namespace cv { namespace gpu { namespace imgproc
float b = 0.f; float b = 0.f;
float c = 0.f; float c = 0.f;
const unsigned int j_begin = max(x - block_size, 0); int offset1 = -(block_size / 2);
const unsigned int i_begin = max(y - block_size, 0); int offset2 = offset1 + block_size;
const unsigned int j_end = min(x + block_size + 1, cols);
const unsigned int i_end = min(y + block_size + 1, rows); unsigned int j_begin = max(x + offset1, 0);
unsigned int i_begin = max(y + offset1, 0);
unsigned int j_end = min(x + offset2, cols);
unsigned int i_end = min(y + offset2, rows);
for (unsigned int i = i_begin; i < i_end; ++i) for (unsigned int i = i_begin; i < i_end; ++i)
{ {
...@@ -509,7 +512,7 @@ namespace cv { namespace gpu { namespace imgproc ...@@ -509,7 +512,7 @@ namespace cv { namespace gpu { namespace imgproc
dim3 threads(32, 8); dim3 threads(32, 8);
dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));
cornerHarris_kernel<<<grid, threads>>>(cols, rows, block_size / 2, k, Dx, Dy, dst); cornerHarris_kernel<<<grid, threads>>>(cols, rows, block_size, k, Dx, Dy, dst);
cudaSafeCall(cudaThreadSynchronize()); cudaSafeCall(cudaThreadSynchronize());
} }
...@@ -527,10 +530,13 @@ namespace cv { namespace gpu { namespace imgproc ...@@ -527,10 +530,13 @@ namespace cv { namespace gpu { namespace imgproc
float b = 0.f; float b = 0.f;
float c = 0.f; float c = 0.f;
const unsigned int j_begin = max(x - block_size, 0); int offset1 = -(block_size / 2);
const unsigned int i_begin = max(y - block_size, 0); int offset2 = offset1 + block_size;
const unsigned int j_end = min(x + block_size + 1, cols);
const unsigned int i_end = min(y + block_size + 1, rows); unsigned int j_begin = max(x + offset1, 0);
unsigned int i_begin = max(y + offset1, 0);
unsigned int j_end = min(x + offset2, cols);
unsigned int i_end = min(y + offset2, rows);
for (unsigned int i = i_begin; i < i_end; ++i) for (unsigned int i = i_begin; i < i_end; ++i)
{ {
...@@ -560,7 +566,7 @@ namespace cv { namespace gpu { namespace imgproc ...@@ -560,7 +566,7 @@ namespace cv { namespace gpu { namespace imgproc
dim3 threads(32, 8); dim3 threads(32, 8);
dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));
cornerMinEigenVal_kernel<<<grid, threads>>>(cols, rows, block_size / 2, Dx, Dy, dst); cornerMinEigenVal_kernel<<<grid, threads>>>(cols, rows, block_size, Dx, Dy, dst);
cudaSafeCall(cudaThreadSynchronize()); cudaSafeCall(cudaThreadSynchronize());
} }
}}} }}}
tests/gpu/src/imgproc_gpu.cpp
View file @ d9489bbe
...@@ -614,12 +614,11 @@ struct CV_GpuCornerHarrisTest: CvTest ...@@ -614,12 +614,11 @@ struct CV_GpuCornerHarrisTest: CvTest
{ {
try try
{ {
int rows = 1 + rand() % 300, cols = 1 + rand() % 300; for (int i = 0; i < 5; ++i)
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, -1)) return;
for (int i = 0; i < 3; ++i)
{ {
rows = 1 + rand() % 300; cols = 1 + rand() % 300; int rows = 1 + rand() % 300, cols = 1 + rand() % 300;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return; if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, -1)) return;
} }
} }
catch (const Exception& e) catch (const Exception& e)
...@@ -645,14 +644,21 @@ struct CV_GpuCornerHarrisTest: CvTest ...@@ -645,14 +644,21 @@ struct CV_GpuCornerHarrisTest: CvTest
cv::gpu::GpuMat dst; cv::gpu::GpuMat dst;
cv::gpu::cornerHarris(cv::gpu::GpuMat(src), dst, blockSize, apertureSize, k); cv::gpu::cornerHarris(cv::gpu::GpuMat(src), dst, blockSize, apertureSize, k);
int asize = apertureSize > 0 ? apertureSize : 3;
cv::Mat dsth = dst; cv::Mat dsth = dst;
for (int i = apertureSize + 2; i < dst.rows - apertureSize - 2; ++i) for (int i = max(blockSize, asize) + 2; i < dst.rows - max(blockSize, asize) - 2; ++i)
{ {
for (int j = apertureSize + 2; j < dst.cols - apertureSize - 2; ++j) for (int j = max(blockSize, asize) + 2; j < dst.cols - max(blockSize, asize) - 2; ++j)
{ {
float a = dst_gold.at<float>(i, j); float a = dst_gold.at<float>(i, j);
float b = dsth.at<float>(i, j); float b = dsth.at<float>(i, j);
if (fabs(a - b) > 1e-3f) return false; if (fabs(a - b) > 1e-3f)
{
ts->printf(CvTS::CONSOLE, "%d %d %f %f %d\n", i, j, a, b, apertureSize);
ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
return false;
};
} }
} }
return true; return true;
...@@ -670,11 +676,10 @@ struct CV_GpuCornerMinEigenValTest: CvTest ...@@ -670,11 +676,10 @@ struct CV_GpuCornerMinEigenValTest: CvTest
{ {
try try
{ {
int rows = 1 + rand() % 300, cols = 1 + rand() % 300;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, -1)) return;
for (int i = 0; i < 3; ++i) for (int i = 0; i < 3; ++i)
{ {
rows = 1 + rand() % 300; cols = 1 + rand() % 300; int rows = 1 + rand() % 300, cols = 1 + rand() % 300;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, -1)) return;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return; if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return;
} }
} }
...@@ -700,14 +705,21 @@ struct CV_GpuCornerMinEigenValTest: CvTest ...@@ -700,14 +705,21 @@ struct CV_GpuCornerMinEigenValTest: CvTest
cv::gpu::GpuMat dst; cv::gpu::GpuMat dst;
cv::gpu::cornerMinEigenVal(cv::gpu::GpuMat(src), dst, blockSize, apertureSize); cv::gpu::cornerMinEigenVal(cv::gpu::GpuMat(src), dst, blockSize, apertureSize);
int asize = apertureSize > 0 ? apertureSize : 3;
cv::Mat dsth = dst; cv::Mat dsth = dst;
for (int i = apertureSize + 2; i < dst.rows - apertureSize - 2; ++i) for (int i = max(blockSize, asize) + 2; i < dst.rows - max(blockSize, asize) - 2; ++i)
{ {
for (int j = apertureSize + 2; j < dst.cols - apertureSize - 2; ++j) for (int j = max(blockSize, asize) + 2; j < dst.cols - max(blockSize, asize) - 2; ++j)
{ {
float a = dst_gold.at<float>(i, j); float a = dst_gold.at<float>(i, j);
float b = dsth.at<float>(i, j); float b = dsth.at<float>(i, j);
if (fabs(a - b) > 1e-3f) return false; if (fabs(a - b) > 1e-3f)
{
ts->printf(CvTS::CONSOLE, "%d %d %f %f %d %d\n", i, j, a, b, apertureSize, blockSize);
ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
return false;
};
} }
} }
return true; return true;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment