//! disabled until fix crash\r
CV_EXPORTS void Canny(const GpuMat& image, GpuMat& edges, double threshold1, double threshold2, int apertureSize = 3);\r
\r
+ //! computes Harris cornerness criteria at each image pixel \r
+ // (does BORDER_CONSTANT interpolation with 0 as the fill value)\r
+ CV_EXPORTS void cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int apertureSize, double k);\r
+\r
//////////////////////////////// Filter Engine ////////////////////////////////\r
\r
/*!\r
{\r
reprojectImageTo3D_caller(disp, xyzw, q, stream);\r
}\r
+\r
+/////////////////////////////////////////// Corner Harris /////////////////////////////////////////////////\r
+\r
+ __global__ void cornerHarris_kernel(const int cols, const int rows, const int block_size, const float k, const PtrStep Dx, const PtrStep Dy, PtrStep dst)\r
+ {\r
+ const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;\r
+ const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;\r
+\r
+ if (x < cols && y < rows)\r
+ {\r
+ float a = 0.f;\r
+ float b = 0.f;\r
+ float c = 0.f;\r
+\r
+ const unsigned int j_begin = max(x - block_size, 0);\r
+ const unsigned int i_begin = max(y - block_size, 0);\r
+ const unsigned int j_end = min(x + block_size + 1, cols);\r
+ const unsigned int i_end = min(y + block_size + 1, rows);\r
+\r
+ for (unsigned int i = i_begin; i < i_end; ++i)\r
+ {\r
+ const float* dx_row = (const float*)Dx.ptr(i);\r
+ const float* dy_row = (const float*)Dy.ptr(i);\r
+ for (unsigned int j = j_begin; j < j_end; ++j)\r
+ {\r
+ float dx = dx_row[j];\r
+ float dy = dy_row[j];\r
+ a += dx * dx;\r
+ b += dx * dy;\r
+ c += dy * dy;\r
+ }\r
+ }\r
+\r
+ ((float*)dst.ptr(y))[x] = a * c - b * b - k * (a + c) * (a + c);\r
+ }\r
+ }\r
+\r
+ void cornerHarris_caller(const int block_size, const float k, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst)\r
+ {\r
+ const int rows = Dx.rows;\r
+ const int cols = Dx.cols;\r
+\r
+ dim3 threads(32, 8);\r
+ dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
+\r
+ cornerHarris_kernel<<<grid, threads>>>(cols, rows, block_size / 2, k, Dx, Dy, dst);\r
+ cudaSafeCall(cudaThreadSynchronize());\r
+ }\r
}}}\r
void cv::gpu::histEven(const GpuMat&, GpuMat*, int*, int*, int*) { throw_nogpu(); }\r
void cv::gpu::histRange(const GpuMat&, GpuMat&, const GpuMat&) { throw_nogpu(); }\r
void cv::gpu::histRange(const GpuMat&, GpuMat*, const GpuMat*) { throw_nogpu(); }\r
+void cv::gpu::cornerHarris(const GpuMat&, GpuMat&, int, int, double) { throw_nogpu(); }\r
+\r
\r
#else /* !defined (HAVE_CUDA) */\r
\r
hist_callers[src.depth()](src, hist, levels);\r
}\r
\r
+namespace cv { namespace gpu { namespace imgproc {\r
+\r
+ void cornerHarris_caller(const int block_size, const float k, const DevMem2D Dx, const DevMem2D Dy, DevMem2D dst);\r
+\r
+}}}\r
+\r
+void cv::gpu::cornerHarris(const GpuMat& src, GpuMat& dst, int blockSize, int apertureSize, double k)\r
+{\r
+ CV_Assert(src.type() == CV_32F);\r
+\r
+ double scale = (double)(1 << ((apertureSize > 0 ? apertureSize : 3) - 1)) * blockSize;\r
+ if (apertureSize < 0) scale *= 2.;\r
+ scale = 1./scale;\r
+\r
+ GpuMat Dx, Dy;\r
+ if (apertureSize > 0)\r
+ {\r
+ Sobel(src, Dx, CV_32F, 1, 0, apertureSize, scale);\r
+ Sobel(src, Dy, CV_32F, 0, 1, apertureSize, scale);\r
+ }\r
+ else\r
+ {\r
+ Scharr(src, Dx, CV_32F, 1, 0, scale);\r
+ Scharr(src, Dy, CV_32F, 0, 1, scale);\r
+ }\r
+\r
+ dst.create(src.size(), CV_32F);\r
+ imgproc::cornerHarris_caller(blockSize, (float)k, Dx, Dy, dst);\r
+}\r
+\r
+\r
#endif /* !defined (HAVE_CUDA) */\r
}\r
}\r
\r
+////////////////////////////////////////////////////////////////////////\r
+// Corner Harris feature detector\r
+\r
+struct CV_GpuCornerHarrisTest: CvTest \r
+{\r
+ CV_GpuCornerHarrisTest(): CvTest("GPU-CornerHarrisTest", "cornerHarris") {}\r
+\r
+ void run(int)\r
+ {\r
+ try\r
+ {\r
+ int rows = 1 + rand() % 300, cols = 1 + rand() % 300;\r
+ if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, -1)) return;\r
+ for (int i = 0; i < 3; ++i)\r
+ {\r
+ rows = 1 + rand() % 300; cols = 1 + rand() % 300;\r
+ if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return;\r
+ }\r
+ }\r
+ catch (const Exception& e)\r
+ {\r
+ if (!check_and_treat_gpu_exception(e, ts)) throw;\r
+ return;\r
+ }\r
+ }\r
+\r
+ bool compareToCpuTest(int rows, int cols, int depth, int blockSize, int apertureSize)\r
+ {\r
+ RNG rng;\r
+ cv::Mat src(rows, cols, depth);\r
+ if (depth == CV_32F) \r
+ rng.fill(src, RNG::UNIFORM, cv::Scalar(0), cv::Scalar(1));\r
+\r
+ double k = 0.1;\r
+ int borderType = BORDER_DEFAULT;\r
+\r
+ cv::Mat dst_gold;\r
+ cv::cornerHarris(src, dst_gold, blockSize, apertureSize, k, borderType); \r
+\r
+ cv::gpu::GpuMat dst;\r
+ cv::gpu::cornerHarris(cv::gpu::GpuMat(src), dst, blockSize, apertureSize, k);\r
+\r
+ cv::Mat dsth = dst;\r
+ for (int i = apertureSize + 2; i < dst.rows - apertureSize - 2; ++i)\r
+ {\r
+ for (int j = apertureSize + 2; j < dst.cols - apertureSize - 2; ++j)\r
+ {\r
+ float a = dst_gold.at<float>(i, j);\r
+ float b = dsth.at<float>(i, j);\r
+ if (fabs(a - b) > 1e-3f) return false;\r
+ }\r
+ }\r
+ return true;\r
+ }\r
+};\r
+\r
/////////////////////////////////////////////////////////////////////////////\r
/////////////////// tests registration /////////////////////////////////////\r
/////////////////////////////////////////////////////////////////////////////\r
CV_GpuNppImageCannyTest CV_GpuNppImageCanny_test;\r
CV_GpuCvtColorTest CV_GpuCvtColor_test;\r
CV_GpuHistogramsTest CV_GpuHistograms_test;\r
+CV_GpuCornerHarrisTest CV_GpuCornerHarris_test;\r
+\r