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_CUDAARITHM) 49 50 Ptr<GeneralizedHoughBallard> cv::cuda::createGeneralizedHoughBallard() { throw_no_cuda(); return Ptr<GeneralizedHoughBallard>(); } 51 52 Ptr<GeneralizedHoughGuil> cv::cuda::createGeneralizedHoughGuil() { throw_no_cuda(); return Ptr<GeneralizedHoughGuil>(); } 53 54 #else /* !defined (HAVE_CUDA) */ 55 56 namespace cv { namespace cuda { namespace device 57 { 58 namespace ght 59 { 60 template <typename T> 61 int buildEdgePointList_gpu(PtrStepSzb edges, PtrStepSzb dx, PtrStepSzb dy, unsigned int* coordList, float* thetaList); 62 void buildRTable_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount, 63 PtrStepSz<short2> r_table, int* r_sizes, 64 short2 templCenter, int levels); 65 66 void Ballard_Pos_calcHist_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount, 67 PtrStepSz<short2> r_table, const int* r_sizes, 68 PtrStepSzi hist, 69 float dp, int levels); 70 int Ballard_Pos_findPosInHist_gpu(PtrStepSzi hist, float4* out, int3* votes, int maxSize, float dp, int threshold); 71 72 void Guil_Full_setTemplFeatures(PtrStepb p1_pos, PtrStepb p1_theta, PtrStepb p2_pos, PtrStepb d12, PtrStepb r1, PtrStepb r2); 73 void Guil_Full_setImageFeatures(PtrStepb p1_pos, PtrStepb p1_theta, PtrStepb p2_pos, PtrStepb d12, PtrStepb r1, PtrStepb r2); 74 void Guil_Full_buildTemplFeatureList_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount, 75 int* sizes, int maxSize, 76 float xi, float angleEpsilon, int levels, 77 float2 center, float maxDist); 78 void Guil_Full_buildImageFeatureList_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount, 79 int* sizes, int maxSize, 80 float xi, float angleEpsilon, int levels, 81 float2 center, float maxDist); 82 void Guil_Full_calcOHist_gpu(const int* templSizes, const int* imageSizes, int* OHist, 83 float minAngle, float maxAngle, float angleStep, int angleRange, 84 int levels, int tMaxSize); 85 void Guil_Full_calcSHist_gpu(const int* templSizes, const int* imageSizes, int* SHist, 86 float angle, float angleEpsilon, 87 float minScale, float maxScale, float iScaleStep, int scaleRange, 88 int levels, int tMaxSize); 89 void Guil_Full_calcPHist_gpu(const int* templSizes, const int* imageSizes, PtrStepSzi PHist, 90 float angle, float angleEpsilon, float scale, 91 float dp, 92 int levels, int tMaxSize); 93 int Guil_Full_findPosInHist_gpu(PtrStepSzi hist, float4* out, int3* votes, int curSize, int maxSize, 94 float angle, int angleVotes, float scale, int scaleVotes, 95 float dp, int threshold); 96 } 97 }}} 98 99 // common 100 101 namespace 102 { 103 class GeneralizedHoughBase 104 { 105 protected: 106 GeneralizedHoughBase(); 107 virtual ~GeneralizedHoughBase() {} 108 109 void setTemplateImpl(InputArray templ, Point templCenter); 110 void setTemplateImpl(InputArray edges, InputArray dx, InputArray dy, Point templCenter); 111 112 void detectImpl(InputArray image, OutputArray positions, OutputArray votes); 113 void detectImpl(InputArray edges, InputArray dx, InputArray dy, OutputArray positions, OutputArray votes); 114 115 void buildEdgePointList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy); 116 117 virtual void processTempl() = 0; 118 virtual void processImage() = 0; 119 120 int cannyLowThresh_; 121 int cannyHighThresh_; 122 double minDist_; 123 double dp_; 124 int maxBufferSize_; 125 126 Size templSize_; 127 Point templCenter_; 128 GpuMat templEdges_; 129 GpuMat templDx_; 130 GpuMat templDy_; 131 132 Size imageSize_; 133 GpuMat imageEdges_; 134 GpuMat imageDx_; 135 GpuMat imageDy_; 136 137 GpuMat edgePointList_; 138 139 GpuMat outBuf_; 140 int posCount_; 141 142 private: 143 #ifdef HAVE_OPENCV_CUDAFILTERS 144 void calcEdges(InputArray src, GpuMat& edges, GpuMat& dx, GpuMat& dy); 145 #endif 146 147 void filterMinDist(); 148 void convertTo(OutputArray positions, OutputArray votes); 149 150 #ifdef HAVE_OPENCV_CUDAFILTERS 151 Ptr<cuda::CannyEdgeDetector> canny_; 152 Ptr<cuda::Filter> filterDx_; 153 Ptr<cuda::Filter> filterDy_; 154 #endif 155 156 std::vector<float4> oldPosBuf_; 157 std::vector<int3> oldVoteBuf_; 158 std::vector<float4> newPosBuf_; 159 std::vector<int3> newVoteBuf_; 160 std::vector<int> indexies_; 161 }; 162 163 GeneralizedHoughBase::GeneralizedHoughBase() 164 { 165 cannyLowThresh_ = 50; 166 cannyHighThresh_ = 100; 167 minDist_ = 1.0; 168 dp_ = 1.0; 169 170 maxBufferSize_ = 10000; 171 172 #ifdef HAVE_OPENCV_CUDAFILTERS 173 canny_ = cuda::createCannyEdgeDetector(cannyLowThresh_, cannyHighThresh_); 174 filterDx_ = cuda::createSobelFilter(CV_8UC1, CV_32S, 1, 0); 175 filterDy_ = cuda::createSobelFilter(CV_8UC1, CV_32S, 0, 1); 176 #endif 177 } 178 179 #ifdef HAVE_OPENCV_CUDAFILTERS 180 void GeneralizedHoughBase::calcEdges(InputArray _src, GpuMat& edges, GpuMat& dx, GpuMat& dy) 181 { 182 GpuMat src = _src.getGpuMat(); 183 184 CV_Assert( src.type() == CV_8UC1 ); 185 CV_Assert( cannyLowThresh_ > 0 && cannyLowThresh_ < cannyHighThresh_ ); 186 187 ensureSizeIsEnough(src.size(), CV_32SC1, dx); 188 ensureSizeIsEnough(src.size(), CV_32SC1, dy); 189 190 filterDx_->apply(src, dx); 191 filterDy_->apply(src, dy); 192 193 ensureSizeIsEnough(src.size(), CV_8UC1, edges); 194 195 canny_->setLowThreshold(cannyLowThresh_); 196 canny_->setHighThreshold(cannyHighThresh_); 197 canny_->detect(dx, dy, edges); 198 } 199 #endif 200 201 void GeneralizedHoughBase::setTemplateImpl(InputArray templ, Point templCenter) 202 { 203 #ifndef HAVE_OPENCV_CUDAFILTERS 204 (void) templ; 205 (void) templCenter; 206 throw_no_cuda(); 207 #else 208 calcEdges(templ, templEdges_, templDx_, templDy_); 209 210 if (templCenter == Point(-1, -1)) 211 templCenter = Point(templEdges_.cols / 2, templEdges_.rows / 2); 212 213 templSize_ = templEdges_.size(); 214 templCenter_ = templCenter; 215 216 processTempl(); 217 #endif 218 } 219 220 void GeneralizedHoughBase::setTemplateImpl(InputArray edges, InputArray dx, InputArray dy, Point templCenter) 221 { 222 edges.getGpuMat().copyTo(templEdges_); 223 dx.getGpuMat().copyTo(templDx_); 224 dy.getGpuMat().copyTo(templDy_); 225 226 CV_Assert( templEdges_.type() == CV_8UC1 ); 227 CV_Assert( templDx_.type() == CV_32FC1 && templDx_.size() == templEdges_.size() ); 228 CV_Assert( templDy_.type() == templDx_.type() && templDy_.size() == templEdges_.size() ); 229 230 if (templCenter == Point(-1, -1)) 231 templCenter = Point(templEdges_.cols / 2, templEdges_.rows / 2); 232 233 templSize_ = templEdges_.size(); 234 templCenter_ = templCenter; 235 236 processTempl(); 237 } 238 239 void GeneralizedHoughBase::detectImpl(InputArray image, OutputArray positions, OutputArray votes) 240 { 241 #ifndef HAVE_OPENCV_CUDAFILTERS 242 (void) image; 243 (void) positions; 244 (void) votes; 245 throw_no_cuda(); 246 #else 247 calcEdges(image, imageEdges_, imageDx_, imageDy_); 248 249 imageSize_ = imageEdges_.size(); 250 251 posCount_ = 0; 252 253 processImage(); 254 255 if (posCount_ == 0) 256 { 257 positions.release(); 258 if (votes.needed()) 259 votes.release(); 260 } 261 else 262 { 263 if (minDist_ > 1) 264 filterMinDist(); 265 convertTo(positions, votes); 266 } 267 #endif 268 } 269 270 void GeneralizedHoughBase::detectImpl(InputArray edges, InputArray dx, InputArray dy, OutputArray positions, OutputArray votes) 271 { 272 edges.getGpuMat().copyTo(imageEdges_); 273 dx.getGpuMat().copyTo(imageDx_); 274 dy.getGpuMat().copyTo(imageDy_); 275 276 CV_Assert( imageEdges_.type() == CV_8UC1 ); 277 CV_Assert( imageDx_.type() == CV_32FC1 && imageDx_.size() == imageEdges_.size() ); 278 CV_Assert( imageDy_.type() == imageDx_.type() && imageDy_.size() == imageEdges_.size() ); 279 280 imageSize_ = imageEdges_.size(); 281 282 posCount_ = 0; 283 284 processImage(); 285 286 if (posCount_ == 0) 287 { 288 positions.release(); 289 if (votes.needed()) 290 votes.release(); 291 } 292 else 293 { 294 if (minDist_ > 1) 295 filterMinDist(); 296 convertTo(positions, votes); 297 } 298 } 299 300 void GeneralizedHoughBase::buildEdgePointList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy) 301 { 302 using namespace cv::cuda::device::ght; 303 304 typedef int (*func_t)(PtrStepSzb edges, PtrStepSzb dx, PtrStepSzb dy, unsigned int* coordList, float* thetaList); 305 static const func_t funcs[] = 306 { 307 0, 308 0, 309 0, 310 buildEdgePointList_gpu<short>, 311 buildEdgePointList_gpu<int>, 312 buildEdgePointList_gpu<float>, 313 0 314 }; 315 316 CV_Assert( edges.type() == CV_8UC1 ); 317 CV_Assert( dx.size() == edges.size() ); 318 CV_Assert( dy.type() == dx.type() && dy.size() == edges.size() ); 319 320 const func_t func = funcs[dx.depth()]; 321 CV_Assert( func != 0 ); 322 323 edgePointList_.cols = (int) (edgePointList_.step / sizeof(int)); 324 ensureSizeIsEnough(2, edges.size().area(), CV_32SC1, edgePointList_); 325 326 edgePointList_.cols = func(edges, dx, dy, edgePointList_.ptr<unsigned int>(0), edgePointList_.ptr<float>(1)); 327 } 328 329 struct IndexCmp 330 { 331 const int3* aux; 332 333 explicit IndexCmp(const int3* _aux) : aux(_aux) {} 334 335 bool operator ()(int l1, int l2) const 336 { 337 return aux[l1].x > aux[l2].x; 338 } 339 }; 340 341 void GeneralizedHoughBase::filterMinDist() 342 { 343 oldPosBuf_.resize(posCount_); 344 oldVoteBuf_.resize(posCount_); 345 346 cudaSafeCall( cudaMemcpy(&oldPosBuf_[0], outBuf_.ptr(0), posCount_ * sizeof(float4), cudaMemcpyDeviceToHost) ); 347 cudaSafeCall( cudaMemcpy(&oldVoteBuf_[0], outBuf_.ptr(1), posCount_ * sizeof(int3), cudaMemcpyDeviceToHost) ); 348 349 indexies_.resize(posCount_); 350 for (int i = 0; i < posCount_; ++i) 351 indexies_[i] = i; 352 std::sort(indexies_.begin(), indexies_.end(), IndexCmp(&oldVoteBuf_[0])); 353 354 newPosBuf_.clear(); 355 newVoteBuf_.clear(); 356 newPosBuf_.reserve(posCount_); 357 newVoteBuf_.reserve(posCount_); 358 359 const int cellSize = cvRound(minDist_); 360 const int gridWidth = (imageSize_.width + cellSize - 1) / cellSize; 361 const int gridHeight = (imageSize_.height + cellSize - 1) / cellSize; 362 363 std::vector< std::vector<Point2f> > grid(gridWidth * gridHeight); 364 365 const double minDist2 = minDist_ * minDist_; 366 367 for (int i = 0; i < posCount_; ++i) 368 { 369 const int ind = indexies_[i]; 370 371 Point2f p(oldPosBuf_[ind].x, oldPosBuf_[ind].y); 372 373 bool good = true; 374 375 const int xCell = static_cast<int>(p.x / cellSize); 376 const int yCell = static_cast<int>(p.y / cellSize); 377 378 int x1 = xCell - 1; 379 int y1 = yCell - 1; 380 int x2 = xCell + 1; 381 int y2 = yCell + 1; 382 383 // boundary check 384 x1 = std::max(0, x1); 385 y1 = std::max(0, y1); 386 x2 = std::min(gridWidth - 1, x2); 387 y2 = std::min(gridHeight - 1, y2); 388 389 for (int yy = y1; yy <= y2; ++yy) 390 { 391 for (int xx = x1; xx <= x2; ++xx) 392 { 393 const std::vector<Point2f>& m = grid[yy * gridWidth + xx]; 394 395 for(size_t j = 0; j < m.size(); ++j) 396 { 397 const Point2f d = p - m[j]; 398 399 if (d.ddot(d) < minDist2) 400 { 401 good = false; 402 goto break_out; 403 } 404 } 405 } 406 } 407 408 break_out: 409 410 if(good) 411 { 412 grid[yCell * gridWidth + xCell].push_back(p); 413 414 newPosBuf_.push_back(oldPosBuf_[ind]); 415 newVoteBuf_.push_back(oldVoteBuf_[ind]); 416 } 417 } 418 419 posCount_ = static_cast<int>(newPosBuf_.size()); 420 cudaSafeCall( cudaMemcpy(outBuf_.ptr(0), &newPosBuf_[0], posCount_ * sizeof(float4), cudaMemcpyHostToDevice) ); 421 cudaSafeCall( cudaMemcpy(outBuf_.ptr(1), &newVoteBuf_[0], posCount_ * sizeof(int3), cudaMemcpyHostToDevice) ); 422 } 423 424 void GeneralizedHoughBase::convertTo(OutputArray positions, OutputArray votes) 425 { 426 ensureSizeIsEnough(1, posCount_, CV_32FC4, positions); 427 GpuMat(1, posCount_, CV_32FC4, outBuf_.ptr(0), outBuf_.step).copyTo(positions); 428 429 if (votes.needed()) 430 { 431 ensureSizeIsEnough(1, posCount_, CV_32FC3, votes); 432 GpuMat(1, posCount_, CV_32FC4, outBuf_.ptr(1), outBuf_.step).copyTo(votes); 433 } 434 } 435 } 436 437 // GeneralizedHoughBallard 438 439 namespace 440 { 441 class GeneralizedHoughBallardImpl : public GeneralizedHoughBallard, private GeneralizedHoughBase 442 { 443 public: 444 GeneralizedHoughBallardImpl(); 445 446 void setTemplate(InputArray templ, Point templCenter) { setTemplateImpl(templ, templCenter); } 447 void setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter) { setTemplateImpl(edges, dx, dy, templCenter); } 448 449 void detect(InputArray image, OutputArray positions, OutputArray votes) { detectImpl(image, positions, votes); } 450 void detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions, OutputArray votes) { detectImpl(edges, dx, dy, positions, votes); } 451 452 void setCannyLowThresh(int cannyLowThresh) { cannyLowThresh_ = cannyLowThresh; } 453 int getCannyLowThresh() const { return cannyLowThresh_; } 454 455 void setCannyHighThresh(int cannyHighThresh) { cannyHighThresh_ = cannyHighThresh; } 456 int getCannyHighThresh() const { return cannyHighThresh_; } 457 458 void setMinDist(double minDist) { minDist_ = minDist; } 459 double getMinDist() const { return minDist_; } 460 461 void setDp(double dp) { dp_ = dp; } 462 double getDp() const { return dp_; } 463 464 void setMaxBufferSize(int maxBufferSize) { maxBufferSize_ = maxBufferSize; } 465 int getMaxBufferSize() const { return maxBufferSize_; } 466 467 void setLevels(int levels) { levels_ = levels; } 468 int getLevels() const { return levels_; } 469 470 void setVotesThreshold(int votesThreshold) { votesThreshold_ = votesThreshold; } 471 int getVotesThreshold() const { return votesThreshold_; } 472 473 private: 474 void processTempl(); 475 void processImage(); 476 477 void calcHist(); 478 void findPosInHist(); 479 480 int levels_; 481 int votesThreshold_; 482 483 GpuMat r_table_; 484 GpuMat r_sizes_; 485 486 GpuMat hist_; 487 }; 488 489 GeneralizedHoughBallardImpl::GeneralizedHoughBallardImpl() 490 { 491 levels_ = 360; 492 votesThreshold_ = 100; 493 } 494 495 void GeneralizedHoughBallardImpl::processTempl() 496 { 497 using namespace cv::cuda::device::ght; 498 499 CV_Assert( levels_ > 0 ); 500 501 buildEdgePointList(templEdges_, templDx_, templDy_); 502 503 ensureSizeIsEnough(levels_ + 1, maxBufferSize_, CV_16SC2, r_table_); 504 ensureSizeIsEnough(1, levels_ + 1, CV_32SC1, r_sizes_); 505 r_sizes_.setTo(Scalar::all(0)); 506 507 if (edgePointList_.cols > 0) 508 { 509 buildRTable_gpu(edgePointList_.ptr<unsigned int>(0), edgePointList_.ptr<float>(1), edgePointList_.cols, 510 r_table_, r_sizes_.ptr<int>(), make_short2(templCenter_.x, templCenter_.y), levels_); 511 cuda::min(r_sizes_, maxBufferSize_, r_sizes_); 512 } 513 } 514 515 void GeneralizedHoughBallardImpl::processImage() 516 { 517 calcHist(); 518 findPosInHist(); 519 } 520 521 void GeneralizedHoughBallardImpl::calcHist() 522 { 523 using namespace cv::cuda::device::ght; 524 525 CV_Assert( levels_ > 0 && r_table_.rows == (levels_ + 1) && r_sizes_.cols == (levels_ + 1) ); 526 CV_Assert( dp_ > 0.0); 527 528 const double idp = 1.0 / dp_; 529 530 buildEdgePointList(imageEdges_, imageDx_, imageDy_); 531 532 ensureSizeIsEnough(cvCeil(imageSize_.height * idp) + 2, cvCeil(imageSize_.width * idp) + 2, CV_32SC1, hist_); 533 hist_.setTo(Scalar::all(0)); 534 535 if (edgePointList_.cols > 0) 536 { 537 Ballard_Pos_calcHist_gpu(edgePointList_.ptr<unsigned int>(0), edgePointList_.ptr<float>(1), edgePointList_.cols, 538 r_table_, r_sizes_.ptr<int>(), 539 hist_, 540 (float)dp_, levels_); 541 } 542 } 543 544 void GeneralizedHoughBallardImpl::findPosInHist() 545 { 546 using namespace cv::cuda::device::ght; 547 548 CV_Assert( votesThreshold_ > 0 ); 549 550 ensureSizeIsEnough(2, maxBufferSize_, CV_32FC4, outBuf_); 551 552 posCount_ = Ballard_Pos_findPosInHist_gpu(hist_, outBuf_.ptr<float4>(0), outBuf_.ptr<int3>(1), maxBufferSize_, (float)dp_, votesThreshold_); 553 } 554 } 555 556 Ptr<GeneralizedHoughBallard> cv::cuda::createGeneralizedHoughBallard() 557 { 558 return makePtr<GeneralizedHoughBallardImpl>(); 559 } 560 561 // GeneralizedHoughGuil 562 563 namespace 564 { 565 class GeneralizedHoughGuilImpl : public GeneralizedHoughGuil, private GeneralizedHoughBase 566 { 567 public: 568 GeneralizedHoughGuilImpl(); 569 570 void setTemplate(InputArray templ, Point templCenter) { setTemplateImpl(templ, templCenter); } 571 void setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter) { setTemplateImpl(edges, dx, dy, templCenter); } 572 573 void detect(InputArray image, OutputArray positions, OutputArray votes) { detectImpl(image, positions, votes); } 574 void detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions, OutputArray votes) { detectImpl(edges, dx, dy, positions, votes); } 575 576 void setCannyLowThresh(int cannyLowThresh) { cannyLowThresh_ = cannyLowThresh; } 577 int getCannyLowThresh() const { return cannyLowThresh_; } 578 579 void setCannyHighThresh(int cannyHighThresh) { cannyHighThresh_ = cannyHighThresh; } 580 int getCannyHighThresh() const { return cannyHighThresh_; } 581 582 void setMinDist(double minDist) { minDist_ = minDist; } 583 double getMinDist() const { return minDist_; } 584 585 void setDp(double dp) { dp_ = dp; } 586 double getDp() const { return dp_; } 587 588 void setMaxBufferSize(int maxBufferSize) { maxBufferSize_ = maxBufferSize; } 589 int getMaxBufferSize() const { return maxBufferSize_; } 590 591 void setXi(double xi) { xi_ = xi; } 592 double getXi() const { return xi_; } 593 594 void setLevels(int levels) { levels_ = levels; } 595 int getLevels() const { return levels_; } 596 597 void setAngleEpsilon(double angleEpsilon) { angleEpsilon_ = angleEpsilon; } 598 double getAngleEpsilon() const { return angleEpsilon_; } 599 600 void setMinAngle(double minAngle) { minAngle_ = minAngle; } 601 double getMinAngle() const { return minAngle_; } 602 603 void setMaxAngle(double maxAngle) { maxAngle_ = maxAngle; } 604 double getMaxAngle() const { return maxAngle_; } 605 606 void setAngleStep(double angleStep) { angleStep_ = angleStep; } 607 double getAngleStep() const { return angleStep_; } 608 609 void setAngleThresh(int angleThresh) { angleThresh_ = angleThresh; } 610 int getAngleThresh() const { return angleThresh_; } 611 612 void setMinScale(double minScale) { minScale_ = minScale; } 613 double getMinScale() const { return minScale_; } 614 615 void setMaxScale(double maxScale) { maxScale_ = maxScale; } 616 double getMaxScale() const { return maxScale_; } 617 618 void setScaleStep(double scaleStep) { scaleStep_ = scaleStep; } 619 double getScaleStep() const { return scaleStep_; } 620 621 void setScaleThresh(int scaleThresh) { scaleThresh_ = scaleThresh; } 622 int getScaleThresh() const { return scaleThresh_; } 623 624 void setPosThresh(int posThresh) { posThresh_ = posThresh; } 625 int getPosThresh() const { return posThresh_; } 626 627 private: 628 void processTempl(); 629 void processImage(); 630 631 double xi_; 632 int levels_; 633 double angleEpsilon_; 634 635 double minAngle_; 636 double maxAngle_; 637 double angleStep_; 638 int angleThresh_; 639 640 double minScale_; 641 double maxScale_; 642 double scaleStep_; 643 int scaleThresh_; 644 645 int posThresh_; 646 647 struct Feature 648 { 649 GpuMat p1_pos; 650 GpuMat p1_theta; 651 GpuMat p2_pos; 652 653 GpuMat d12; 654 655 GpuMat r1; 656 GpuMat r2; 657 658 GpuMat sizes; 659 int maxSize; 660 661 void create(int levels, int maxCapacity, bool isTempl); 662 }; 663 664 typedef void (*set_func_t)(PtrStepb p1_pos, PtrStepb p1_theta, PtrStepb p2_pos, PtrStepb d12, PtrStepb r1, PtrStepb r2); 665 typedef void (*build_func_t)(const unsigned int* coordList, const float* thetaList, int pointsCount, 666 int* sizes, int maxSize, 667 float xi, float angleEpsilon, int levels, 668 float2 center, float maxDist); 669 670 void buildFeatureList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, Feature& features, 671 set_func_t set_func, build_func_t build_func, bool isTempl, Point2d center = Point2d()); 672 673 void calcOrientation(); 674 void calcScale(double angle); 675 void calcPosition(double angle, int angleVotes, double scale, int scaleVotes); 676 677 Feature templFeatures_; 678 Feature imageFeatures_; 679 680 std::vector< std::pair<double, int> > angles_; 681 std::vector< std::pair<double, int> > scales_; 682 683 GpuMat hist_; 684 std::vector<int> h_buf_; 685 }; 686 687 double toRad(double a) 688 { 689 return a * CV_PI / 180.0; 690 } 691 692 double clampAngle(double a) 693 { 694 double res = a; 695 696 while (res > 360.0) 697 res -= 360.0; 698 while (res < 0) 699 res += 360.0; 700 701 return res; 702 } 703 704 bool angleEq(double a, double b, double eps = 1.0) 705 { 706 return (fabs(clampAngle(a - b)) <= eps); 707 } 708 709 GeneralizedHoughGuilImpl::GeneralizedHoughGuilImpl() 710 { 711 maxBufferSize_ = 1000; 712 713 xi_ = 90.0; 714 levels_ = 360; 715 angleEpsilon_ = 1.0; 716 717 minAngle_ = 0.0; 718 maxAngle_ = 360.0; 719 angleStep_ = 1.0; 720 angleThresh_ = 15000; 721 722 minScale_ = 0.5; 723 maxScale_ = 2.0; 724 scaleStep_ = 0.05; 725 scaleThresh_ = 1000; 726 727 posThresh_ = 100; 728 } 729 730 void GeneralizedHoughGuilImpl::processTempl() 731 { 732 using namespace cv::cuda::device::ght; 733 734 buildFeatureList(templEdges_, templDx_, templDy_, templFeatures_, 735 Guil_Full_setTemplFeatures, Guil_Full_buildTemplFeatureList_gpu, 736 true, templCenter_); 737 738 h_buf_.resize(templFeatures_.sizes.cols); 739 cudaSafeCall( cudaMemcpy(&h_buf_[0], templFeatures_.sizes.data, h_buf_.size() * sizeof(int), cudaMemcpyDeviceToHost) ); 740 templFeatures_.maxSize = *std::max_element(h_buf_.begin(), h_buf_.end()); 741 } 742 743 void GeneralizedHoughGuilImpl::processImage() 744 { 745 using namespace cv::cuda::device::ght; 746 747 CV_Assert( levels_ > 0 ); 748 CV_Assert( templFeatures_.sizes.cols == levels_ + 1 ); 749 CV_Assert( minAngle_ >= 0.0 && minAngle_ < maxAngle_ && maxAngle_ <= 360.0 ); 750 CV_Assert( angleStep_ > 0.0 && angleStep_ < 360.0 ); 751 CV_Assert( angleThresh_ > 0 ); 752 CV_Assert( minScale_ > 0.0 && minScale_ < maxScale_ ); 753 CV_Assert( scaleStep_ > 0.0 ); 754 CV_Assert( scaleThresh_ > 0 ); 755 CV_Assert( dp_ > 0.0 ); 756 CV_Assert( posThresh_ > 0 ); 757 758 const double iAngleStep = 1.0 / angleStep_; 759 const int angleRange = cvCeil((maxAngle_ - minAngle_) * iAngleStep); 760 761 const double iScaleStep = 1.0 / scaleStep_; 762 const int scaleRange = cvCeil((maxScale_ - minScale_) * iScaleStep); 763 764 const double idp = 1.0 / dp_; 765 const int histRows = cvCeil(imageSize_.height * idp); 766 const int histCols = cvCeil(imageSize_.width * idp); 767 768 ensureSizeIsEnough(histRows + 2, std::max(angleRange + 1, std::max(scaleRange + 1, histCols + 2)), CV_32SC1, hist_); 769 h_buf_.resize(std::max(angleRange + 1, scaleRange + 1)); 770 771 ensureSizeIsEnough(2, maxBufferSize_, CV_32FC4, outBuf_); 772 773 buildFeatureList(imageEdges_, imageDx_, imageDy_, imageFeatures_, 774 Guil_Full_setImageFeatures, Guil_Full_buildImageFeatureList_gpu, 775 false); 776 777 calcOrientation(); 778 779 for (size_t i = 0; i < angles_.size(); ++i) 780 { 781 const double angle = angles_[i].first; 782 const int angleVotes = angles_[i].second; 783 784 calcScale(angle); 785 786 for (size_t j = 0; j < scales_.size(); ++j) 787 { 788 const double scale = scales_[j].first; 789 const int scaleVotes = scales_[j].second; 790 791 calcPosition(angle, angleVotes, scale, scaleVotes); 792 } 793 } 794 } 795 796 void GeneralizedHoughGuilImpl::Feature::create(int levels, int maxCapacity, bool isTempl) 797 { 798 if (!isTempl) 799 { 800 ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC2, p1_pos); 801 ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC2, p2_pos); 802 } 803 804 ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC1, p1_theta); 805 806 ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC1, d12); 807 808 if (isTempl) 809 { 810 ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC2, r1); 811 ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC2, r2); 812 } 813 814 ensureSizeIsEnough(1, levels + 1, CV_32SC1, sizes); 815 sizes.setTo(Scalar::all(0)); 816 817 maxSize = 0; 818 } 819 820 void GeneralizedHoughGuilImpl::buildFeatureList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, Feature& features, 821 set_func_t set_func, build_func_t build_func, bool isTempl, Point2d center) 822 { 823 CV_Assert( levels_ > 0 ); 824 825 const double maxDist = sqrt((double) templSize_.width * templSize_.width + templSize_.height * templSize_.height) * maxScale_; 826 827 features.create(levels_, maxBufferSize_, isTempl); 828 set_func(features.p1_pos, features.p1_theta, features.p2_pos, features.d12, features.r1, features.r2); 829 830 buildEdgePointList(edges, dx, dy); 831 832 if (edgePointList_.cols > 0) 833 { 834 build_func(edgePointList_.ptr<unsigned int>(0), edgePointList_.ptr<float>(1), edgePointList_.cols, 835 features.sizes.ptr<int>(), maxBufferSize_, (float)xi_, (float)angleEpsilon_, levels_, make_float2((float)center.x, (float)center.y), (float)maxDist); 836 } 837 } 838 839 void GeneralizedHoughGuilImpl::calcOrientation() 840 { 841 using namespace cv::cuda::device::ght; 842 843 const double iAngleStep = 1.0 / angleStep_; 844 const int angleRange = cvCeil((maxAngle_ - minAngle_) * iAngleStep); 845 846 hist_.setTo(Scalar::all(0)); 847 Guil_Full_calcOHist_gpu(templFeatures_.sizes.ptr<int>(), imageFeatures_.sizes.ptr<int>(0), hist_.ptr<int>(), 848 (float)minAngle_, (float)maxAngle_, (float)angleStep_, angleRange, levels_, templFeatures_.maxSize); 849 cudaSafeCall( cudaMemcpy(&h_buf_[0], hist_.data, h_buf_.size() * sizeof(int), cudaMemcpyDeviceToHost) ); 850 851 angles_.clear(); 852 853 for (int n = 0; n < angleRange; ++n) 854 { 855 if (h_buf_[n] >= angleThresh_) 856 { 857 const double angle = minAngle_ + n * angleStep_; 858 angles_.push_back(std::make_pair(angle, h_buf_[n])); 859 } 860 } 861 } 862 863 void GeneralizedHoughGuilImpl::calcScale(double angle) 864 { 865 using namespace cv::cuda::device::ght; 866 867 const double iScaleStep = 1.0 / scaleStep_; 868 const int scaleRange = cvCeil((maxScale_ - minScale_) * iScaleStep); 869 870 hist_.setTo(Scalar::all(0)); 871 Guil_Full_calcSHist_gpu(templFeatures_.sizes.ptr<int>(), imageFeatures_.sizes.ptr<int>(0), hist_.ptr<int>(), 872 (float)angle, (float)angleEpsilon_, (float)minScale_, (float)maxScale_, 873 (float)iScaleStep, scaleRange, levels_, templFeatures_.maxSize); 874 cudaSafeCall( cudaMemcpy(&h_buf_[0], hist_.data, h_buf_.size() * sizeof(int), cudaMemcpyDeviceToHost) ); 875 876 scales_.clear(); 877 878 for (int s = 0; s < scaleRange; ++s) 879 { 880 if (h_buf_[s] >= scaleThresh_) 881 { 882 const double scale = minScale_ + s * scaleStep_; 883 scales_.push_back(std::make_pair(scale, h_buf_[s])); 884 } 885 } 886 } 887 888 void GeneralizedHoughGuilImpl::calcPosition(double angle, int angleVotes, double scale, int scaleVotes) 889 { 890 using namespace cv::cuda::device::ght; 891 892 hist_.setTo(Scalar::all(0)); 893 Guil_Full_calcPHist_gpu(templFeatures_.sizes.ptr<int>(), imageFeatures_.sizes.ptr<int>(0), hist_, 894 (float)angle, (float)angleEpsilon_, (float)scale, (float)dp_, levels_, templFeatures_.maxSize); 895 896 posCount_ = Guil_Full_findPosInHist_gpu(hist_, outBuf_.ptr<float4>(0), outBuf_.ptr<int3>(1), 897 posCount_, maxBufferSize_, (float)angle, angleVotes, 898 (float)scale, scaleVotes, (float)dp_, posThresh_); 899 } 900 } 901 902 Ptr<GeneralizedHoughGuil> cv::cuda::createGeneralizedHoughGuil() 903 { 904 return makePtr<GeneralizedHoughGuilImpl>(); 905 } 906 907 #endif /* !defined (HAVE_CUDA) */ 908
