modules/gpu/src/hough.cpp

   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 std;
  46 using namespace cv;
  47 using namespace cv::gpu;
  48
  49 #if !defined (HAVE_CUDA)
  50
  51 void cv::gpu::HoughLines(const GpuMat&, GpuMat&, float, float, int, bool, int) { throw_nogpu(); }
  52 void cv::gpu::HoughLines(const GpuMat&, GpuMat&, HoughLinesBuf&, float, float, int, bool, int) { throw_nogpu(); }
  53 void cv::gpu::HoughLinesDownload(const GpuMat&, OutputArray, OutputArray) { throw_nogpu(); }
  54
  55 void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, int, float, float, int, int, int, int, int) { throw_nogpu(); }
  56 void cv::gpu::HoughCircles(const GpuMat&, GpuMat&, HoughCirclesBuf&, int, float, float, int, int, int, int, int) { throw_nogpu(); }
  57 void cv::gpu::HoughCirclesDownload(const GpuMat&, OutputArray) { throw_nogpu(); }
  58
  59 Ptr<GeneralizedHough_GPU> cv::gpu::GeneralizedHough_GPU::create(int) { throw_nogpu(); return Ptr<GeneralizedHough_GPU>(); }
  60 cv::gpu::GeneralizedHough_GPU::~GeneralizedHough_GPU() {}
  61 void cv::gpu::GeneralizedHough_GPU::setTemplate(const GpuMat&, int, Point) { throw_nogpu(); }
  62 void cv::gpu::GeneralizedHough_GPU::setTemplate(const GpuMat&, const GpuMat&, const GpuMat&, Point) { throw_nogpu(); }
  63 void cv::gpu::GeneralizedHough_GPU::detect(const GpuMat&, GpuMat&, int) { throw_nogpu(); }
  64 void cv::gpu::GeneralizedHough_GPU::detect(const GpuMat&, const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); }
  65 void cv::gpu::GeneralizedHough_GPU::download(const GpuMat&, OutputArray, OutputArray) { throw_nogpu(); }
  66 void cv::gpu::GeneralizedHough_GPU::release() {}
  67
  68 #else /* !defined (HAVE_CUDA) */
  69
  70 namespace cv { namespace gpu { namespace device
  71 {
  72     namespace hough
  73     {
  74         int buildPointList_gpu(PtrStepSzb src, unsigned int* list);
  75     }
  76 }}}
  77
  78 //////////////////////////////////////////////////////////
  79 // HoughLines
  80
  81 namespace cv { namespace gpu { namespace device
  82 {
  83     namespace hough
  84     {
  85         void linesAccum_gpu(const unsigned int* list, int count, PtrStepSzi accum, float rho, float theta, size_t sharedMemPerBlock, bool has20);
  86         int linesGetResult_gpu(PtrStepSzi accum, float2* out, int* votes, int maxSize, float rho, float theta, int threshold, bool doSort);
  87     }
  88 }}}
  89
  90 void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, float rho, float theta, int threshold, bool doSort, int maxLines)
  91 {
  92     HoughLinesBuf buf;
  93     HoughLines(src, lines, buf, rho, theta, threshold, doSort, maxLines);
  94 }
  95
  96 void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, HoughLinesBuf& buf, float rho, float theta, int threshold, bool doSort, int maxLines)
  97 {
  98     using namespace cv::gpu::device::hough;
  99
 100     CV_Assert(src.type() == CV_8UC1);
 101     CV_Assert(src.cols < std::numeric_limits<unsigned short>::max());
 102     CV_Assert(src.rows < std::numeric_limits<unsigned short>::max());
 103
 104     ensureSizeIsEnough(1, src.size().area(), CV_32SC1, buf.list);
 105     unsigned int* srcPoints = buf.list.ptr<unsigned int>();
 106
 107     const int pointsCount = buildPointList_gpu(src, srcPoints);
 108     if (pointsCount == 0)
 109     {
 110         lines.release();
 111         return;
 112     }
 113
 114     const int numangle = cvRound(CV_PI / theta);
 115     const int numrho = cvRound(((src.cols + src.rows) * 2 + 1) / rho);
 116     CV_Assert(numangle > 0 && numrho > 0);
 117
 118     ensureSizeIsEnough(numangle + 2, numrho + 2, CV_32SC1, buf.accum);
 119     buf.accum.setTo(Scalar::all(0));
 120
 121     DeviceInfo devInfo;
 122     linesAccum_gpu(srcPoints, pointsCount, buf.accum, rho, theta, devInfo.sharedMemPerBlock(), devInfo.supports(FEATURE_SET_COMPUTE_20));
 123
 124     ensureSizeIsEnough(2, maxLines, CV_32FC2, lines);
 125
 126     int linesCount = linesGetResult_gpu(buf.accum, lines.ptr<float2>(0), lines.ptr<int>(1), maxLines, rho, theta, threshold, doSort);
 127     if (linesCount > 0)
 128         lines.cols = linesCount;
 129     else
 130         lines.release();
 131 }
 132
 133 void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, OutputArray h_votes_)
 134 {
 135     if (d_lines.empty())
 136     {
 137         h_lines_.release();
 138         if (h_votes_.needed())
 139             h_votes_.release();
 140         return;
 141     }
 142
 143     CV_Assert(d_lines.rows == 2 && d_lines.type() == CV_32FC2);
 144
 145     h_lines_.create(1, d_lines.cols, CV_32FC2);
 146     Mat h_lines = h_lines_.getMat();
 147     d_lines.row(0).download(h_lines);
 148
 149     if (h_votes_.needed())
 150     {
 151         h_votes_.create(1, d_lines.cols, CV_32SC1);
 152         Mat h_votes = h_votes_.getMat();
 153         GpuMat d_votes(1, d_lines.cols, CV_32SC1, const_cast<int*>(d_lines.ptr<int>(1)));
 154         d_votes.download(h_votes);
 155     }
 156 }
 157
 158 //////////////////////////////////////////////////////////
 159 // HoughCircles
 160
 161 namespace cv { namespace gpu { namespace device
 162 {
 163     namespace hough
 164     {
 165         void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp);
 166         int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold);
 167         int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count,
 168                                    float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20);
 169     }
 170 }}}
 171
 172 void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, int method, float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles)
 173 {
 174     HoughCirclesBuf buf;
 175     HoughCircles(src, circles, buf, method, dp, minDist, cannyThreshold, votesThreshold, minRadius, maxRadius, maxCircles);
 176 }
 177
 178 void cv::gpu::HoughCircles(const GpuMat& src, GpuMat& circles, HoughCirclesBuf& buf, int method,
 179                            float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles)
 180 {
 181     using namespace cv::gpu::device::hough;
 182
 183     CV_Assert(src.type() == CV_8UC1);
 184     CV_Assert(src.cols < std::numeric_limits<unsigned short>::max());
 185     CV_Assert(src.rows < std::numeric_limits<unsigned short>::max());
 186     CV_Assert(method == CV_HOUGH_GRADIENT);
 187     CV_Assert(dp > 0);
 188     CV_Assert(minRadius > 0 && maxRadius > minRadius);
 189     CV_Assert(cannyThreshold > 0);
 190     CV_Assert(votesThreshold > 0);
 191     CV_Assert(maxCircles > 0);
 192
 193     const float idp = 1.0f / dp;
 194
 195     cv::gpu::Canny(src, buf.cannyBuf, buf.edges, std::max(cannyThreshold / 2, 1), cannyThreshold);
 196
 197     ensureSizeIsEnough(2, src.size().area(), CV_32SC1, buf.list);
 198     unsigned int* srcPoints = buf.list.ptr<unsigned int>(0);
 199     unsigned int* centers = buf.list.ptr<unsigned int>(1);
 200
 201     const int pointsCount = buildPointList_gpu(buf.edges, srcPoints);
 202     if (pointsCount == 0)
 203     {
 204         circles.release();
 205         return;
 206     }
 207
 208     ensureSizeIsEnough(cvCeil(src.rows * idp) + 2, cvCeil(src.cols * idp) + 2, CV_32SC1, buf.accum);
 209     buf.accum.setTo(Scalar::all(0));
 210
 211     circlesAccumCenters_gpu(srcPoints, pointsCount, buf.cannyBuf.dx, buf.cannyBuf.dy, buf.accum, minRadius, maxRadius, idp);
 212
 213     int centersCount = buildCentersList_gpu(buf.accum, centers, votesThreshold);
 214     if (centersCount == 0)
 215     {
 216         circles.release();
 217         return;
 218     }
 219
 220     if (minDist > 1)
 221     {
 222         cv::AutoBuffer<ushort2> oldBuf_(centersCount);
 223         cv::AutoBuffer<ushort2> newBuf_(centersCount);
 224         int newCount = 0;
 225
 226         ushort2* oldBuf = oldBuf_;
 227         ushort2* newBuf = newBuf_;
 228
 229         cudaSafeCall( cudaMemcpy(oldBuf, centers, centersCount * sizeof(ushort2), cudaMemcpyDeviceToHost) );
 230
 231         const int cellSize = cvRound(minDist);
 232         const int gridWidth = (src.cols + cellSize - 1) / cellSize;
 233         const int gridHeight = (src.rows + cellSize - 1) / cellSize;
 234
 235         std::vector< std::vector<ushort2> > grid(gridWidth * gridHeight);
 236
 237         const float minDist2 = minDist * minDist;
 238
 239         for (int i = 0; i < centersCount; ++i)
 240         {
 241             ushort2 p = oldBuf[i];
 242
 243             bool good = true;
 244
 245             int xCell = static_cast<int>(p.x / cellSize);
 246             int yCell = static_cast<int>(p.y / cellSize);
 247
 248             int x1 = xCell - 1;
 249             int y1 = yCell - 1;
 250             int x2 = xCell + 1;
 251             int y2 = yCell + 1;
 252
 253             // boundary check
 254             x1 = std::max(0, x1);
 255             y1 = std::max(0, y1);
 256             x2 = std::min(gridWidth - 1, x2);
 257             y2 = std::min(gridHeight - 1, y2);
 258
 259             for (int yy = y1; yy <= y2; ++yy)
 260             {
 261                 for (int xx = x1; xx <= x2; ++xx)
 262                 {
 263                     vector<ushort2>& m = grid[yy * gridWidth + xx];
 264
 265                     for(size_t j = 0; j < m.size(); ++j)
 266                     {
 267                         float dx = (float)(p.x - m[j].x);
 268                         float dy = (float)(p.y - m[j].y);
 269
 270                         if (dx * dx + dy * dy < minDist2)
 271                         {
 272                             good = false;
 273                             goto break_out;
 274                         }
 275                     }
 276                 }
 277             }
 278
 279             break_out:
 280
 281             if(good)
 282             {
 283                 grid[yCell * gridWidth + xCell].push_back(p);
 284
 285                 newBuf[newCount++] = p;
 286             }
 287         }
 288
 289         cudaSafeCall( cudaMemcpy(centers, newBuf, newCount * sizeof(unsigned int), cudaMemcpyHostToDevice) );
 290         centersCount = newCount;
 291     }
 292
 293     ensureSizeIsEnough(1, maxCircles, CV_32FC3, circles);
 294
 295     DeviceInfo devInfo;
 296     const int circlesCount = circlesAccumRadius_gpu(centers, centersCount, srcPoints, pointsCount, circles.ptr<float3>(), maxCircles,
 297                                                     dp, minRadius, maxRadius, votesThreshold, devInfo.supports(FEATURE_SET_COMPUTE_20));
 298
 299     if (circlesCount > 0)
 300         circles.cols = circlesCount;
 301     else
 302         circles.release();
 303 }
 304
 305 void cv::gpu::HoughCirclesDownload(const GpuMat& d_circles, cv::OutputArray h_circles_)
 306 {
 307     if (d_circles.empty())
 308     {
 309         h_circles_.release();
 310         return;
 311     }
 312
 313     CV_Assert(d_circles.rows == 1 && d_circles.type() == CV_32FC3);
 314
 315     h_circles_.create(1, d_circles.cols, CV_32FC3);
 316     Mat h_circles = h_circles_.getMat();
 317     d_circles.download(h_circles);
 318 }
 319
 320 //////////////////////////////////////////////////////////
 321 // GeneralizedHough
 322
 323 namespace cv { namespace gpu { namespace device
 324 {
 325     namespace hough
 326     {
 327         template <typename T>
 328         int buildEdgePointList_gpu(PtrStepSzb edges, PtrStepSzb dx, PtrStepSzb dy, unsigned int* coordList, float* thetaList);
 329         void buildRTable_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
 330                              PtrStepSz<short2> r_table, int* r_sizes,
 331                              short2 templCenter, int levels);
 332
 333         void GHT_Ballard_Pos_calcHist_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
 334                                           PtrStepSz<short2> r_table, const int* r_sizes,
 335                                           PtrStepSzi hist,
 336                                           float dp, int levels);
 337         int GHT_Ballard_Pos_findPosInHist_gpu(PtrStepSzi hist, float4* out, int3* votes, int maxSize, float dp, int threshold);
 338
 339         void GHT_Ballard_PosScale_calcHist_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
 340                                                PtrStepSz<short2> r_table, const int* r_sizes,
 341                                                PtrStepi hist, int rows, int cols,
 342                                                float minScale, float scaleStep, int scaleRange,
 343                                                float dp, int levels);
 344         int GHT_Ballard_PosScale_findPosInHist_gpu(PtrStepi hist, int rows, int cols, int scaleRange, float4* out, int3* votes, int maxSize,
 345                                                    float minScale, float scaleStep, float dp, int threshold);
 346
 347         void GHT_Ballard_PosRotation_calcHist_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
 348                                                   PtrStepSz<short2> r_table, const int* r_sizes,
 349                                                   PtrStepi hist, int rows, int cols,
 350                                                   float minAngle, float angleStep, int angleRange,
 351                                                   float dp, int levels);
 352         int GHT_Ballard_PosRotation_findPosInHist_gpu(PtrStepi hist, int rows, int cols, int angleRange, float4* out, int3* votes, int maxSize,
 353                                                       float minAngle, float angleStep, float dp, int threshold);
 354
 355         void GHT_Guil_Full_setTemplFeatures(PtrStepb p1_pos, PtrStepb p1_theta, PtrStepb p2_pos, PtrStepb d12, PtrStepb r1, PtrStepb r2);
 356         void GHT_Guil_Full_setImageFeatures(PtrStepb p1_pos, PtrStepb p1_theta, PtrStepb p2_pos, PtrStepb d12, PtrStepb r1, PtrStepb r2);
 357         void GHT_Guil_Full_buildTemplFeatureList_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
 358                                                      int* sizes, int maxSize,
 359                                                      float xi, float angleEpsilon, int levels,
 360                                                      float2 center, float maxDist);
 361         void GHT_Guil_Full_buildImageFeatureList_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
 362                                                      int* sizes, int maxSize,
 363                                                      float xi, float angleEpsilon, int levels,
 364                                                      float2 center, float maxDist);
 365         void GHT_Guil_Full_calcOHist_gpu(const int* templSizes, const int* imageSizes, int* OHist,
 366                                          float minAngle, float maxAngle, float angleStep, int angleRange,
 367                                          int levels, int tMaxSize);
 368         void GHT_Guil_Full_calcSHist_gpu(const int* templSizes, const int* imageSizes, int* SHist,
 369                                          float angle, float angleEpsilon,
 370                                          float minScale, float maxScale, float iScaleStep, int scaleRange,
 371                                          int levels, int tMaxSize);
 372         void GHT_Guil_Full_calcPHist_gpu(const int* templSizes, const int* imageSizes, PtrStepSzi PHist,
 373                                          float angle, float angleEpsilon, float scale,
 374                                          float dp,
 375                                          int levels, int tMaxSize);
 376         int GHT_Guil_Full_findPosInHist_gpu(PtrStepSzi hist, float4* out, int3* votes, int curSize, int maxSize,
 377                                              float angle, int angleVotes, float scale, int scaleVotes,
 378                                              float dp, int threshold);
 379     }
 380 }}}
 381
 382 namespace
 383 {
 384     /////////////////////////////////////
 385     // Common
 386
 387     template <typename T, class A> void releaseVector(vector<T, A>& v)
 388     {
 389         vector<T, A> empty;
 390         empty.swap(v);
 391     }
 392
 393     class GHT_Pos : public GeneralizedHough_GPU
 394     {
 395     public:
 396         GHT_Pos();
 397
 398     protected:
 399         void setTemplateImpl(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, Point templCenter);
 400         void detectImpl(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, GpuMat& positions);
 401         void releaseImpl();
 402
 403         virtual void processTempl() = 0;
 404         virtual void processImage() = 0;
 405
 406         void buildEdgePointList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy);
 407         void filterMinDist();
 408         void convertTo(GpuMat& positions);
 409
 410         int maxSize;
 411         double minDist;
 412
 413         Size templSize;
 414         Point templCenter;
 415         GpuMat templEdges;
 416         GpuMat templDx;
 417         GpuMat templDy;
 418
 419         Size imageSize;
 420         GpuMat imageEdges;
 421         GpuMat imageDx;
 422         GpuMat imageDy;
 423
 424         GpuMat edgePointList;
 425
 426         GpuMat outBuf;
 427         int posCount;
 428
 429         vector<float4> oldPosBuf;
 430         vector<int3> oldVoteBuf;
 431         vector<float4> newPosBuf;
 432         vector<int3> newVoteBuf;
 433         vector<int> indexies;
 434     };
 435
 436     GHT_Pos::GHT_Pos()
 437     {
 438         maxSize = 10000;
 439         minDist = 1.0;
 440     }
 441
 442     void GHT_Pos::setTemplateImpl(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, Point templCenter_)
 443     {
 444         templSize = edges.size();
 445         templCenter = templCenter_;
 446
 447         ensureSizeIsEnough(templSize, edges.type(), templEdges);
 448         ensureSizeIsEnough(templSize, dx.type(), templDx);
 449         ensureSizeIsEnough(templSize, dy.type(), templDy);
 450
 451         edges.copyTo(templEdges);
 452         dx.copyTo(templDx);
 453         dy.copyTo(templDy);
 454
 455         processTempl();
 456     }
 457
 458     void GHT_Pos::detectImpl(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, GpuMat& positions)
 459     {
 460         imageSize = edges.size();
 461
 462         ensureSizeIsEnough(imageSize, edges.type(), imageEdges);
 463         ensureSizeIsEnough(imageSize, dx.type(), imageDx);
 464         ensureSizeIsEnough(imageSize, dy.type(), imageDy);
 465
 466         edges.copyTo(imageEdges);
 467         dx.copyTo(imageDx);
 468         dy.copyTo(imageDy);
 469
 470         posCount = 0;
 471
 472         processImage();
 473
 474         if (posCount == 0)
 475             positions.release();
 476         else
 477         {
 478             if (minDist > 1)
 479                 filterMinDist();
 480             convertTo(positions);
 481         }
 482     }
 483
 484     void GHT_Pos::releaseImpl()
 485     {
 486         templSize = Size();
 487         templCenter = Point(-1, -1);
 488         templEdges.release();
 489         templDx.release();
 490         templDy.release();
 491
 492         imageSize = Size();
 493         imageEdges.release();
 494         imageDx.release();
 495         imageDy.release();
 496
 497         edgePointList.release();
 498
 499         outBuf.release();
 500         posCount = 0;
 501
 502         releaseVector(oldPosBuf);
 503         releaseVector(oldVoteBuf);
 504         releaseVector(newPosBuf);
 505         releaseVector(newVoteBuf);
 506         releaseVector(indexies);
 507     }
 508
 509     void GHT_Pos::buildEdgePointList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy)
 510     {
 511         using namespace cv::gpu::device::hough;
 512
 513         typedef int (*func_t)(PtrStepSzb edges, PtrStepSzb dx, PtrStepSzb dy, unsigned int* coordList, float* thetaList);
 514         static const func_t funcs[] =
 515         {
 516             0,
 517             0,
 518             0,
 519             buildEdgePointList_gpu<short>,
 520             buildEdgePointList_gpu<int>,
 521             buildEdgePointList_gpu<float>,
 522             0
 523         };
 524
 525         CV_Assert(edges.type() == CV_8UC1);
 526         CV_Assert(dx.size() == edges.size());
 527         CV_Assert(dy.type() == dx.type() && dy.size() == edges.size());
 528
 529         const func_t func = funcs[dx.depth()];
 530         CV_Assert(func != 0);
 531
 532         edgePointList.cols = edgePointList.step / sizeof(int);
 533         ensureSizeIsEnough(2, edges.size().area(), CV_32SC1, edgePointList);
 534
 535         edgePointList.cols = func(edges, dx, dy, edgePointList.ptr<unsigned int>(0), edgePointList.ptr<float>(1));
 536     }
 537
 538     #define votes_cmp_gt(l1, l2) (aux[l1].x > aux[l2].x)
 539     static CV_IMPLEMENT_QSORT_EX( sortIndexies, int, votes_cmp_gt, const int3* )
 540
 541     void GHT_Pos::filterMinDist()
 542     {
 543         oldPosBuf.resize(posCount);
 544         oldVoteBuf.resize(posCount);
 545
 546         cudaSafeCall( cudaMemcpy(&oldPosBuf[0], outBuf.ptr(0), posCount * sizeof(float4), cudaMemcpyDeviceToHost) );
 547         cudaSafeCall( cudaMemcpy(&oldVoteBuf[0], outBuf.ptr(1), posCount * sizeof(int3), cudaMemcpyDeviceToHost) );
 548
 549         indexies.resize(posCount);
 550         for (int i = 0; i < posCount; ++i)
 551             indexies[i] = i;
 552         sortIndexies(&indexies[0], posCount, &oldVoteBuf[0]);
 553
 554         newPosBuf.clear();
 555         newVoteBuf.clear();
 556         newPosBuf.reserve(posCount);
 557         newVoteBuf.reserve(posCount);
 558
 559         const int cellSize = cvRound(minDist);
 560         const int gridWidth = (imageSize.width + cellSize - 1) / cellSize;
 561         const int gridHeight = (imageSize.height + cellSize - 1) / cellSize;
 562
 563         vector< vector<Point2f> > grid(gridWidth * gridHeight);
 564
 565         const double minDist2 = minDist * minDist;
 566
 567         for (int i = 0; i < posCount; ++i)
 568         {
 569             const int ind = indexies[i];
 570
 571             Point2f p(oldPosBuf[ind].x, oldPosBuf[ind].y);
 572
 573             bool good = true;
 574
 575             const int xCell = static_cast<int>(p.x / cellSize);
 576             const int yCell = static_cast<int>(p.y / cellSize);
 577
 578             int x1 = xCell - 1;
 579             int y1 = yCell - 1;
 580             int x2 = xCell + 1;
 581             int y2 = yCell + 1;
 582
 583             // boundary check
 584             x1 = std::max(0, x1);
 585             y1 = std::max(0, y1);
 586             x2 = std::min(gridWidth - 1, x2);
 587             y2 = std::min(gridHeight - 1, y2);
 588
 589             for (int yy = y1; yy <= y2; ++yy)
 590             {
 591                 for (int xx = x1; xx <= x2; ++xx)
 592                 {
 593                     const vector<Point2f>& m = grid[yy * gridWidth + xx];
 594
 595                     for(size_t j = 0; j < m.size(); ++j)
 596                     {
 597                         const Point2f d = p - m[j];
 598
 599                         if (d.ddot(d) < minDist2)
 600                         {
 601                             good = false;
 602                             goto break_out;
 603                         }
 604                     }
 605                 }
 606             }
 607
 608             break_out:
 609
 610             if(good)
 611             {
 612                 grid[yCell * gridWidth + xCell].push_back(p);
 613
 614                 newPosBuf.push_back(oldPosBuf[ind]);
 615                 newVoteBuf.push_back(oldVoteBuf[ind]);
 616             }
 617         }
 618
 619         posCount = static_cast<int>(newPosBuf.size());
 620         cudaSafeCall( cudaMemcpy(outBuf.ptr(0), &newPosBuf[0], posCount * sizeof(float4), cudaMemcpyHostToDevice) );
 621         cudaSafeCall( cudaMemcpy(outBuf.ptr(1), &newVoteBuf[0], posCount * sizeof(int3), cudaMemcpyHostToDevice) );
 622     }
 623
 624     void GHT_Pos::convertTo(GpuMat& positions)
 625     {
 626         ensureSizeIsEnough(2, posCount, CV_32FC4, positions);
 627         GpuMat(2, posCount, CV_32FC4, outBuf.data, outBuf.step).copyTo(positions);
 628     }
 629
 630     /////////////////////////////////////
 631     // POSITION Ballard
 632
 633     class GHT_Ballard_Pos : public GHT_Pos
 634     {
 635     public:
 636         AlgorithmInfo* info() const;
 637
 638         GHT_Ballard_Pos();
 639
 640     protected:
 641         void releaseImpl();
 642
 643         void processTempl();
 644         void processImage();
 645
 646         virtual void calcHist();
 647         virtual void findPosInHist();
 648
 649         int levels;
 650         int votesThreshold;
 651         double dp;
 652
 653         GpuMat r_table;
 654         GpuMat r_sizes;
 655
 656         GpuMat hist;
 657     };
 658
 659     CV_INIT_ALGORITHM(GHT_Ballard_Pos, "GeneralizedHough_GPU.POSITION",
 660                       obj.info()->addParam(obj, "maxSize", obj.maxSize, false, 0, 0,
 661                                            "Maximal size of inner buffers.");
 662                       obj.info()->addParam(obj, "minDist", obj.minDist, false, 0, 0,
 663                                            "Minimum distance between the centers of the detected objects.");
 664                       obj.info()->addParam(obj, "levels", obj.levels, false, 0, 0,
 665                                            "R-Table levels.");
 666                       obj.info()->addParam(obj, "votesThreshold", obj.votesThreshold, false, 0, 0,
 667                                            "The accumulator threshold for the template centers at the detection stage. The smaller it is, the more false positions may be detected.");
 668                       obj.info()->addParam(obj, "dp", obj.dp, false, 0, 0,
 669                                            "Inverse ratio of the accumulator resolution to the image resolution."));
 670
 671     GHT_Ballard_Pos::GHT_Ballard_Pos()
 672     {
 673         levels = 360;
 674         votesThreshold = 100;
 675         dp = 1.0;
 676     }
 677
 678     void GHT_Ballard_Pos::releaseImpl()
 679     {
 680         GHT_Pos::releaseImpl();
 681
 682         r_table.release();
 683         r_sizes.release();
 684
 685         hist.release();
 686     }
 687
 688     void GHT_Ballard_Pos::processTempl()
 689     {
 690         using namespace cv::gpu::device::hough;
 691
 692         CV_Assert(levels > 0);
 693
 694         buildEdgePointList(templEdges, templDx, templDy);
 695
 696         ensureSizeIsEnough(levels + 1, maxSize, CV_16SC2, r_table);
 697         ensureSizeIsEnough(1, levels + 1, CV_32SC1, r_sizes);
 698         r_sizes.setTo(Scalar::all(0));
 699
 700         if (edgePointList.cols > 0)
 701         {
 702             buildRTable_gpu(edgePointList.ptr<unsigned int>(0), edgePointList.ptr<float>(1), edgePointList.cols,
 703                             r_table, r_sizes.ptr<int>(), make_short2(templCenter.x, templCenter.y), levels);
 704             min(r_sizes, maxSize, r_sizes);
 705         }
 706     }
 707
 708     void GHT_Ballard_Pos::processImage()
 709     {
 710         calcHist();
 711         findPosInHist();
 712     }
 713
 714     void GHT_Ballard_Pos::calcHist()
 715     {
 716         using namespace cv::gpu::device::hough;
 717
 718         CV_Assert(levels > 0 && r_table.rows == (levels + 1) && r_sizes.cols == (levels + 1));
 719         CV_Assert(dp > 0.0);
 720
 721         const double idp = 1.0 / dp;
 722
 723         buildEdgePointList(imageEdges, imageDx, imageDy);
 724
 725         ensureSizeIsEnough(cvCeil(imageSize.height * idp) + 2, cvCeil(imageSize.width * idp) + 2, CV_32SC1, hist);
 726         hist.setTo(Scalar::all(0));
 727
 728         if (edgePointList.cols > 0)
 729         {
 730             GHT_Ballard_Pos_calcHist_gpu(edgePointList.ptr<unsigned int>(0), edgePointList.ptr<float>(1), edgePointList.cols,
 731                                          r_table, r_sizes.ptr<int>(),
 732                                          hist,
 733                                          dp, levels);
 734         }
 735     }
 736
 737     void GHT_Ballard_Pos::findPosInHist()
 738     {
 739         using namespace cv::gpu::device::hough;
 740
 741         CV_Assert(votesThreshold > 0);
 742
 743         ensureSizeIsEnough(2, maxSize, CV_32FC4, outBuf);
 744
 745         posCount = GHT_Ballard_Pos_findPosInHist_gpu(hist, outBuf.ptr<float4>(0), outBuf.ptr<int3>(1), maxSize, dp, votesThreshold);
 746     }
 747
 748     /////////////////////////////////////
 749     // POSITION & SCALE
 750
 751     class GHT_Ballard_PosScale : public GHT_Ballard_Pos
 752     {
 753     public:
 754         AlgorithmInfo* info() const;
 755
 756         GHT_Ballard_PosScale();
 757
 758     protected:
 759         void calcHist();
 760         void findPosInHist();
 761
 762         double minScale;
 763         double maxScale;
 764         double scaleStep;
 765     };
 766
 767     CV_INIT_ALGORITHM(GHT_Ballard_PosScale, "GeneralizedHough_GPU.POSITION_SCALE",
 768                       obj.info()->addParam(obj, "maxSize", obj.maxSize, false, 0, 0,
 769                                            "Maximal size of inner buffers.");
 770                       obj.info()->addParam(obj, "minDist", obj.minDist, false, 0, 0,
 771                                            "Minimum distance between the centers of the detected objects.");
 772                       obj.info()->addParam(obj, "levels", obj.levels, false, 0, 0,
 773                                            "R-Table levels.");
 774                       obj.info()->addParam(obj, "votesThreshold", obj.votesThreshold, false, 0, 0,
 775                                            "The accumulator threshold for the template centers at the detection stage. The smaller it is, the more false positions may be detected.");
 776                       obj.info()->addParam(obj, "dp", obj.dp, false, 0, 0,
 777                                            "Inverse ratio of the accumulator resolution to the image resolution.");
 778                       obj.info()->addParam(obj, "minScale", obj.minScale, false, 0, 0,
 779                                            "Minimal scale to detect.");
 780                       obj.info()->addParam(obj, "maxScale", obj.maxScale, false, 0, 0,
 781                                            "Maximal scale to detect.");
 782                       obj.info()->addParam(obj, "scaleStep", obj.scaleStep, false, 0, 0,
 783                                            "Scale step."));
 784
 785     GHT_Ballard_PosScale::GHT_Ballard_PosScale()
 786     {
 787         minScale = 0.5;
 788         maxScale = 2.0;
 789         scaleStep = 0.05;
 790     }
 791
 792     void GHT_Ballard_PosScale::calcHist()
 793     {
 794         using namespace cv::gpu::device::hough;
 795
 796         CV_Assert(levels > 0 && r_table.rows == (levels + 1) && r_sizes.cols == (levels + 1));
 797         CV_Assert(dp > 0.0);
 798         CV_Assert(minScale > 0.0 && minScale < maxScale);
 799         CV_Assert(scaleStep > 0.0);
 800
 801         const double idp = 1.0 / dp;
 802         const int scaleRange = cvCeil((maxScale - minScale) / scaleStep);
 803         const int rows = cvCeil(imageSize.height * idp);
 804         const int cols = cvCeil(imageSize.width * idp);
 805
 806         buildEdgePointList(imageEdges, imageDx, imageDy);
 807
 808         ensureSizeIsEnough((scaleRange + 2) * (rows + 2), cols + 2, CV_32SC1, hist);
 809         hist.setTo(Scalar::all(0));
 810
 811         if (edgePointList.cols > 0)
 812         {
 813             GHT_Ballard_PosScale_calcHist_gpu(edgePointList.ptr<unsigned int>(0), edgePointList.ptr<float>(1), edgePointList.cols,
 814                                               r_table, r_sizes.ptr<int>(),
 815                                               hist, rows, cols,
 816                                               minScale, scaleStep, scaleRange, dp, levels);
 817         }
 818     }
 819
 820     void GHT_Ballard_PosScale::findPosInHist()
 821     {
 822         using namespace cv::gpu::device::hough;
 823
 824         CV_Assert(votesThreshold > 0);
 825
 826         const double idp = 1.0 / dp;
 827         const int scaleRange = cvCeil((maxScale - minScale) / scaleStep);
 828         const int rows = cvCeil(imageSize.height * idp);
 829         const int cols = cvCeil(imageSize.width * idp);
 830
 831         ensureSizeIsEnough(2, maxSize, CV_32FC4, outBuf);
 832
 833         posCount =  GHT_Ballard_PosScale_findPosInHist_gpu(hist, rows, cols, scaleRange, outBuf.ptr<float4>(0), outBuf.ptr<int3>(1), maxSize, minScale, scaleStep, dp, votesThreshold);
 834     }
 835
 836     /////////////////////////////////////
 837     // POSITION & Rotation
 838
 839     class GHT_Ballard_PosRotation : public GHT_Ballard_Pos
 840     {
 841     public:
 842         AlgorithmInfo* info() const;
 843
 844         GHT_Ballard_PosRotation();
 845
 846     protected:
 847         void calcHist();
 848         void findPosInHist();
 849
 850         double minAngle;
 851         double maxAngle;
 852         double angleStep;
 853     };
 854
 855     CV_INIT_ALGORITHM(GHT_Ballard_PosRotation, "GeneralizedHough_GPU.POSITION_ROTATION",
 856                       obj.info()->addParam(obj, "maxSize", obj.maxSize, false, 0, 0,
 857                                            "Maximal size of inner buffers.");
 858                       obj.info()->addParam(obj, "minDist", obj.minDist, false, 0, 0,
 859                                            "Minimum distance between the centers of the detected objects.");
 860                       obj.info()->addParam(obj, "levels", obj.levels, false, 0, 0,
 861                                            "R-Table levels.");
 862                       obj.info()->addParam(obj, "votesThreshold", obj.votesThreshold, false, 0, 0,
 863                                            "The accumulator threshold for the template centers at the detection stage. The smaller it is, the more false positions may be detected.");
 864                       obj.info()->addParam(obj, "dp", obj.dp, false, 0, 0,
 865                                            "Inverse ratio of the accumulator resolution to the image resolution.");
 866                       obj.info()->addParam(obj, "minAngle", obj.minAngle, false, 0, 0,
 867                                            "Minimal rotation angle to detect in degrees.");
 868                       obj.info()->addParam(obj, "maxAngle", obj.maxAngle, false, 0, 0,
 869                                            "Maximal rotation angle to detect in degrees.");
 870                       obj.info()->addParam(obj, "angleStep", obj.angleStep, false, 0, 0,
 871                                            "Angle step in degrees."));
 872
 873     GHT_Ballard_PosRotation::GHT_Ballard_PosRotation()
 874     {
 875         minAngle = 0.0;
 876         maxAngle = 360.0;
 877         angleStep = 1.0;
 878     }
 879
 880     void GHT_Ballard_PosRotation::calcHist()
 881     {
 882         using namespace cv::gpu::device::hough;
 883
 884         CV_Assert(levels > 0 && r_table.rows == (levels + 1) && r_sizes.cols == (levels + 1));
 885         CV_Assert(dp > 0.0);
 886         CV_Assert(minAngle >= 0.0 && minAngle < maxAngle && maxAngle <= 360.0);
 887         CV_Assert(angleStep > 0.0 && angleStep < 360.0);
 888
 889         const double idp = 1.0 / dp;
 890         const int angleRange = cvCeil((maxAngle - minAngle) / angleStep);
 891         const int rows = cvCeil(imageSize.height * idp);
 892         const int cols = cvCeil(imageSize.width * idp);
 893
 894         buildEdgePointList(imageEdges, imageDx, imageDy);
 895
 896         ensureSizeIsEnough((angleRange + 2) * (rows + 2), cols + 2, CV_32SC1, hist);
 897         hist.setTo(Scalar::all(0));
 898
 899         if (edgePointList.cols > 0)
 900         {
 901             GHT_Ballard_PosRotation_calcHist_gpu(edgePointList.ptr<unsigned int>(0), edgePointList.ptr<float>(1), edgePointList.cols,
 902                                                  r_table, r_sizes.ptr<int>(),
 903                                                  hist, rows, cols,
 904                                                  minAngle, angleStep, angleRange, dp, levels);
 905         }
 906     }
 907
 908     void GHT_Ballard_PosRotation::findPosInHist()
 909     {
 910         using namespace cv::gpu::device::hough;
 911
 912         CV_Assert(votesThreshold > 0);
 913
 914         const double idp = 1.0 / dp;
 915         const int angleRange = cvCeil((maxAngle - minAngle) / angleStep);
 916         const int rows = cvCeil(imageSize.height * idp);
 917         const int cols = cvCeil(imageSize.width * idp);
 918
 919         ensureSizeIsEnough(2, maxSize, CV_32FC4, outBuf);
 920
 921         posCount = GHT_Ballard_PosRotation_findPosInHist_gpu(hist, rows, cols, angleRange, outBuf.ptr<float4>(0), outBuf.ptr<int3>(1), maxSize, minAngle, angleStep, dp, votesThreshold);
 922     }
 923
 924     /////////////////////////////////////////
 925     // POSITION & SCALE & ROTATION
 926
 927     double toRad(double a)
 928     {
 929         return a * CV_PI / 180.0;
 930     }
 931
 932     double clampAngle(double a)
 933     {
 934         double res = a;
 935
 936         while (res > 360.0)
 937             res -= 360.0;
 938         while (res < 0)
 939             res += 360.0;
 940
 941         return res;
 942     }
 943
 944     bool angleEq(double a, double b, double eps = 1.0)
 945     {
 946         return (fabs(clampAngle(a - b)) <= eps);
 947     }
 948
 949     class GHT_Guil_Full : public GHT_Pos
 950     {
 951     public:
 952         AlgorithmInfo* info() const;
 953
 954         GHT_Guil_Full();
 955
 956     protected:
 957         void releaseImpl();
 958
 959         void processTempl();
 960         void processImage();
 961
 962         struct Feature
 963         {
 964             GpuMat p1_pos;
 965             GpuMat p1_theta;
 966             GpuMat p2_pos;
 967
 968             GpuMat d12;
 969
 970             GpuMat r1;
 971             GpuMat r2;
 972
 973             GpuMat sizes;
 974             int maxSize;
 975
 976             void create(int levels, int maxCapacity, bool isTempl);
 977             void release();
 978         };
 979
 980         typedef void (*set_func_t)(PtrStepb p1_pos, PtrStepb p1_theta, PtrStepb p2_pos, PtrStepb d12, PtrStepb r1, PtrStepb r2);
 981         typedef void (*build_func_t)(const unsigned int* coordList, const float* thetaList, int pointsCount,
 982                                      int* sizes, int maxSize,
 983                                      float xi, float angleEpsilon, int levels,
 984                                      float2 center, float maxDist);
 985
 986         void buildFeatureList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, Feature& features,
 987                               set_func_t set_func, build_func_t build_func, bool isTempl, Point2d center = Point2d());
 988
 989         void calcOrientation();
 990         void calcScale(double angle);
 991         void calcPosition(double angle, int angleVotes, double scale, int scaleVotes);
 992
 993         double xi;
 994         int levels;
 995         double angleEpsilon;
 996
 997         double minAngle;
 998         double maxAngle;
 999         double angleStep;
