Skip to content

O
opencv
Commit c019d06d authored by Vadim Pisarevsky's avatar Vadim Pisarevsky Committed by OpenCV Buildbot
Browse files
Options

Merge pull request #805 from pengx17:master_canny_fix

parents 5a2c5227 5eb551b3
Hide whitespace changes
Inline Side-by-side
Showing with 21 additions and 245 deletions
modules/ocl/include/opencv2/ocl.hpp
View file @ c019d06d
......@@ -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<FilterEngine_GPU> filterDX, filterDY;
};
///////////////////////////////////////// Hough Transform /////////////////////////////////////////
......
modules/ocl/src/canny.cpp
View file @ c019d06d
......@@ -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<float>(low_thresh), static_cast<float>(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<float>(low_thresh), static_cast<float>(high_thresh));
}
......
modules/ocl/src/opencl/imgproc_canny.cl
View file @ c019d06d
......@@ -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
......
modules/ocl/test/test_canny.cpp
View file @ c019d06d
......@@ -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<cv::ocl::Info> 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
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