• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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) || !defined(HAVE_OPENCV_CUDAFILTERS)
49 
createHoughCirclesDetector(float,float,int,int,int,int,int)50 Ptr<cuda::HoughCirclesDetector> cv::cuda::createHoughCirclesDetector(float, float, int, int, int, int, int) { throw_no_cuda(); return Ptr<HoughCirclesDetector>(); }
51 
52 #else /* !defined (HAVE_CUDA) */
53 
54 namespace cv { namespace cuda { namespace device
55 {
56     namespace hough
57     {
58         int buildPointList_gpu(PtrStepSzb src, unsigned int* list);
59     }
60 
61     namespace hough_circles
62     {
63         void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp);
64         int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold);
65         int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count,
66                                    float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20);
67     }
68 }}}
69 
70 namespace
71 {
72     class HoughCirclesDetectorImpl : public HoughCirclesDetector
73     {
74     public:
75         HoughCirclesDetectorImpl(float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles);
76 
77         void detect(InputArray src, OutputArray circles, Stream& stream);
78 
setDp(float dp)79         void setDp(float dp) { dp_ = dp; }
getDp() const80         float getDp() const { return dp_; }
81 
setMinDist(float minDist)82         void setMinDist(float minDist) { minDist_ = minDist; }
getMinDist() const83         float getMinDist() const { return minDist_; }
84 
setCannyThreshold(int cannyThreshold)85         void setCannyThreshold(int cannyThreshold) { cannyThreshold_ = cannyThreshold; }
getCannyThreshold() const86         int getCannyThreshold() const { return cannyThreshold_; }
87 
setVotesThreshold(int votesThreshold)88         void setVotesThreshold(int votesThreshold) { votesThreshold_ = votesThreshold; }
getVotesThreshold() const89         int getVotesThreshold() const { return votesThreshold_; }
90 
setMinRadius(int minRadius)91         void setMinRadius(int minRadius) { minRadius_ = minRadius; }
getMinRadius() const92         int getMinRadius() const { return minRadius_; }
93 
setMaxRadius(int maxRadius)94         void setMaxRadius(int maxRadius) { maxRadius_ = maxRadius; }
getMaxRadius() const95         int getMaxRadius() const { return maxRadius_; }
96 
setMaxCircles(int maxCircles)97         void setMaxCircles(int maxCircles) { maxCircles_ = maxCircles; }
getMaxCircles() const98         int getMaxCircles() const { return maxCircles_; }
99 
write(FileStorage & fs) const100         void write(FileStorage& fs) const
101         {
102             fs << "name" << "HoughCirclesDetector_CUDA"
103             << "dp" << dp_
104             << "minDist" << minDist_
105             << "cannyThreshold" << cannyThreshold_
106             << "votesThreshold" << votesThreshold_
107             << "minRadius" << minRadius_
108             << "maxRadius" << maxRadius_
109             << "maxCircles" << maxCircles_;
110         }
111 
read(const FileNode & fn)112         void read(const FileNode& fn)
113         {
114             CV_Assert( String(fn["name"]) == "HoughCirclesDetector_CUDA" );
115             dp_ = (float)fn["dp"];
116             minDist_ = (float)fn["minDist"];
117             cannyThreshold_ = (int)fn["cannyThreshold"];
118             votesThreshold_ = (int)fn["votesThreshold"];
119             minRadius_ = (int)fn["minRadius"];
120             maxRadius_ = (int)fn["maxRadius"];
121             maxCircles_ = (int)fn["maxCircles"];
122         }
123 
124     private:
125         float dp_;
126         float minDist_;
127         int cannyThreshold_;
128         int votesThreshold_;
129         int minRadius_;
130         int maxRadius_;
131         int maxCircles_;
132 
133         GpuMat dx_, dy_;
134         GpuMat edges_;
135         GpuMat accum_;
136         Mat tt; //CPU copy of accum_
137         GpuMat list_;
138         GpuMat result_;
139         Ptr<cuda::Filter> filterDx_;
140         Ptr<cuda::Filter> filterDy_;
141         Ptr<cuda::CannyEdgeDetector> canny_;
142     };
143 
centersCompare(Vec3f a,Vec3f b)144     bool centersCompare(Vec3f a, Vec3f b) {return (a[2] > b[2]);}
145 
HoughCirclesDetectorImpl(float dp,float minDist,int cannyThreshold,int votesThreshold,int minRadius,int maxRadius,int maxCircles)146     HoughCirclesDetectorImpl::HoughCirclesDetectorImpl(float dp, float minDist, int cannyThreshold, int votesThreshold,
147                                                        int minRadius, int maxRadius, int maxCircles) :
148         dp_(dp), minDist_(minDist), cannyThreshold_(cannyThreshold), votesThreshold_(votesThreshold),
149         minRadius_(minRadius), maxRadius_(maxRadius), maxCircles_(maxCircles)
150     {
151         canny_ = cuda::createCannyEdgeDetector(std::max(cannyThreshold_ / 2, 1), cannyThreshold_);
152 
153         filterDx_ = cuda::createSobelFilter(CV_8UC1, CV_32S, 1, 0);
154         filterDy_ = cuda::createSobelFilter(CV_8UC1, CV_32S, 0, 1);
155     }
156 
detect(InputArray _src,OutputArray circles,Stream & stream)157     void HoughCirclesDetectorImpl::detect(InputArray _src, OutputArray circles, Stream& stream)
158     {
159         // TODO : implement async version
160         (void) stream;
161 
162         using namespace cv::cuda::device::hough;
163         using namespace cv::cuda::device::hough_circles;
164 
165         GpuMat src = _src.getGpuMat();
166 
167         CV_Assert( src.type() == CV_8UC1 );
168         CV_Assert( src.cols < std::numeric_limits<unsigned short>::max() );
169         CV_Assert( src.rows < std::numeric_limits<unsigned short>::max() );
170         CV_Assert( dp_ > 0 );
171         CV_Assert( minRadius_ > 0 && maxRadius_ > minRadius_ );
172         CV_Assert( cannyThreshold_ > 0 );
173         CV_Assert( votesThreshold_ > 0 );
174         CV_Assert( maxCircles_ > 0 );
175 
176         const float idp = 1.0f / dp_;
177 
178         filterDx_->apply(src, dx_);
179         filterDy_->apply(src, dy_);
180 
181         canny_->setLowThreshold(std::max(cannyThreshold_ / 2, 1));
182         canny_->setHighThreshold(cannyThreshold_);
183 
184         canny_->detect(dx_, dy_, edges_);
185 
186         ensureSizeIsEnough(2, src.size().area(), CV_32SC1, list_);
187         unsigned int* srcPoints = list_.ptr<unsigned int>(0);
188         unsigned int* centers = list_.ptr<unsigned int>(1);
189 
190         const int pointsCount = buildPointList_gpu(edges_, srcPoints);
191         if (pointsCount == 0)
192         {
193             circles.release();
194             return;
195         }
196 
197         ensureSizeIsEnough(cvCeil(src.rows * idp) + 2, cvCeil(src.cols * idp) + 2, CV_32SC1, accum_);
198         accum_.setTo(Scalar::all(0));
199 
200         circlesAccumCenters_gpu(srcPoints, pointsCount, dx_, dy_, accum_, minRadius_, maxRadius_, idp);
201 
202         accum_.download(tt);
203 
204         int centersCount = buildCentersList_gpu(accum_, centers, votesThreshold_);
205         if (centersCount == 0)
206         {
207             circles.release();
208             return;
209         }
210 
211         if (minDist_ > 1)
212         {
213             AutoBuffer<ushort2> oldBuf_(centersCount);
214             AutoBuffer<ushort2> newBuf_(centersCount);
215             int newCount = 0;
216 
217             ushort2* oldBuf = oldBuf_;
218             ushort2* newBuf = newBuf_;
219 
220             cudaSafeCall( cudaMemcpy(oldBuf, centers, centersCount * sizeof(ushort2), cudaMemcpyDeviceToHost) );
221 
222             const int cellSize = cvRound(minDist_);
223             const int gridWidth = (src.cols + cellSize - 1) / cellSize;
224             const int gridHeight = (src.rows + cellSize - 1) / cellSize;
225 
226             std::vector< std::vector<ushort2> > grid(gridWidth * gridHeight);
227 
228             const float minDist2 = minDist_ * minDist_;
229 
230             std::vector<Vec3f> sortBuf;
231             for(int i=0; i<centersCount; i++){
232                 Vec3f temp;
233                 temp[0] = oldBuf[i].x;
234                 temp[1] = oldBuf[i].y;
235                 temp[2] = tt.at<int>(temp[1]+1, temp[0]+1);
236                 sortBuf.push_back(temp);
237             }
238             std::sort(sortBuf.begin(), sortBuf.end(), centersCompare);
239 
240             for (int i = 0; i < centersCount; ++i)
241             {
242                 ushort2 p;
243                 p.x = sortBuf[i][0];
244                 p.y = sortBuf[i][1];
245 
246                 bool good = true;
247 
248                 int xCell = static_cast<int>(p.x / cellSize);
249                 int yCell = static_cast<int>(p.y / cellSize);
250 
251                 int x1 = xCell - 1;
252                 int y1 = yCell - 1;
253                 int x2 = xCell + 1;
254                 int y2 = yCell + 1;
255 
256                 // boundary check
257                 x1 = std::max(0, x1);
258                 y1 = std::max(0, y1);
259                 x2 = std::min(gridWidth - 1, x2);
260                 y2 = std::min(gridHeight - 1, y2);
261 
262                 for (int yy = y1; yy <= y2; ++yy)
263                 {
264                     for (int xx = x1; xx <= x2; ++xx)
265                     {
266                         std::vector<ushort2>& m = grid[yy * gridWidth + xx];
267 
268                         for(size_t j = 0; j < m.size(); ++j)
269                         {
270                             float dx = (float)(p.x - m[j].x);
271                             float dy = (float)(p.y - m[j].y);
272 
273                             if (dx * dx + dy * dy < minDist2)
274                             {
275                                 good = false;
276                                 goto break_out;
277                             }
278                         }
279                     }
280                 }
281 
282                 break_out:
283 
284                 if(good)
285                 {
286                     grid[yCell * gridWidth + xCell].push_back(p);
287 
288                     newBuf[newCount++] = p;
289                 }
290             }
291 
292             cudaSafeCall( cudaMemcpy(centers, newBuf, newCount * sizeof(unsigned int), cudaMemcpyHostToDevice) );
293             centersCount = newCount;
294         }
295 
296         ensureSizeIsEnough(1, maxCircles_, CV_32FC3, result_);
297 
298         int circlesCount = circlesAccumRadius_gpu(centers, centersCount, srcPoints, pointsCount, result_.ptr<float3>(), maxCircles_,
299                                                   dp_, minRadius_, maxRadius_, votesThreshold_, deviceSupports(FEATURE_SET_COMPUTE_20));
300 
301         if (circlesCount == 0)
302         {
303             circles.release();
304             return;
305         }
306 
307         result_.cols = circlesCount;
308         result_.copyTo(circles);
309     }
310 }
311 
createHoughCirclesDetector(float dp,float minDist,int cannyThreshold,int votesThreshold,int minRadius,int maxRadius,int maxCircles)312 Ptr<HoughCirclesDetector> cv::cuda::createHoughCirclesDetector(float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles)
313 {
314     return makePtr<HoughCirclesDetectorImpl>(dp, minDist, cannyThreshold, votesThreshold, minRadius, maxRadius, maxCircles);
315 }
316 
317 #endif /* !defined (HAVE_CUDA) */
318