1000         int angleThresh;
1001
1002         double minScale;
1003         double maxScale;
1004         double scaleStep;
1005         int scaleThresh;
1006
1007         double dp;
1008         int posThresh;
1009
1010         Feature templFeatures;
1011         Feature imageFeatures;
1012
1013         vector< pair<double, int> > angles;
1014         vector< pair<double, int> > scales;
1015
1016         GpuMat hist;
1017         vector<int> h_buf;
1018     };
1019
1020     CV_INIT_ALGORITHM(GHT_Guil_Full, "GeneralizedHough_GPU.POSITION_SCALE_ROTATION",
1021                       obj.info()->addParam(obj, "minDist", obj.minDist, false, 0, 0,
1022                                            "Minimum distance between the centers of the detected objects.");
1023                       obj.info()->addParam(obj, "maxSize", obj.maxSize, false, 0, 0,
1024                                            "Maximal size of inner buffers.");
1025                       obj.info()->addParam(obj, "xi", obj.xi, false, 0, 0,
1026                                            "Angle difference in degrees between two points in feature.");
1027                       obj.info()->addParam(obj, "levels", obj.levels, false, 0, 0,
1028                                            "Feature table levels.");
1029                       obj.info()->addParam(obj, "angleEpsilon", obj.angleEpsilon, false, 0, 0,
1030                                            "Maximal difference between angles that treated as equal.");
1031                       obj.info()->addParam(obj, "minAngle", obj.minAngle, false, 0, 0,
1032                                            "Minimal rotation angle to detect in degrees.");
1033                       obj.info()->addParam(obj, "maxAngle", obj.maxAngle, false, 0, 0,
1034                                            "Maximal rotation angle to detect in degrees.");
1035                       obj.info()->addParam(obj, "angleStep", obj.angleStep, false, 0, 0,
1036                                            "Angle step in degrees.");
1037                       obj.info()->addParam(obj, "angleThresh", obj.angleThresh, false, 0, 0,
1038                                            "Angle threshold.");
1039                       obj.info()->addParam(obj, "minScale", obj.minScale, false, 0, 0,
1040                                            "Minimal scale to detect.");
1041                       obj.info()->addParam(obj, "maxScale", obj.maxScale, false, 0, 0,
1042                                            "Maximal scale to detect.");
1043                       obj.info()->addParam(obj, "scaleStep", obj.scaleStep, false, 0, 0,
1044                                            "Scale step.");
1045                       obj.info()->addParam(obj, "scaleThresh", obj.scaleThresh, false, 0, 0,
1046                                            "Scale threshold.");
1047                       obj.info()->addParam(obj, "dp", obj.dp, false, 0, 0,
1048                                            "Inverse ratio of the accumulator resolution to the image resolution.");
1049                       obj.info()->addParam(obj, "posThresh", obj.posThresh, false, 0, 0,
1050                                            "Position threshold."));
1051
1052     GHT_Guil_Full::GHT_Guil_Full()
1053     {
1054         maxSize = 1000;
1055         xi = 90.0;
1056         levels = 360;
1057         angleEpsilon = 1.0;
1058
1059         minAngle = 0.0;
1060         maxAngle = 360.0;
1061         angleStep = 1.0;
1062         angleThresh = 15000;
1063
1064         minScale = 0.5;
1065         maxScale = 2.0;
1066         scaleStep = 0.05;
1067         scaleThresh = 1000;
1068
1069         dp = 1.0;
1070         posThresh = 100;
1071     }
1072
1073     void GHT_Guil_Full::releaseImpl()
1074     {
1075         GHT_Pos::releaseImpl();
1076
1077         templFeatures.release();
1078         imageFeatures.release();
1079
1080         releaseVector(angles);
1081         releaseVector(scales);
1082
1083         hist.release();
1084         releaseVector(h_buf);
1085     }
1086
1087     void GHT_Guil_Full::processTempl()
1088     {
1089         using namespace cv::gpu::device::hough;
1090
1091         buildFeatureList(templEdges, templDx, templDy, templFeatures,
1092             GHT_Guil_Full_setTemplFeatures, GHT_Guil_Full_buildTemplFeatureList_gpu,
1093             true, templCenter);
1094
1095         h_buf.resize(templFeatures.sizes.cols);
1096         cudaSafeCall( cudaMemcpy(&h_buf[0], templFeatures.sizes.data, h_buf.size() * sizeof(int), cudaMemcpyDeviceToHost) );
1097         templFeatures.maxSize = *max_element(h_buf.begin(), h_buf.end());
1098     }
1099
1100     void GHT_Guil_Full::processImage()
1101     {
1102         using namespace cv::gpu::device::hough;
1103
1104         CV_Assert(levels > 0);
1105         CV_Assert(templFeatures.sizes.cols == levels + 1);
1106         CV_Assert(minAngle >= 0.0 && minAngle < maxAngle && maxAngle <= 360.0);
1107         CV_Assert(angleStep > 0.0 && angleStep < 360.0);
1108         CV_Assert(angleThresh > 0);
1109         CV_Assert(minScale > 0.0 && minScale < maxScale);
1110         CV_Assert(scaleStep > 0.0);
1111         CV_Assert(scaleThresh > 0);
1112         CV_Assert(dp > 0.0);
1113         CV_Assert(posThresh > 0);
1114
1115         const double iAngleStep = 1.0 / angleStep;
1116         const int angleRange = cvCeil((maxAngle - minAngle) * iAngleStep);
1117
1118         const double iScaleStep = 1.0 / scaleStep;
1119         const int scaleRange = cvCeil((maxScale - minScale) * iScaleStep);
1120
1121         const double idp = 1.0 / dp;
1122         const int histRows = cvCeil(imageSize.height * idp);
1123         const int histCols = cvCeil(imageSize.width * idp);
1124
1125         ensureSizeIsEnough(histRows + 2, std::max(angleRange + 1, std::max(scaleRange + 1, histCols + 2)), CV_32SC1, hist);
1126         h_buf.resize(std::max(angleRange + 1, scaleRange + 1));
1127
1128         ensureSizeIsEnough(2, maxSize, CV_32FC4, outBuf);
1129
1130         buildFeatureList(imageEdges, imageDx, imageDy, imageFeatures,
1131             GHT_Guil_Full_setImageFeatures, GHT_Guil_Full_buildImageFeatureList_gpu,
1132             false);
1133
1134         calcOrientation();
1135
1136         for (size_t i = 0; i < angles.size(); ++i)
1137         {
1138             const double angle = angles[i].first;
1139             const int angleVotes = angles[i].second;
1140
1141             calcScale(angle);
1142
1143             for (size_t j = 0; j < scales.size(); ++j)
1144             {
1145                 const double scale = scales[j].first;
1146                 const int scaleVotes = scales[j].second;
1147
1148                 calcPosition(angle, angleVotes, scale, scaleVotes);
1149             }
1150         }
1151     }
1152
1153     void GHT_Guil_Full::Feature::create(int levels, int maxCapacity, bool isTempl)
1154     {
1155         if (!isTempl)
1156         {
1157             ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC2, p1_pos);
1158             ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC2, p2_pos);
1159         }
1160
1161         ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC1, p1_theta);
1162
1163         ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC1, d12);
1164
1165         if (isTempl)
1166         {
1167             ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC2, r1);
1168             ensureSizeIsEnough(levels + 1, maxCapacity, CV_32FC2, r2);
1169         }
1170
1171         ensureSizeIsEnough(1, levels + 1, CV_32SC1, sizes);
1172         sizes.setTo(Scalar::all(0));
1173
1174         maxSize = 0;
1175     }
1176
1177     void GHT_Guil_Full::Feature::release()
1178     {
1179         p1_pos.release();
1180         p1_theta.release();
1181         p2_pos.release();
1182
1183         d12.release();
1184
1185         r1.release();
1186         r2.release();
1187
1188         sizes.release();
1189
1190         maxSize = 0;
1191     }
1192
1193     void GHT_Guil_Full::buildFeatureList(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, Feature& features,
1194                                          set_func_t set_func, build_func_t build_func, bool isTempl, Point2d center)
1195     {
1196         CV_Assert(levels > 0);
1197
1198         const double maxDist = sqrt((double) templSize.width * templSize.width + templSize.height * templSize.height) * maxScale;
1199
1200         features.create(levels, maxSize, isTempl);
1201         set_func(features.p1_pos, features.p1_theta, features.p2_pos, features.d12, features.r1, features.r2);
1202
1203         buildEdgePointList(edges, dx, dy);
1204
1205         if (edgePointList.cols > 0)
1206         {
1207             build_func(edgePointList.ptr<unsigned int>(0), edgePointList.ptr<float>(1), edgePointList.cols,
1208                 features.sizes.ptr<int>(), maxSize, xi, angleEpsilon, levels, make_float2(center.x, center.y), maxDist);
1209         }
1210     }
1211
1212     void GHT_Guil_Full::calcOrientation()
1213     {
1214         using namespace cv::gpu::device::hough;
1215
1216         const double iAngleStep = 1.0 / angleStep;
1217         const int angleRange = cvCeil((maxAngle - minAngle) * iAngleStep);
1218
1219         hist.setTo(Scalar::all(0));
1220         GHT_Guil_Full_calcOHist_gpu(templFeatures.sizes.ptr<int>(), imageFeatures.sizes.ptr<int>(0),
1221             hist.ptr<int>(), minAngle, maxAngle, angleStep, angleRange, levels, templFeatures.maxSize);
1222         cudaSafeCall( cudaMemcpy(&h_buf[0], hist.data, h_buf.size() * sizeof(int), cudaMemcpyDeviceToHost) );
1223
1224         angles.clear();
1225
1226         for (int n = 0; n < angleRange; ++n)
1227         {
1228             if (h_buf[n] >= angleThresh)
1229             {
1230                 const double angle = minAngle + n * angleStep;
1231                 angles.push_back(make_pair(angle, h_buf[n]));
1232             }
1233         }
1234     }
1235
1236     void GHT_Guil_Full::calcScale(double angle)
1237     {
1238         using namespace cv::gpu::device::hough;
1239
1240         const double iScaleStep = 1.0 / scaleStep;
1241         const int scaleRange = cvCeil((maxScale - minScale) * iScaleStep);
1242
1243         hist.setTo(Scalar::all(0));
1244         GHT_Guil_Full_calcSHist_gpu(templFeatures.sizes.ptr<int>(), imageFeatures.sizes.ptr<int>(0),
1245             hist.ptr<int>(), angle, angleEpsilon, minScale, maxScale, iScaleStep, scaleRange, levels, templFeatures.maxSize);
1246         cudaSafeCall( cudaMemcpy(&h_buf[0], hist.data, h_buf.size() * sizeof(int), cudaMemcpyDeviceToHost) );
1247
1248         scales.clear();
1249
1250         for (int s = 0; s < scaleRange; ++s)
1251         {
1252             if (h_buf[s] >= scaleThresh)
1253             {
1254                 const double scale = minScale + s * scaleStep;
1255                 scales.push_back(make_pair(scale, h_buf[s]));
1256             }
1257         }
1258     }
1259
1260     void GHT_Guil_Full::calcPosition(double angle, int angleVotes, double scale, int scaleVotes)
1261     {
1262         using namespace cv::gpu::device::hough;
1263
1264         hist.setTo(Scalar::all(0));
1265         GHT_Guil_Full_calcPHist_gpu(templFeatures.sizes.ptr<int>(), imageFeatures.sizes.ptr<int>(0),
1266             hist, angle, angleEpsilon, scale, dp, levels, templFeatures.maxSize);
1267
1268         posCount = GHT_Guil_Full_findPosInHist_gpu(hist, outBuf.ptr<float4>(0), outBuf.ptr<int3>(1),
1269             posCount, maxSize, angle, angleVotes, scale, scaleVotes, dp, posThresh);
1270     }
1271 }
1272
1273 Ptr<GeneralizedHough_GPU> cv::gpu::GeneralizedHough_GPU::create(int method)
1274 {
1275     switch (method)
1276     {
1277     case GHT_POSITION:
1278         CV_Assert( !GHT_Ballard_Pos_info_auto.name().empty() );
1279         return new GHT_Ballard_Pos();
1280
1281     case (GHT_POSITION | GHT_SCALE):
1282         CV_Assert( !GHT_Ballard_PosScale_info_auto.name().empty() );
1283         return new GHT_Ballard_PosScale();
1284
1285     case (GHT_POSITION | GHT_ROTATION):
1286         CV_Assert( !GHT_Ballard_PosRotation_info_auto.name().empty() );
1287         return new GHT_Ballard_PosRotation();
1288
1289     case (GHT_POSITION | GHT_SCALE | GHT_ROTATION):
1290         CV_Assert( !GHT_Guil_Full_info_auto.name().empty() );
1291         return new GHT_Guil_Full();
1292     }
1293
1294     CV_Error(CV_StsBadArg, "Unsupported method");
1295     return Ptr<GeneralizedHough_GPU>();
1296 }
1297
1298 cv::gpu::GeneralizedHough_GPU::~GeneralizedHough_GPU()
1299 {
1300 }
1301
1302 void cv::gpu::GeneralizedHough_GPU::setTemplate(const GpuMat& templ, int cannyThreshold, Point templCenter)
1303 {
1304     CV_Assert(templ.type() == CV_8UC1);
1305     CV_Assert(cannyThreshold > 0);
1306
1307     ensureSizeIsEnough(templ.size(), CV_8UC1, edges_);
1308     Canny(templ, cannyBuf_, edges_, cannyThreshold / 2, cannyThreshold);
1309
1310     if (templCenter == Point(-1, -1))
1311         templCenter = Point(templ.cols / 2, templ.rows / 2);
1312
1313     setTemplateImpl(edges_, cannyBuf_.dx, cannyBuf_.dy, templCenter);
1314 }
1315
1316 void cv::gpu::GeneralizedHough_GPU::setTemplate(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, Point templCenter)
1317 {
1318     if (templCenter == Point(-1, -1))
1319         templCenter = Point(edges.cols / 2, edges.rows / 2);
1320
1321     setTemplateImpl(edges, dx, dy, templCenter);
1322 }
1323
1324 void cv::gpu::GeneralizedHough_GPU::detect(const GpuMat& image, GpuMat& positions, int cannyThreshold)
1325 {
1326     CV_Assert(image.type() == CV_8UC1);
1327     CV_Assert(cannyThreshold > 0);
1328
1329     ensureSizeIsEnough(image.size(), CV_8UC1, edges_);
1330     Canny(image, cannyBuf_, edges_, cannyThreshold / 2, cannyThreshold);
1331
1332     detectImpl(edges_, cannyBuf_.dx, cannyBuf_.dy, positions);
1333 }
1334
1335 void cv::gpu::GeneralizedHough_GPU::detect(const GpuMat& edges, const GpuMat& dx, const GpuMat& dy, GpuMat& positions)
1336 {
1337     detectImpl(edges, dx, dy, positions);
1338 }
1339
1340 void cv::gpu::GeneralizedHough_GPU::download(const GpuMat& d_positions, OutputArray h_positions_, OutputArray h_votes_)
1341 {
1342     if (d_positions.empty())
1343     {
1344         h_positions_.release();
1345         if (h_votes_.needed())
1346             h_votes_.release();
1347         return;
1348     }
1349
1350     CV_Assert(d_positions.rows == 2 && d_positions.type() == CV_32FC4);
1351
1352     h_positions_.create(1, d_positions.cols, CV_32FC4);
1353     Mat h_positions = h_positions_.getMat();
1354     d_positions.row(0).download(h_positions);
1355
1356     if (h_votes_.needed())
1357     {
1358         h_votes_.create(1, d_positions.cols, CV_32SC3);
1359         Mat h_votes = h_votes_.getMat();
1360         GpuMat d_votes(1, d_positions.cols, CV_32SC3, const_cast<int3*>(d_positions.ptr<int3>(1)));
1361         d_votes.download(h_votes);
1362     }
1363 }
1364
1365 void cv::gpu::GeneralizedHough_GPU::release()
1366 {
1367     edges_.release();
1368     cannyBuf_.release();
1369     releaseImpl();
1370 }
1371
1372 #endif /* !defined (HAVE_CUDA) */