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 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 79 void setDp(float dp) { dp_ = dp; } 80 float getDp() const { return dp_; } 81 82 void setMinDist(float minDist) { minDist_ = minDist; } 83 float getMinDist() const { return minDist_; } 84 85 void setCannyThreshold(int cannyThreshold) { cannyThreshold_ = cannyThreshold; } 86 int getCannyThreshold() const { return cannyThreshold_; } 87 88 void setVotesThreshold(int votesThreshold) { votesThreshold_ = votesThreshold; } 89 int getVotesThreshold() const { return votesThreshold_; } 90 91 void setMinRadius(int minRadius) { minRadius_ = minRadius; } 92 int getMinRadius() const { return minRadius_; } 93 94 void setMaxRadius(int maxRadius) { maxRadius_ = maxRadius; } 95 int getMaxRadius() const { return maxRadius_; } 96 97 void setMaxCircles(int maxCircles) { maxCircles_ = maxCircles; } 98 int getMaxCircles() const { return maxCircles_; } 99 100 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 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 144 bool centersCompare(Vec3f a, Vec3f b) {return (a[2] > b[2]);} 145 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 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 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
