diff --git a/modules/ocl/include/opencv2/ocl/ocl.hpp b/modules/ocl/include/opencv2/ocl/ocl.hpp index c5075d811a..a58cce3b3d 100644 --- a/modules/ocl/include/opencv2/ocl/ocl.hpp +++ b/modules/ocl/include/opencv2/ocl/ocl.hpp @@ -895,6 +895,34 @@ namespace cv // Supports TM_SQDIFF, TM_CCORR for type 32FC1 and 32FC4 CV_EXPORTS void matchTemplate(const oclMat& image, const oclMat& templ, oclMat& result, int method, MatchTemplateBuf& buf); + ///////////////////////////////////////////// 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() {} + explicit CannyBuf(const Size& image_size, int apperture_size = 3) {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 trackBuf1, trackBuf2; + oclMat counter; + Ptr filterDX, filterDY; + }; + #ifdef HAVE_CLAMDFFT ///////////////////////////////////////// clAmdFft related ///////////////////////////////////////// // the two functions must be called before/after run any fft library functions. diff --git a/modules/ocl/src/canny.cpp b/modules/ocl/src/canny.cpp new file mode 100644 index 0000000000..f41d953dc7 --- /dev/null +++ b/modules/ocl/src/canny.cpp @@ -0,0 +1,419 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other oclMaterials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include +#include "precomp.hpp" + +using namespace cv; +using namespace cv::ocl; +using namespace std; + +#if !defined (HAVE_OPENCL) +void cv::ocl::Canny(const oclMat& image, oclMat& edges, double low_thresh, double high_thresh, int apperture_size = 3, bool L2gradient = false) { throw_nogpu(); } +void cv::ocl::Canny(const oclMat& image, CannyBuf& buf, oclMat& edges, double low_thresh, double high_thresh, int apperture_size = 3, bool L2gradient = false){ throw_nogpu(); } +void cv::ocl::Canny(const oclMat& dx, const oclMat& dy, oclMat& edges, double low_thresh, double high_thresh, bool L2gradient = false){ throw_nogpu(); } +void cv::ocl::Canny(const oclMat& dx, const oclMat& dy, CannyBuf& buf, oclMat& edges, double low_thresh, double high_thresh, bool L2gradient = false){ throw_nogpu(); } +#else + +namespace cv +{ + namespace ocl + { + ///////////////////////////OpenCL kernel strings/////////////////////////// + extern const char *imgproc_canny; + } +} + +cv::ocl::CannyBuf::CannyBuf(const oclMat& dx_, const oclMat& dy_) : dx(dx_), dy(dy_) +{ + CV_Assert(dx_.type() == CV_32SC1 && dy_.type() == CV_32SC1 && dx_.size() == dy_.size()); + + create(dx_.size(), -1); +} + +void cv::ocl::CannyBuf::create(const Size& image_size, int apperture_size) +{ + dx.create(image_size, CV_32SC1); + dy.create(image_size, CV_32SC1); + + if(apperture_size == 3) + { + dx_buf.create(image_size, CV_32SC1); + dy_buf.create(image_size, CV_32SC1); + } + else if(apperture_size > 0) + { + Mat kx, ky; + if (!filterDX) + { + filterDX = createDerivFilter_GPU(CV_32F, CV_32F, 1, 0, apperture_size, BORDER_REPLICATE); + } + if (!filterDY) + { + filterDY = createDerivFilter_GPU(CV_32F, CV_32F, 0, 1, apperture_size, BORDER_REPLICATE); + } + } + edgeBuf.create(image_size.height + 2, image_size.width + 2, CV_32FC1); + + trackBuf1.create(1, image_size.width * image_size.height, CV_16UC2); + trackBuf2.create(1, image_size.width * image_size.height, CV_16UC2); + + counter.create(1,1, CV_32SC1); +} + +void cv::ocl::CannyBuf::release() +{ + dx.release(); + dy.release(); + dx_buf.release(); + dy_buf.release(); + edgeBuf.release(); + trackBuf1.release(); + trackBuf2.release(); + counter.release(); +} + +namespace cv { namespace ocl { + namespace canny + { + void calcSobelRowPass_gpu(const oclMat& src, oclMat& dx_buf, oclMat& dy_buf, int rows, int cols); + + void calcMagnitude_gpu(const oclMat& dx_buf, const oclMat& dy_buf, oclMat& dx, oclMat& dy, oclMat& mag, int rows, int cols, bool L2Grad); + void calcMagnitude_gpu(const oclMat& dx, const oclMat& dy, oclMat& mag, int rows, int cols, bool L2Grad); + + void calcMap_gpu(oclMat& dx, oclMat& dy, oclMat& mag, oclMat& map, int rows, int cols, float low_thresh, float high_thresh); + + void edgesHysteresisLocal_gpu(oclMat& map, oclMat& st1, oclMat& counter, int rows, int cols); + + void edgesHysteresisGlobal_gpu(oclMat& map, oclMat& st1, oclMat& st2, oclMat& counter, int rows, int cols); + + void getEdges_gpu(oclMat& map, oclMat& dst, int rows, int cols); + } +}}// cv::ocl + +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); + + edgesHysteresisLocal_gpu(buf.edgeBuf, buf.trackBuf1, buf.counter, dst.rows, dst.cols); + + edgesHysteresisGlobal_gpu(buf.edgeBuf, buf.trackBuf1, buf.trackBuf2, buf.counter, dst.rows, dst.cols); + + getEdges_gpu(buf.edgeBuf, dst, dst.rows, dst.cols); + } +} + +void cv::ocl::Canny(const oclMat& src, oclMat& dst, double low_thresh, double high_thresh, int apperture_size, bool L2gradient) +{ + CannyBuf buf(src.size(), apperture_size); + Canny(src, buf, dst, low_thresh, high_thresh, apperture_size, L2gradient); +} + +void cv::ocl::Canny(const oclMat& src, CannyBuf& buf, oclMat& dst, double low_thresh, double high_thresh, int apperture_size, bool L2gradient) +{ + using namespace ::cv::ocl::canny; + + CV_Assert(src.type() == CV_8UC1); + + if( low_thresh > high_thresh ) + std::swap( low_thresh, high_thresh ); + + dst.create(src.size(), CV_8U); + dst.setTo(Scalar::all(0)); + + buf.create(src.size(), apperture_size); + buf.edgeBuf.setTo(Scalar::all(0)); + buf.counter.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); + } + else + { + // FIXME: + // current ocl implementation requires the src and dst having same type + // convertTo is time consuming so this may be optimized later. + oclMat src_omat32f = src; + src.convertTo(src_omat32f, CV_32F); // FIXME + + buf.filterDX->apply(src_omat32f, buf.dx); + buf.filterDY->apply(src_omat32f, buf.dy); + + buf.dx.convertTo(buf.dx, CV_32S); // FIXME + buf.dy.convertTo(buf.dy, CV_32S); // FIXME + + calcMagnitude_gpu(buf.dx, buf.dy, buf.edgeBuf, src.rows, src.cols, L2gradient); + } + CannyCaller(buf, dst, static_cast(low_thresh), static_cast(high_thresh)); +} +void cv::ocl::Canny(const oclMat& dx, const oclMat& dy, oclMat& dst, double low_thresh, double high_thresh, bool L2gradient) +{ + CannyBuf buf(dx, dy); + Canny(dx, dy, buf, dst, low_thresh, high_thresh, L2gradient); +} + +void cv::ocl::Canny(const oclMat& dx, const oclMat& dy, CannyBuf& buf, oclMat& dst, double low_thresh, double high_thresh, bool L2gradient) +{ + using namespace ::cv::ocl::canny; + + CV_Assert(dx.type() == CV_32SC1 && dy.type() == CV_32SC1 && dx.size() == dy.size()); + + if( low_thresh > high_thresh ) + std::swap( low_thresh, high_thresh); + + dst.create(dx.size(), CV_8U); + dst.setTo(Scalar::all(0)); + + buf.dx = dx; buf.dy = dy; + buf.create(dx.size(), -1); + buf.edgeBuf.setTo(Scalar::all(0)); + buf.counter.setTo(Scalar::all(0)); + calcMagnitude_gpu(buf.dx, buf.dy, buf.edgeBuf, dx.rows, dx.cols, L2gradient); + + CannyCaller(buf, dst, static_cast(low_thresh), static_cast(high_thresh)); +} + +void canny::calcSobelRowPass_gpu(const oclMat& src, oclMat& dx_buf, oclMat& dy_buf, int rows, int cols) +{ + Context *clCxt = src.clCxt; + string kernelName = "calcSobelRowPass"; + vector< pair > args; + + args.push_back( make_pair( sizeof(cl_mem), (void *)&src.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&dx_buf.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&dy_buf.data)); + args.push_back( make_pair( sizeof(cl_int), (void *)&rows)); + args.push_back( make_pair( sizeof(cl_int), (void *)&cols)); + args.push_back( make_pair( sizeof(cl_int), (void *)&src.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&src.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx_buf.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx_buf.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy_buf.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy_buf.offset)); + + size_t globalThreads[3] = {cols, rows, 1}; + size_t localThreads[3] = {16, 16, 1}; + openCLExecuteKernel(clCxt, &imgproc_canny, kernelName, globalThreads, localThreads, args, -1, -1); +} + +void canny::calcMagnitude_gpu(const oclMat& dx_buf, const oclMat& dy_buf, oclMat& dx, oclMat& dy, oclMat& mag, int rows, int cols, bool L2Grad) +{ + Context *clCxt = dx_buf.clCxt; + string kernelName = "calcMagnitude_buf"; + vector< pair > args; + + args.push_back( make_pair( sizeof(cl_mem), (void *)&dx_buf.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&dy_buf.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&dx.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&dy.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&mag.data)); + args.push_back( make_pair( sizeof(cl_int), (void *)&rows)); + args.push_back( make_pair( sizeof(cl_int), (void *)&cols)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx_buf.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx_buf.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy_buf.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy_buf.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&mag.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&mag.offset)); + + size_t globalThreads[3] = {cols, rows, 1}; + size_t localThreads[3] = {16, 16, 1}; + + char build_options [15] = ""; + if(L2Grad) + { + strcat(build_options, "-D L2GRAD"); + } + openCLExecuteKernel(clCxt, &imgproc_canny, kernelName, globalThreads, localThreads, args, -1, -1, build_options); +} +void canny::calcMagnitude_gpu(const oclMat& dx, const oclMat& dy, oclMat& mag, int rows, int cols, bool L2Grad) +{ + Context *clCxt = dx.clCxt; + string kernelName = "calcMagnitude"; + vector< pair > args; + + args.push_back( make_pair( sizeof(cl_mem), (void *)&dx.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&dy.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&mag.data)); + args.push_back( make_pair( sizeof(cl_int), (void *)&rows)); + args.push_back( make_pair( sizeof(cl_int), (void *)&cols)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&mag.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&mag.offset)); + + size_t globalThreads[3] = {cols, rows, 1}; + size_t localThreads[3] = {16, 16, 1}; + + char build_options [15] = ""; + if(L2Grad) + { + strcat(build_options, "-D L2GRAD"); + } + openCLExecuteKernel(clCxt, &imgproc_canny, kernelName, globalThreads, localThreads, args, -1, -1, build_options); +} + +void canny::calcMap_gpu(oclMat& dx, oclMat& dy, oclMat& mag, oclMat& map, int rows, int cols, float low_thresh, float high_thresh) +{ + Context *clCxt = dx.clCxt; + + vector< pair > args; + + args.push_back( make_pair( sizeof(cl_mem), (void *)&dx.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&dy.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&mag.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&map.data)); + args.push_back( make_pair( sizeof(cl_int), (void *)&rows)); + args.push_back( make_pair( sizeof(cl_int), (void *)&cols)); + args.push_back( make_pair( sizeof(cl_float), (void *)&low_thresh)); + args.push_back( make_pair( sizeof(cl_float), (void *)&high_thresh)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dx.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dy.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&mag.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&mag.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&map.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&map.offset)); + +#if CALCMAP_FIXED + size_t globalThreads[3] = {cols, rows, 1}; + string kernelName = "calcMap"; + size_t localThreads[3] = {16, 16, 1}; +#else + size_t globalThreads[3] = {cols, rows, 1}; + string kernelName = "calcMap_2"; + size_t localThreads[3] = {256, 1, 1}; +#endif + openCLExecuteKernel(clCxt, &imgproc_canny, kernelName, globalThreads, localThreads, args, -1, -1); +} + +void canny::edgesHysteresisLocal_gpu(oclMat& map, oclMat& st1, oclMat& counter, int rows, int cols) +{ + Context *clCxt = map.clCxt; + string kernelName = "edgesHysteresisLocal"; + vector< pair > args; + + args.push_back( make_pair( sizeof(cl_mem), (void *)&map.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&st1.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&counter.data)); + args.push_back( make_pair( sizeof(cl_int), (void *)&rows)); + args.push_back( make_pair( sizeof(cl_int), (void *)&cols)); + args.push_back( make_pair( sizeof(cl_int), (void *)&map.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&map.offset)); + + size_t globalThreads[3] = {cols, rows, 1}; + size_t localThreads[3] = {16, 16, 1}; + + openCLExecuteKernel(clCxt, &imgproc_canny, kernelName, globalThreads, localThreads, args, -1, -1); +} + +void canny::edgesHysteresisGlobal_gpu(oclMat& map, oclMat& st1, oclMat& st2, oclMat& counter, int rows, int cols) +{ + unsigned int count = Mat(counter).at(0); + + Context *clCxt = map.clCxt; + string kernelName = "edgesHysteresisGlobal"; + vector< pair > args; + size_t localThreads[3] = {128, 1, 1}; + +#define DIVUP(a, b) ((a)+(b)-1)/(b) + + while(count > 0) + { + counter.setTo(0); + args.clear(); + size_t globalThreads[3] = {std::min(count, 65535u) * 128, DIVUP(count, 65535), 1}; + args.push_back( make_pair( sizeof(cl_mem), (void *)&map.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&st1.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&st2.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&counter.data)); + args.push_back( make_pair( sizeof(cl_int), (void *)&rows)); + args.push_back( make_pair( sizeof(cl_int), (void *)&cols)); + args.push_back( make_pair( sizeof(cl_int), (void *)&count)); + args.push_back( make_pair( sizeof(cl_int), (void *)&map.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&map.offset)); + + openCLExecuteKernel(clCxt, &imgproc_canny, kernelName, globalThreads, localThreads, args, -1, -1); + count = Mat(counter).at(0); + std::swap(st1, st2); + } +#undef DIVUP +} + +void canny::getEdges_gpu(oclMat& map, oclMat& dst, int rows, int cols) +{ + Context *clCxt = map.clCxt; + string kernelName = "getEdges"; + vector< pair > args; + + args.push_back( make_pair( sizeof(cl_mem), (void *)&map.data)); + args.push_back( make_pair( sizeof(cl_mem), (void *)&dst.data)); + args.push_back( make_pair( sizeof(cl_int), (void *)&rows)); + args.push_back( make_pair( sizeof(cl_int), (void *)&cols)); + args.push_back( make_pair( sizeof(cl_int), (void *)&map.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&map.offset)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dst.step)); + args.push_back( make_pair( sizeof(cl_int), (void *)&dst.offset)); + + size_t globalThreads[3] = {cols, rows, 1}; + size_t localThreads[3] = {16, 16, 1}; + + openCLExecuteKernel(clCxt, &imgproc_canny, kernelName, globalThreads, localThreads, args, -1, -1); +} + +#endif // HAVE_OPENCL diff --git a/modules/ocl/src/kernels/imgproc_canny.cl b/modules/ocl/src/kernels/imgproc_canny.cl new file mode 100644 index 0000000000..da64d157f5 --- /dev/null +++ b/modules/ocl/src/kernels/imgproc_canny.cl @@ -0,0 +1,798 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other oclMaterials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#pragma OPENCL EXTENSION cl_amd_printf : enable +#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable +#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable + +#ifdef L2GRAD +inline float calc(int x, int y) +{ + return sqrt((float)(x * x + y * y)); +} +#else +inline float calc(int x, int y) +{ + return (float)abs(x) + abs(y); +} +#endif // + +// Smoothing perpendicular to the derivative direction with a triangle filter +// only support 3x3 Sobel kernel +// h (-1) = 1, h (0) = 2, h (1) = 1 +// h'(-1) = -1, h'(0) = 0, h'(1) = 1 +// thus sobel 2D operator can be calculated as: +// h'(x, y) = h'(x)h(y) for x direction +// +// src input 8bit single channel image data +// dx_buf output dx buffer +// dy_buf output dy buffer +__kernel +void calcSobelRowPass +( + __global const uchar * src, + __global int * dx_buf, + __global int * dy_buf, + int rows, + int cols, + int src_step, + int src_offset, + int dx_buf_step, + int dx_buf_offset, + int dy_buf_step, + int dy_buf_offset +) +{ + //src_step /= sizeof(*src); + //src_offset /= sizeof(*src); + dx_buf_step /= sizeof(*dx_buf); + dx_buf_offset /= sizeof(*dx_buf); + dy_buf_step /= sizeof(*dy_buf); + dy_buf_offset /= sizeof(*dy_buf); + + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + int lidx = get_local_id(0); + int lidy = get_local_id(1); + + __local int smem[16][18]; + + if(gidy < rows) + { + smem[lidy][lidx + 1] = src[gidx + gidy * src_step + src_offset]; + if(lidx == 0) + { + smem[lidy][0] = src[max(gidx - 1, 0) + gidy * src_step + src_offset]; + smem[lidy][17] = src[min(gidx + 16, cols - 1) + gidy * src_step + src_offset]; + } + barrier(CLK_LOCAL_MEM_FENCE); + if(gidx < cols) + { + dx_buf[gidx + gidy * dx_buf_step + dx_buf_offset] = + -smem[lidy][lidx] + smem[lidy][lidx + 2]; + dy_buf[gidx + gidy * dy_buf_step + dy_buf_offset] = + smem[lidy][lidx] + 2 * smem[lidy][lidx + 1] + smem[lidy][lidx + 2]; + } + } +} + +// calculate the magnitude of the filter pass combining both x and y directions +// This is the buffered version(3x3 sobel) +// +// dx_buf dx buffer, calculated from calcSobelRowPass +// dy_buf dy buffer, calculated from calcSobelRowPass +// dx direvitive in x direction output +// dy direvitive in y direction output +// mag magnitude direvitive of xy output +__kernel +void calcMagnitude_buf +( + __global const int * dx_buf, + __global const int * dy_buf, + __global int * dx, + __global int * dy, + __global float * mag, + int rows, + int cols, + int dx_buf_step, + int dx_buf_offset, + int dy_buf_step, + int dy_buf_offset, + int dx_step, + int dx_offset, + int dy_step, + int dy_offset, + int mag_step, + int mag_offset +) +{ + dx_buf_step /= sizeof(*dx_buf); + dx_buf_offset /= sizeof(*dx_buf); + dy_buf_step /= sizeof(*dy_buf); + dy_buf_offset /= sizeof(*dy_buf); + dx_step /= sizeof(*dx); + dx_offset /= sizeof(*dx); + dy_step /= sizeof(*dy); + dy_offset /= sizeof(*dy); + mag_step /= sizeof(*mag); + mag_offset /= sizeof(*mag); + + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + int lidx = get_local_id(0); + int lidy = get_local_id(1); + + __local int sdx[18][16]; + __local int sdy[18][16]; + + if(gidx < cols) + { + sdx[lidy + 1][lidx] = dx_buf[gidx + gidy * dx_buf_step + dx_buf_offset]; + sdy[lidy + 1][lidx] = dy_buf[gidx + gidy * dy_buf_step + dy_buf_offset]; + if(lidy == 0) + { + sdx[0][lidx] = dx_buf[gidx + max(gidy - 1, 0) * dx_buf_step + dx_buf_offset]; + sdx[17][lidx] = dx_buf[gidx + min(gidy + 16, rows - 1) * dx_buf_step + dx_buf_offset]; + + sdy[0][lidx] = dy_buf[gidx + max(gidy - 1, 0) * dy_buf_step + dy_buf_offset]; + sdy[17][lidx] = dy_buf[gidx + min(gidy + 16, rows - 1) * dy_buf_step + dy_buf_offset]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + if(gidy < rows) + { + int x = sdx[lidy][lidx] + 2 * sdx[lidy + 1][lidx] + sdx[lidy + 2][lidx]; + int y = -sdy[lidy][lidx] + sdy[lidy + 2][lidx]; + + dx[gidx + gidy * dx_step + dx_offset] = x; + dy[gidx + gidy * dy_step + dy_offset] = y; + + mag[(gidx + 1) + (gidy + 1) * mag_step + mag_offset] = calc(x, y); + } + } +} + +// calculate the magnitude of the filter pass combining both x and y directions +// This is the non-buffered version(non-3x3 sobel) +// +// dx_buf dx buffer, calculated from calcSobelRowPass +// dy_buf dy buffer, calculated from calcSobelRowPass +// dx direvitive in x direction output +// dy direvitive in y direction output +// mag magnitude direvitive of xy output +__kernel +void calcMagnitude +( + __global const int * dx, + __global const int * dy, + __global float * mag, + int rows, + int cols, + int dx_step, + int dx_offset, + int dy_step, + int dy_offset, + int mag_step, + int mag_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); + + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + if(gidy < rows && gidx < cols) + { + mag[(gidx + 1) + (gidy + 1) * mag_step + mag_offset] = + calc( + dx[gidx + gidy * dx_step + dx_offset], + dy[gidx + gidy * dy_step + dy_offset] + ); + } +} + +////////////////////////////////////////////////////////////////////////////////////////// +// 0.4142135623730950488016887242097 is tan(22.5) +#define CANNY_SHIFT 15 +#define TG22 (int)(0.4142135623730950488016887242097*(1< 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; + } +} + +// 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 + +////////////////////////////////////////////////////////////////////////////////////////// +// do Hysteresis for pixel whose edge type is 1 +// +// If candidate pixel (edge type is 1) has a neighbour pixel (in 3x3 area) with type 2, it is believed to be part of an edge and +// marked as edge. Each thread will iterate for 16 times to connect local edges. +// Candidate pixel being identified as edge will then be tested if there is nearby potiential edge points. If there is, counter will +// be incremented by 1 and the point location is stored. These potiential candidates will be processed further in next kernel. +// +// map raw edge type results calculated from calcMap. +// st the potiential edge points found in this kernel call +// counter the number of potiential edge points +__kernel +void edgesHysteresisLocal +( + __global int * map, + __global ushort2 * st, + volatile __global unsigned int * counter, + int rows, + int cols, + int map_step, + int map_offset +) +{ + map_step /= sizeof(*map); + map_offset /= sizeof(*map); + + __local int smem[18][18]; + + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + int lidx = get_local_id(0); + int lidy = get_local_id(1); + + int grp_idx = get_global_id(0) & 0xFFFFF0; + int grp_idy = get_global_id(1) & 0xFFFFF0; + + int tid = lidx + lidy * 16; + int lx = tid % 18; + int ly = tid / 18; + if(ly < 14) + { + smem[ly][lx] = map[grp_idx + lx + (grp_idy + ly) * map_step + map_offset]; + } + if(ly < 4 && grp_idy + ly + 14 <= rows && grp_idx + lx <= cols) + { + smem[ly + 14][lx] = map[grp_idx + lx + (grp_idy + ly + 14) * map_step + map_offset]; + } + + barrier(CLK_LOCAL_MEM_FENCE); + + if(gidy < rows && gidx < cols) + { + int n; + + #pragma unroll + for (int k = 0; k < 16; ++k) + { + n = 0; + + if (smem[lidy + 1][lidx + 1] == 1) + { + n += smem[lidy ][lidx ] == 2; + n += smem[lidy ][lidx + 1] == 2; + n += smem[lidy ][lidx + 2] == 2; + + n += smem[lidy + 1][lidx ] == 2; + n += smem[lidy + 1][lidx + 2] == 2; + + n += smem[lidy + 2][lidx ] == 2; + n += smem[lidy + 2][lidx + 1] == 2; + n += smem[lidy + 2][lidx + 2] == 2; + } + + if (n > 0) + smem[lidy + 1][lidx + 1] = 2; + } + + const int e = smem[lidy + 1][lidx + 1]; + map[gidx + 1 + (gidy + 1) * map_step] = e; + + n = 0; + if(e == 2) + { + n += smem[lidy ][lidx ] == 1; + n += smem[lidy ][lidx + 1] == 1; + n += smem[lidy ][lidx + 2] == 1; + + n += smem[lidy + 1][lidx ] == 1; + n += smem[lidy + 1][lidx + 2] == 1; + + n += smem[lidy + 2][lidx ] == 1; + n += smem[lidy + 2][lidx + 1] == 1; + n += smem[lidy + 2][lidx + 2] == 1; + } + + if(n > 0) + { + unsigned int ind = atomic_inc(counter); + st[ind] = (ushort2)(gidx + 1, gidy + 1); + } + } +} + +__constant int c_dx[8] = {-1, 0, 1, -1, 1, -1, 0, 1}; +__constant c_dy[8] = {-1, -1, -1, 0, 0, 1, 1, 1}; + +#define stack_size 512 +__kernel +void edgesHysteresisGlobal +( + __global int * map, + __global ushort2 * st1, + __global ushort2 * st2, + volatile __global int * counter, + int rows, + int cols, + int count, + int map_step, + int map_offset +) +{ + + map_step /= sizeof(*map); + map_offset /= sizeof(*map); + + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + int lidx = get_local_id(0); + int lidy = get_local_id(1); + + int grp_idx = get_group_id(0); + int grp_idy = get_group_id(1); + + volatile __local unsigned int s_counter; + __local unsigned int s_ind; + + __local ushort2 s_st[stack_size]; + + if(lidx == 0) + { + s_counter = 0; + } + barrier(CLK_LOCAL_MEM_FENCE); + + int ind = grp_idy * get_num_groups(0) + grp_idx; + + if(ind < count) + { + ushort2 pos = st1[ind]; + if (pos.x > 0 && pos.x <= cols && pos.y > 0 && pos.y <= rows) + { + if (lidx < 8) + { + pos.x += c_dx[lidx]; + pos.y += c_dy[lidx]; + + if (map[pos.x + pos.y * map_step] == 1) + { + map[pos.x + pos.y * map_step] = 2; + + ind = atomic_inc(&s_counter); + + s_st[ind] = pos; + } + } + barrier(CLK_LOCAL_MEM_FENCE); + + while (s_counter > 0 && s_counter <= stack_size - get_num_groups(0)) + { + const int subTaskIdx = lidx >> 3; + const int portion = min(s_counter, get_num_groups(0) >> 3); + + pos.x = pos.y = 0; + + if (subTaskIdx < portion) + pos = s_st[s_counter - 1 - subTaskIdx]; + barrier(CLK_LOCAL_MEM_FENCE); + + if (lidx == 0) + s_counter -= portion; + barrier(CLK_LOCAL_MEM_FENCE); + + if (pos.x > 0 && pos.x <= cols && pos.y > 0 && pos.y <= rows) + { + pos.x += c_dx[lidx & 7]; + pos.y += c_dy[lidx & 7]; + + if (map[pos.x + map_offset + pos.y * map_step] == 1) + { + map[pos.x + map_offset + pos.y * map_step] = 2; + + ind = atomic_inc(&s_counter); + + s_st[ind] = pos; + } + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + if (s_counter > 0) + { + if (lidx == 0) + { + ind = atomic_add(counter, s_counter); + s_ind = ind - s_counter; + } + barrier(CLK_LOCAL_MEM_FENCE); + + ind = s_ind; + + for (int i = lidx; i < s_counter; i += get_num_groups(0)) + { + st2[ind + i] = s_st[i]; + } + } + } + } +} +#undef stack_size + +//Get the edge result. egde type of value 2 will be marked as an edge point and set to 255. Otherwise 0. +// map edge type mappings +// dst edge output +__kernel +void getEdges +( + __global const int * map, + __global uchar * dst, + int rows, + int cols, + int map_step, + int map_offset, + int dst_step, + int dst_offset +) +{ + map_step /= sizeof(*map); + map_offset /= sizeof(*map); + //dst_step /= sizeof(*dst); + //dst_offset /= sizeof(*dst); + + int gidx = get_global_id(0); + int gidy = get_global_id(1); + + if(gidy < rows && gidx < cols) + { + //dst[gidx + gidy * dst_step] = map[gidx + 1 + (gidy + 1) * map_step] == 2 ? 255: 0; + dst[gidx + gidy * dst_step] = (uchar)(-(map[gidx + 1 + (gidy + 1) * map_step] / 2)); + } +} diff --git a/modules/ocl/test/test_canny.cpp b/modules/ocl/test/test_canny.cpp new file mode 100644 index 0000000000..c6aae6346c --- /dev/null +++ b/modules/ocl/test/test_canny.cpp @@ -0,0 +1,112 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other oclMaterials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "precomp.hpp" + +#ifdef WIN32 +#define FILTER_IMAGE "C:/Users/Public/Pictures/Sample Pictures/Penguins.jpg" +#else +#define FILTER_IMAGE "/Users/Test/Valve_original.PNG" // user need to specify a valid image path +#endif +#define SHOW_RESULT 0 + +//////////////////////////////////////////////////////// +// Canny + +IMPLEMENT_PARAM_CLASS(AppertureSize, int); +IMPLEMENT_PARAM_CLASS(L2gradient, bool); + +PARAM_TEST_CASE(Canny, AppertureSize, L2gradient) +{ + int apperture_size; + 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); + } +}; + +TEST_P(Canny, Accuracy) +{ + cv::Mat img = readImage(FILTER_IMAGE, cv::IMREAD_GRAYSCALE); + ASSERT_FALSE(img.empty()); + + double low_thresh = 50.0; + double high_thresh = 100.0; + + cv::resize(img, img, cv::Size(512, 384)); + cv::ocl::oclMat ocl_img = cv::ocl::oclMat(img); + + 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(ocl_ImgProc, Canny, testing::Combine( + testing::Values(AppertureSize(3), AppertureSize(5)), + testing::Values(L2gradient(false), L2gradient(true))));