diff --git a/modules/ocl/include/opencv2/ocl.hpp b/modules/ocl/include/opencv2/ocl.hpp index 3145c60981..8d92546b06 100644 --- a/modules/ocl/include/opencv2/ocl.hpp +++ b/modules/ocl/include/opencv2/ocl.hpp @@ -861,68 +861,36 @@ namespace cv ///////////////////////////////////////////// Canny ///////////////////////////////////////////// - struct CV_EXPORTS CannyBuf; - - //! compute edges of the input image using Canny operator - // Support CV_8UC1 only - CV_EXPORTS void Canny(const oclMat &image, oclMat &edges, double low_thresh, double high_thresh, int apperture_size = 3, bool L2gradient = false); - CV_EXPORTS void Canny(const oclMat &image, CannyBuf &buf, oclMat &edges, double low_thresh, double high_thresh, int apperture_size = 3, bool L2gradient = false); - CV_EXPORTS void Canny(const oclMat &dx, const oclMat &dy, oclMat &edges, double low_thresh, double high_thresh, bool L2gradient = false); - CV_EXPORTS void Canny(const oclMat &dx, const oclMat &dy, CannyBuf &buf, oclMat &edges, double low_thresh, double high_thresh, bool L2gradient = false); - - struct CV_EXPORTS CannyBuf - { - CannyBuf() : counter(NULL) {} - ~CannyBuf() { release(); } - explicit CannyBuf(const Size &image_size, int apperture_size = 3) : counter(NULL) - { - create(image_size, apperture_size); - } - CannyBuf(const oclMat &dx_, const oclMat &dy_); - - - void create(const Size &image_size, int apperture_size = 3); - - - void release(); - - oclMat dx, dy; - oclMat dx_buf, dy_buf; - - oclMat edgeBuf; - + oclMat magBuf, mapBuf; oclMat trackBuf1, trackBuf2; - void *counter; - Ptr filterDX, filterDY; - }; ///////////////////////////////////////// Hough Transform ///////////////////////////////////////// diff --git a/modules/ocl/src/canny.cpp b/modules/ocl/src/canny.cpp index 9381416950..e06d29904e 100644 --- a/modules/ocl/src/canny.cpp +++ b/modules/ocl/src/canny.cpp @@ -86,7 +86,8 @@ void cv::ocl::CannyBuf::create(const Size &image_size, int apperture_size) filterDY = createDerivFilter_GPU(CV_8U, CV_32S, 0, 1, apperture_size, BORDER_REPLICATE); } } - ensureSizeIsEnough(image_size.height + 2, image_size.width + 2, CV_32FC1, edgeBuf); + ensureSizeIsEnough(image_size.height + 2, image_size.width + 2, CV_32FC1, magBuf); + ensureSizeIsEnough(image_size.height + 2, image_size.width + 2, CV_32FC1, mapBuf); ensureSizeIsEnough(1, image_size.width * image_size.height, CV_16UC2, trackBuf1); ensureSizeIsEnough(1, image_size.width * image_size.height, CV_16UC2, trackBuf2); @@ -107,10 +108,15 @@ void cv::ocl::CannyBuf::release() dy.release(); dx_buf.release(); dy_buf.release(); - edgeBuf.release(); + magBuf.release(); + mapBuf.release(); trackBuf1.release(); trackBuf2.release(); - openCLFree(counter); + if(counter) + { + openCLFree(counter); + counter = NULL; + } } namespace cv @@ -140,13 +146,13 @@ namespace void CannyCaller(CannyBuf &buf, oclMat &dst, float low_thresh, float high_thresh) { using namespace ::cv::ocl::canny; - calcMap_gpu(buf.dx, buf.dy, buf.edgeBuf, buf.edgeBuf, dst.rows, dst.cols, low_thresh, high_thresh); + calcMap_gpu(buf.dx, buf.dy, buf.magBuf, buf.mapBuf, dst.rows, dst.cols, low_thresh, high_thresh); - edgesHysteresisLocal_gpu(buf.edgeBuf, buf.trackBuf1, buf.counter, dst.rows, dst.cols); + edgesHysteresisLocal_gpu(buf.mapBuf, buf.trackBuf1, buf.counter, dst.rows, dst.cols); - edgesHysteresisGlobal_gpu(buf.edgeBuf, buf.trackBuf1, buf.trackBuf2, buf.counter, dst.rows, dst.cols); + edgesHysteresisGlobal_gpu(buf.mapBuf, buf.trackBuf1, buf.trackBuf2, buf.counter, dst.rows, dst.cols); - getEdges_gpu(buf.edgeBuf, dst, dst.rows, dst.cols); + getEdges_gpu(buf.mapBuf, dst, dst.rows, dst.cols); } } @@ -169,20 +175,20 @@ void cv::ocl::Canny(const oclMat &src, CannyBuf &buf, oclMat &dst, double low_th dst.setTo(Scalar::all(0)); buf.create(src.size(), apperture_size); - buf.edgeBuf.setTo(Scalar::all(0)); + buf.magBuf.setTo(Scalar::all(0)); if (apperture_size == 3) { calcSobelRowPass_gpu(src, buf.dx_buf, buf.dy_buf, src.rows, src.cols); - calcMagnitude_gpu(buf.dx_buf, buf.dy_buf, buf.dx, buf.dy, buf.edgeBuf, src.rows, src.cols, L2gradient); + calcMagnitude_gpu(buf.dx_buf, buf.dy_buf, buf.dx, buf.dy, buf.magBuf, src.rows, src.cols, L2gradient); } else { buf.filterDX->apply(src, buf.dx); buf.filterDY->apply(src, buf.dy); - calcMagnitude_gpu(buf.dx, buf.dy, buf.edgeBuf, src.rows, src.cols, L2gradient); + calcMagnitude_gpu(buf.dx, buf.dy, buf.magBuf, src.rows, src.cols, L2gradient); } CannyCaller(buf, dst, static_cast(low_thresh), static_cast(high_thresh)); } @@ -207,8 +213,8 @@ void cv::ocl::Canny(const oclMat &dx, const oclMat &dy, CannyBuf &buf, oclMat &d buf.dx = dx; buf.dy = dy; buf.create(dx.size(), -1); - buf.edgeBuf.setTo(Scalar::all(0)); - calcMagnitude_gpu(buf.dx, buf.dy, buf.edgeBuf, dx.rows, dx.cols, L2gradient); + buf.magBuf.setTo(Scalar::all(0)); + calcMagnitude_gpu(buf.dx, buf.dy, buf.magBuf, dx.rows, dx.cols, L2gradient); CannyCaller(buf, dst, static_cast(low_thresh), static_cast(high_thresh)); } diff --git a/modules/ocl/src/opencl/imgproc_canny.cl b/modules/ocl/src/opencl/imgproc_canny.cl index 5ec4465238..839c16d92c 100644 --- a/modules/ocl/src/opencl/imgproc_canny.cl +++ b/modules/ocl/src/opencl/imgproc_canny.cl @@ -360,188 +360,6 @@ __kernel } } -// non local memory version -__kernel - void calcMap_2 - ( - __global const int * dx, - __global const int * dy, - __global const float * mag, - __global int * map, - int rows, - int cols, - float low_thresh, - float high_thresh, - int dx_step, - int dx_offset, - int dy_step, - int dy_offset, - int mag_step, - int mag_offset, - int map_step, - int map_offset - ) -{ - dx_step /= sizeof(*dx); - dx_offset /= sizeof(*dx); - dy_step /= sizeof(*dy); - dy_offset /= sizeof(*dy); - mag_step /= sizeof(*mag); - mag_offset /= sizeof(*mag); - map_step /= sizeof(*map); - map_offset /= sizeof(*map); - - - int gidx = get_global_id(0); - int gidy = get_global_id(1); - - if(gidy < rows && gidx < cols) - { - int x = dx[gidx + gidy * dx_step]; - int y = dy[gidx + gidy * dy_step]; - const int s = (x ^ y) < 0 ? -1 : 1; - const float m = mag[gidx + 1 + (gidy + 1) * mag_step]; - x = abs(x); - y = abs(y); - - // 0 - the pixel can not belong to an edge - // 1 - the pixel might belong to an edge - // 2 - the pixel does belong to an edge - int edge_type = 0; - if(m > low_thresh) - { - const int tg22x = x * TG22; - const int tg67x = tg22x + (x << (1 + CANNY_SHIFT)); - y <<= CANNY_SHIFT; - if(y < tg22x) - { - if(m > mag[gidx + (gidy + 1) * mag_step] && m >= mag[gidx + 2 + (gidy + 1) * mag_step]) - { - edge_type = 1 + (int)(m > high_thresh); - } - } - else if (y > tg67x) - { - if(m > mag[gidx + 1 + gidy* mag_step] && m >= mag[gidx + 1 + (gidy + 2) * mag_step]) - { - edge_type = 1 + (int)(m > high_thresh); - } - } - else - { - if(m > mag[gidx + 1 - s + gidy * mag_step] && m > mag[gidx + 1 + s + (gidy + 2) * mag_step]) - { - edge_type = 1 + (int)(m > high_thresh); - } - } - } - map[gidx + 1 + (gidy + 1) * map_step] = edge_type; - } -} - -// [256, 1, 1] threaded, local memory version -__kernel - void calcMap_3 - ( - __global const int * dx, - __global const int * dy, - __global const float * mag, - __global int * map, - int rows, - int cols, - float low_thresh, - float high_thresh, - int dx_step, - int dx_offset, - int dy_step, - int dy_offset, - int mag_step, - int mag_offset, - int map_step, - int map_offset - ) -{ - dx_step /= sizeof(*dx); - dx_offset /= sizeof(*dx); - dy_step /= sizeof(*dy); - dy_offset /= sizeof(*dy); - mag_step /= sizeof(*mag); - mag_offset /= sizeof(*mag); - map_step /= sizeof(*map); - map_offset /= sizeof(*map); - - __local float smem[18][18]; - - int lidx = get_local_id(0) % 16; - int lidy = get_local_id(0) / 16; - - int grp_pix = get_global_id(0); // identifies which pixel is processing currently in the target block - int grp_ind = get_global_id(1); // identifies which block of pixels is currently processing - - int grp_idx = (grp_ind % (cols/16)) * 16; - int grp_idy = (grp_ind / (cols/16)) * 16; //(grp_ind / (cols/16)) * 16 - - int gidx = grp_idx + lidx; - int gidy = grp_idy + lidy; - - int tid = get_global_id(0) % 256; - int lx = tid % 18; - int ly = tid / 18; - if(ly < 14) - { - smem[ly][lx] = mag[grp_idx + lx + (grp_idy + ly) * mag_step]; - } - if(ly < 4 && grp_idy + ly + 14 <= rows && grp_idx + lx <= cols) - { - smem[ly + 14][lx] = mag[grp_idx + lx + (grp_idy + ly + 14) * mag_step]; - } - - barrier(CLK_LOCAL_MEM_FENCE); - - if(gidy < rows && gidx < cols) - { - int x = dx[gidx + gidy * dx_step]; - int y = dy[gidx + gidy * dy_step]; - const int s = (x ^ y) < 0 ? -1 : 1; - const float m = smem[lidy + 1][lidx + 1]; - x = abs(x); - y = abs(y); - - // 0 - the pixel can not belong to an edge - // 1 - the pixel might belong to an edge - // 2 - the pixel does belong to an edge - int edge_type = 0; - if(m > low_thresh) - { - const int tg22x = x * TG22; - const int tg67x = tg22x + (x << (1 + CANNY_SHIFT)); - y <<= CANNY_SHIFT; - if(y < tg22x) - { - if(m > smem[lidy + 1][lidx] && m >= smem[lidy + 1][lidx + 2]) - { - edge_type = 1 + (int)(m > high_thresh); - } - } - else if (y > tg67x) - { - if(m > smem[lidy][lidx + 1]&& m >= smem[lidy + 2][lidx + 1]) - { - edge_type = 1 + (int)(m > high_thresh); - } - } - else - { - if(m > smem[lidy][lidx + 1 - s]&& m > smem[lidy + 2][lidx + 1 + s]) - { - edge_type = 1 + (int)(m > high_thresh); - } - } - } - map[gidx + 1 + (gidy + 1) * map_step] = edge_type; - } -} - #undef CANNY_SHIFT #undef TG22 diff --git a/modules/ocl/test/test_canny.cpp b/modules/ocl/test/test_canny.cpp index cac6b66f51..e7b9316d8e 100644 --- a/modules/ocl/test/test_canny.cpp +++ b/modules/ocl/test/test_canny.cpp @@ -45,7 +45,6 @@ #include "precomp.hpp" #ifdef HAVE_OPENCL -#define SHOW_RESULT 0 //////////////////////////////////////////////////////// // Canny @@ -59,13 +58,10 @@ PARAM_TEST_CASE(Canny, AppertureSize, L2gradient) bool useL2gradient; cv::Mat edges_gold; - //std::vector oclinfo; virtual void SetUp() { apperture_size = GET_PARAM(0); useL2gradient = GET_PARAM(1); - //int devnums = getDevice(oclinfo); - //CV_Assert(devnums > 0); } }; @@ -83,26 +79,13 @@ TEST_P(Canny, Accuracy) cv::ocl::oclMat edges; cv::ocl::Canny(ocl_img, edges, low_thresh, high_thresh, apperture_size, useL2gradient); - char filename [100]; - sprintf(filename, "G:/Valve_edges_a%d_L2Grad%d.jpg", apperture_size, (int)useL2gradient); - cv::Mat edges_gold; cv::Canny(img, edges_gold, low_thresh, high_thresh, apperture_size, useL2gradient); -#if SHOW_RESULT - cv::Mat edges_x2, ocl_edges(edges); - edges_x2.create(edges.rows, edges.cols * 2, edges.type()); - edges_x2.setTo(0); - cv::add(edges_gold, cv::Mat(edges_x2, cv::Rect(0, 0, edges_gold.cols, edges_gold.rows)), cv::Mat(edges_x2, cv::Rect(0, 0, edges_gold.cols, edges_gold.rows))); - cv::add(ocl_edges, cv::Mat(edges_x2, cv::Rect(edges_gold.cols, 0, edges_gold.cols, edges_gold.rows)), cv::Mat(edges_x2, cv::Rect(edges_gold.cols, 0, edges_gold.cols, edges_gold.rows))); - cv::namedWindow("Canny result (left: cpu, right: ocl)"); - cv::imshow("Canny result (left: cpu, right: ocl)", edges_x2); - cv::waitKey(); -#endif //OUTPUT_RESULT EXPECT_MAT_SIMILAR(edges_gold, edges, 1e-2); } -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Canny, testing::Combine( +INSTANTIATE_TEST_CASE_P(OCL_ImgProc, Canny, testing::Combine( testing::Values(AppertureSize(3), AppertureSize(5)), testing::Values(L2gradient(false), L2gradient(true)))); #endif \ No newline at end of file