1 /*M///////////////////////////////////////////////////////////////////////////////////////
2 //
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4 //
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install,
7 // copy or use the software.
8 //
9 //
10 // License Agreement
11 // For Open Source Computer Vision Library
12 //
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
15 // Third party copyrights are property of their respective owners.
16 //
17 // Redistribution and use in source and binary forms, with or without modification,
18 // are permitted provided that the following conditions are met:
19 //
20 // * Redistribution's of source code must retain the above copyright notice,
21 // this list of conditions and the following disclaimer.
22 //
23 // * Redistribution's in binary form must reproduce the above copyright notice,
24 // this list of conditions and the following disclaimer in the documentation
25 // and/or other materials provided with the distribution.
26 //
27 // * The name of the copyright holders may not be used to endorse or promote products
28 // derived from this software without specific prior written permission.
29 //
30 // This software is provided by the copyright holders and contributors "as is" and
31 // any express or implied warranties, including, but not limited to, the implied
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
34 // indirect, incidental, special, exemplary, or consequential damages
35 // (including, but not limited to, procurement of substitute goods or services;
36 // loss of use, data, or profits; or business interruption) however caused
37 // and on any theory of liability, whether in contract, strict liability,
38 // or tort (including negligence or otherwise) arising in any way out of
39 // the use of this software, even if advised of the possibility of such damage.
40 //
41 //M*/
42
43 #include "precomp.hpp"
44
45 using namespace cv;
46 using namespace cv::cuda;
47
48 #if !defined (HAVE_CUDA) || defined (CUDA_DISABLER)
49
createCannyEdgeDetector(double,double,int,bool)50 Ptr<CannyEdgeDetector> cv::cuda::createCannyEdgeDetector(double, double, int, bool) { throw_no_cuda(); return Ptr<CannyEdgeDetector>(); }
51
52 #else /* !defined (HAVE_CUDA) */
53
54 namespace canny
55 {
56 void calcMagnitude(PtrStepSzb srcWhole, int xoff, int yoff, PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, bool L2Grad, cudaStream_t stream);
57 void calcMagnitude(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, bool L2Grad, cudaStream_t stream);
58
59 void calcMap(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, PtrStepSzi map, float low_thresh, float high_thresh, cudaStream_t stream);
60
61 void edgesHysteresisLocal(PtrStepSzi map, short2* st1, int* d_counter, cudaStream_t stream);
62
63 void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2, int* d_counter, cudaStream_t stream);
64
65 void getEdges(PtrStepSzi map, PtrStepSzb dst, cudaStream_t stream);
66 }
67
68 namespace
69 {
70 class CannyImpl : public CannyEdgeDetector
71 {
72 public:
CannyImpl(double low_thresh,double high_thresh,int apperture_size,bool L2gradient)73 CannyImpl(double low_thresh, double high_thresh, int apperture_size, bool L2gradient) :
74 low_thresh_(low_thresh), high_thresh_(high_thresh), apperture_size_(apperture_size), L2gradient_(L2gradient)
75 {
76 old_apperture_size_ = -1;
77 d_counter = nullptr;
78 }
79
80 void detect(InputArray image, OutputArray edges, Stream& stream);
81 void detect(InputArray dx, InputArray dy, OutputArray edges, Stream& stream);
82
setLowThreshold(double low_thresh)83 void setLowThreshold(double low_thresh) { low_thresh_ = low_thresh; }
getLowThreshold() const84 double getLowThreshold() const { return low_thresh_; }
85
setHighThreshold(double high_thresh)86 void setHighThreshold(double high_thresh) { high_thresh_ = high_thresh; }
getHighThreshold() const87 double getHighThreshold() const { return high_thresh_; }
88
setAppertureSize(int apperture_size)89 void setAppertureSize(int apperture_size) { apperture_size_ = apperture_size; }
getAppertureSize() const90 int getAppertureSize() const { return apperture_size_; }
91
setL2Gradient(bool L2gradient)92 void setL2Gradient(bool L2gradient) { L2gradient_ = L2gradient; }
getL2Gradient() const93 bool getL2Gradient() const { return L2gradient_; }
94
write(FileStorage & fs) const95 void write(FileStorage& fs) const
96 {
97 writeFormat(fs);
98 fs << "name" << "Canny_CUDA"
99 << "low_thresh" << low_thresh_
100 << "high_thresh" << high_thresh_
101 << "apperture_size" << apperture_size_
102 << "L2gradient" << L2gradient_;
103 }
104
read(const FileNode & fn)105 void read(const FileNode& fn)
106 {
107 CV_Assert( String(fn["name"]) == "Canny_CUDA" );
108 low_thresh_ = (double)fn["low_thresh"];
109 high_thresh_ = (double)fn["high_thresh"];
110 apperture_size_ = (int)fn["apperture_size"];
111 L2gradient_ = (int)fn["L2gradient"] != 0;
112 }
113
114 private:
115 void createBuf(Size image_size);
116 void CannyCaller(GpuMat& edges, Stream& stream);
117
118 double low_thresh_;
119 double high_thresh_;
120 int apperture_size_;
121 bool L2gradient_;
122
123 GpuMat dx_, dy_;
124 GpuMat mag_;
125 GpuMat map_;
126 GpuMat st1_, st2_;
127 #ifdef HAVE_OPENCV_CUDAFILTERS
128 Ptr<Filter> filterDX_, filterDY_;
129 #endif
130 int old_apperture_size_;
131
132 int *d_counter;
133 };
134
detect(InputArray _image,OutputArray _edges,Stream & stream)135 void CannyImpl::detect(InputArray _image, OutputArray _edges, Stream& stream)
136 {
137 GpuMat image = _image.getGpuMat();
138
139 CV_Assert( image.type() == CV_8UC1 );
140 CV_Assert( deviceSupports(SHARED_ATOMICS) );
141
142 if (low_thresh_ > high_thresh_)
143 std::swap(low_thresh_, high_thresh_);
144
145 createBuf(image.size());
146
147 _edges.create(image.size(), CV_8UC1);
148 GpuMat edges = _edges.getGpuMat();
149
150 if (apperture_size_ == 3)
151 {
152 Size wholeSize;
153 Point ofs;
154 image.locateROI(wholeSize, ofs);
155 GpuMat srcWhole(wholeSize, image.type(), image.datastart, image.step);
156
157 canny::calcMagnitude(srcWhole, ofs.x, ofs.y, dx_, dy_, mag_, L2gradient_, StreamAccessor::getStream(stream));
158 }
159 else
160 {
161 #ifndef HAVE_OPENCV_CUDAFILTERS
162 throw_no_cuda();
163 #else
164 filterDX_->apply(image, dx_, stream);
165 filterDY_->apply(image, dy_, stream);
166
167 canny::calcMagnitude(dx_, dy_, mag_, L2gradient_, StreamAccessor::getStream(stream));
168 #endif
169 }
170
171 CannyCaller(edges, stream);
172 }
173
detect(InputArray _dx,InputArray _dy,OutputArray _edges,Stream & stream)174 void CannyImpl::detect(InputArray _dx, InputArray _dy, OutputArray _edges, Stream& stream)
175 {
176 GpuMat dx = _dx.getGpuMat();
177 GpuMat dy = _dy.getGpuMat();
178
179 CV_Assert( dx.type() == CV_32SC1 );
180 CV_Assert( dy.type() == dx.type() && dy.size() == dx.size() );
181 CV_Assert( deviceSupports(SHARED_ATOMICS) );
182
183 dx.copyTo(dx_, stream);
184 dy.copyTo(dy_, stream);
185
186 if (low_thresh_ > high_thresh_)
187 std::swap(low_thresh_, high_thresh_);
188
189 createBuf(dx.size());
190
191 _edges.create(dx.size(), CV_8UC1);
192 GpuMat edges = _edges.getGpuMat();
193
194 canny::calcMagnitude(dx_, dy_, mag_, L2gradient_, StreamAccessor::getStream(stream));
195
196 CannyCaller(edges, stream);
197 }
198
createBuf(Size image_size)199 void CannyImpl::createBuf(Size image_size)
200 {
201 CV_Assert(image_size.width < std::numeric_limits<short>::max() && image_size.height < std::numeric_limits<short>::max());
202
203 ensureSizeIsEnough(image_size, CV_32SC1, dx_);
204 ensureSizeIsEnough(image_size, CV_32SC1, dy_);
205
206 #ifdef HAVE_OPENCV_CUDAFILTERS
207 if (apperture_size_ != 3 && apperture_size_ != old_apperture_size_)
208 {
209 filterDX_ = cuda::createDerivFilter(CV_8UC1, CV_32S, 1, 0, apperture_size_, false, 1, BORDER_REPLICATE);
210 filterDY_ = cuda::createDerivFilter(CV_8UC1, CV_32S, 0, 1, apperture_size_, false, 1, BORDER_REPLICATE);
211 old_apperture_size_ = apperture_size_;
212 }
213 #endif
214
215 ensureSizeIsEnough(image_size, CV_32FC1, mag_);
216 ensureSizeIsEnough(image_size, CV_32SC1, map_);
217
218 ensureSizeIsEnough(1, image_size.area(), CV_16SC2, st1_);
219 ensureSizeIsEnough(1, image_size.area(), CV_16SC2, st2_);
220 }
221
CannyCaller(GpuMat & edges,Stream & stream)222 void CannyImpl::CannyCaller(GpuMat& edges, Stream& stream)
223 {
224 map_.setTo(Scalar::all(0), stream);
225
226 canny::calcMap(dx_, dy_, mag_, map_, static_cast<float>(low_thresh_), static_cast<float>(high_thresh_), StreamAccessor::getStream(stream));
227
228 cudaSafeCall( cudaMalloc(&d_counter, sizeof(int)) );
229
230 canny::edgesHysteresisLocal(map_, st1_.ptr<short2>(), d_counter, StreamAccessor::getStream(stream));
231
232 canny::edgesHysteresisGlobal(map_, st1_.ptr<short2>(), st2_.ptr<short2>(), d_counter, StreamAccessor::getStream(stream));
233
234 cudaSafeCall( cudaFree(d_counter) );
235
236 canny::getEdges(map_, edges, StreamAccessor::getStream(stream));
237 }
238 }
239
createCannyEdgeDetector(double low_thresh,double high_thresh,int apperture_size,bool L2gradient)240 Ptr<CannyEdgeDetector> cv::cuda::createCannyEdgeDetector(double low_thresh, double high_thresh, int apperture_size, bool L2gradient)
241 {
242 return makePtr<CannyImpl>(low_thresh, high_thresh, apperture_size, L2gradient);
243 }
244
245 #endif /* !defined (HAVE_CUDA) */
246