Skip to content

Instantly share code, notes, and snippets.

@trdenton
Created September 26, 2018 04:05
Show Gist options
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
  • Save trdenton/3365f63ae7e523dbbcdbf26e8d6d3925 to your computer and use it in GitHub Desktop.
Save trdenton/3365f63ae7e523dbbcdbf26e8d6d3925 to your computer and use it in GitHub Desktop.
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
using namespace cv;
using namespace cv::cuda;
#if !defined (HAVE_CUDA) || defined (CUDA_DISABLER) || !defined(HAVE_OPENCV_CUDAFILTERS)
Ptr<cuda::HoughCirclesDetector> cv::cuda::createHoughCirclesDetector(float, float, int, int, int, int, int) { throw_no_cuda(); return Ptr<HoughCirclesDetector>(); }
#else /* !defined (HAVE_CUDA) */
namespace cv { namespace cuda { namespace device
{
namespace hough
{
int buildPointList_gpu(PtrStepSzb src, unsigned int* list);
}
namespace hough_circles
{
void circlesAccumCenters_gpu(const unsigned int* list, int count, PtrStepi dx, PtrStepi dy, PtrStepSzi accum, int minRadius, int maxRadius, float idp);
int buildCentersList_gpu(PtrStepSzi accum, unsigned int* centers, int threshold);
int circlesAccumRadius_gpu(const unsigned int* centers, int centersCount, const unsigned int* list, int count,
float3* circles, int maxCircles, float dp, int minRadius, int maxRadius, int threshold, bool has20);
}
}}}
namespace
{
class HoughCirclesDetectorImpl : public HoughCirclesDetector
{
public:
HoughCirclesDetectorImpl(float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles);
void detect(InputArray src, OutputArray circles, Stream& stream);
void setDp(float dp) { dp_ = dp; }
float getDp() const { return dp_; }
void setMinDist(float minDist) { minDist_ = minDist; }
float getMinDist() const { return minDist_; }
void setCannyThreshold(int cannyThreshold) { cannyThreshold_ = cannyThreshold; }
int getCannyThreshold() const { return cannyThreshold_; }
void setVotesThreshold(int votesThreshold) { votesThreshold_ = votesThreshold; }
int getVotesThreshold() const { return votesThreshold_; }
void setMinRadius(int minRadius) { minRadius_ = minRadius; }
int getMinRadius() const { return minRadius_; }
void setMaxRadius(int maxRadius) { maxRadius_ = maxRadius; }
int getMaxRadius() const { return maxRadius_; }
void setMaxCircles(int maxCircles) { maxCircles_ = maxCircles; }
int getMaxCircles() const { return maxCircles_; }
void write(FileStorage& fs) const
{
writeFormat(fs);
fs << "name" << "HoughCirclesDetector_CUDA"
<< "dp" << dp_
<< "minDist" << minDist_
<< "cannyThreshold" << cannyThreshold_
<< "votesThreshold" << votesThreshold_
<< "minRadius" << minRadius_
<< "maxRadius" << maxRadius_
<< "maxCircles" << maxCircles_;
}
void read(const FileNode& fn)
{
CV_Assert( String(fn["name"]) == "HoughCirclesDetector_CUDA" );
dp_ = (float)fn["dp"];
minDist_ = (float)fn["minDist"];
cannyThreshold_ = (int)fn["cannyThreshold"];
votesThreshold_ = (int)fn["votesThreshold"];
minRadius_ = (int)fn["minRadius"];
maxRadius_ = (int)fn["maxRadius"];
maxCircles_ = (int)fn["maxCircles"];
}
private:
float dp_;
float minDist_;
int cannyThreshold_;
int votesThreshold_;
int minRadius_;
int maxRadius_;
int maxCircles_;
GpuMat dx_, dy_;
GpuMat edges_;
GpuMat accum_;
Mat tt; //CPU copy of accum_
GpuMat list_;
GpuMat result_;
Ptr<cuda::Filter> filterDx_;
Ptr<cuda::Filter> filterDy_;
Ptr<cuda::CannyEdgeDetector> canny_;
};
bool centersCompare(Vec3f a, Vec3f b) {return (a[2] > b[2]);}
HoughCirclesDetectorImpl::HoughCirclesDetectorImpl(float dp, float minDist, int cannyThreshold, int votesThreshold,
int minRadius, int maxRadius, int maxCircles) :
dp_(dp), minDist_(minDist), cannyThreshold_(cannyThreshold), votesThreshold_(votesThreshold),
minRadius_(minRadius), maxRadius_(maxRadius), maxCircles_(maxCircles)
{
canny_ = cuda::createCannyEdgeDetector(std::max(cannyThreshold_ / 2, 1), cannyThreshold_);
filterDx_ = cuda::createSobelFilter(CV_8UC1, CV_32S, 1, 0);
filterDy_ = cuda::createSobelFilter(CV_8UC1, CV_32S, 0, 1);
}
void HoughCirclesDetectorImpl::detect(InputArray _src, OutputArray circles, Stream& stream)
{
// TODO : implement async version
CV_UNUSED(stream);
using namespace cv::cuda::device::hough;
using namespace cv::cuda::device::hough_circles;
GpuMat src = _src.getGpuMat();
CV_Assert( src.type() == CV_8UC1 );
CV_Assert( src.cols < std::numeric_limits<unsigned short>::max() );
CV_Assert( src.rows < std::numeric_limits<unsigned short>::max() );
CV_Assert( dp_ > 0 );
CV_Assert( minRadius_ > 0 && maxRadius_ > minRadius_ );
CV_Assert( cannyThreshold_ > 0 );
CV_Assert( votesThreshold_ > 0 );
CV_Assert( maxCircles_ > 0 );
const float idp = 1.0f / dp_;
filterDx_->apply(src, dx_);
filterDy_->apply(src, dy_);
canny_->setLowThreshold(std::max(cannyThreshold_ / 2, 1));
canny_->setHighThreshold(cannyThreshold_);
canny_->detect(dx_, dy_, edges_);
ensureSizeIsEnough(2, src.size().area(), CV_32SC1, list_);
unsigned int* srcPoints = list_.ptr<unsigned int>(0);
unsigned int* centers = list_.ptr<unsigned int>(1);
const int pointsCount = buildPointList_gpu(edges_, srcPoints);
if (pointsCount == 0)
{
circles.release();
return;
}
ensureSizeIsEnough(cvCeil(src.rows * idp) + 2, cvCeil(src.cols * idp) + 2, CV_32SC1, accum_);
accum_.setTo(Scalar::all(0));
circlesAccumCenters_gpu(srcPoints, pointsCount, dx_, dy_, accum_, minRadius_, maxRadius_, idp);
accum_.download(tt);
int centersCount = buildCentersList_gpu(accum_, centers, votesThreshold_);
if (centersCount == 0)
{
circles.release();
return;
}
//troy
//fprintf(stderr,"mindist is: %f!\n",minDist_);
//fflush(stderr);
if (minDist_ > 1)
{
AutoBuffer<ushort2> oldBuf_(centersCount);
AutoBuffer<ushort2> newBuf_(centersCount);
int newCount = 0;
ushort2* oldBuf = oldBuf_.data();
ushort2* newBuf = newBuf_.data();
cudaSafeCall( cudaMemcpy(oldBuf, centers, centersCount * sizeof(ushort2), cudaMemcpyDeviceToHost) );
const int cellSize = cvRound(minDist_);
const int gridWidth = (src.cols + cellSize - 1) / cellSize;
const int gridHeight = (src.rows + cellSize - 1) / cellSize;
std::vector< std::vector<ushort2> > grid(gridWidth * gridHeight);
const float minDist2 = minDist_ * minDist_;
std::vector<Vec3f> sortBuf;
for(int i=0; i<centersCount; i++){
Vec3f temp;
temp[0] = oldBuf[i].x;
temp[1] = oldBuf[i].y;
temp[2] = tt.at<int>(temp[1]+1, temp[0]+1);
sortBuf.push_back(temp);
}
std::sort(sortBuf.begin(), sortBuf.end(), centersCompare);
for (int i = 0; i < centersCount; ++i)
{
ushort2 p;
p.x = sortBuf[i][0];
p.y = sortBuf[i][1];
bool good = true;
int xCell = static_cast<int>(p.x / cellSize);
int yCell = static_cast<int>(p.y / cellSize);
int x1 = xCell - 1;
int y1 = yCell - 1;
int x2 = xCell + 1;
int y2 = yCell + 1;
// boundary check
x1 = std::max(0, x1);
y1 = std::max(0, y1);
x2 = std::min(gridWidth - 1, x2);
y2 = std::min(gridHeight - 1, y2);
for (int yy = y1; yy <= y2; ++yy)
{
for (int xx = x1; xx <= x2; ++xx)
{
std::vector<ushort2>& m = grid[yy * gridWidth + xx];
for(size_t j = 0; j < m.size(); ++j)
{
float dx = (float)(p.x - m[j].x);
float dy = (float)(p.y - m[j].y);
float dx2 = dx*dx;
float dy2 = dy*dy;
//troy
//fprintf(stderr,"%d: %f vs %f\n",p.x,minDist2,(dx2+dy2));
//fflush(stderr);
if (dx * dx + dy * dy < minDist2)
{
good = false;
goto break_out;
}
}
}
}
break_out:
if(good)
{
//troy
fprintf(stderr,"adding (%d,%d)!\n",p.x,p.y);
fflush(stderr);
grid[yCell * gridWidth + xCell].push_back(p);
newBuf[newCount++] = p;
}
}
cudaSafeCall( cudaMemcpy(centers, newBuf, newCount * sizeof(unsigned int), cudaMemcpyHostToDevice) );
fprintf(stderr,"centersCount %d -> %d\n",centersCount,newCount);
fflush(stderr);
centersCount = newCount;
}
ensureSizeIsEnough(1, maxCircles_, CV_32FC3, result_);
int circlesCount = circlesAccumRadius_gpu(centers, centersCount, srcPoints, pointsCount, result_.ptr<float3>(), maxCircles_,
dp_, minRadius_, maxRadius_, votesThreshold_, deviceSupports(FEATURE_SET_COMPUTE_20));
fprintf(stderr,"circlesCount: %d\n",circlesCount);
fflush(stderr);
if (circlesCount == 0)
{
circles.release();
return;
}
result_.cols = centersCount;
result_.copyTo(circles);
}
}
Ptr<HoughCirclesDetector> cv::cuda::createHoughCirclesDetector(float dp, float minDist, int cannyThreshold, int votesThreshold, int minRadius, int maxRadius, int maxCircles)
{
return makePtr<HoughCirclesDetectorImpl>(dp, minDist, cannyThreshold, votesThreshold, minRadius, maxRadius, maxCircles);
}
#endif /* !defined (HAVE_CUDA) */
ctop
[[ 8910 252 53]
[ 8910 252 55]
[ 5930 260 23]
[ 5930 260 26]
[ 5930 260 29]
[ 5930 260 35]
[ 5930 260 48]
[ 5930 260 51]
[ 4906 270 25]
[ 4906 270 33]
[ 4906 270 50]
[ 8910 252 21]
[ 8910 252 24]
[ 8910 252 26]
[ 8910 252 30]
[ 8910 252 35]
[ 8910 252 37]
[ 8910 252 47]
[ 8910 252 49]
[11400 240 24]
[11400 240 27]
[11400 240 30]
[11400 240 33]
[11400 240 36]
[11400 240 47]
[11400 240 50]
[ 4906 270 53]
[ 5930 260 53]
[ 5930 260 55]
[11400 240 52]
[11400 240 54]
[11400 240 57]
[ 3928 272 24]
[ 3928 272 26]
[ 3928 272 28]
[ 3928 272 32]
[ 3928 272 34]
[ 3928 272 49]
[13906 224 53]
[13906 224 24]
[13906 224 29]
[13906 224 32]
[13906 224 35]
[13906 224 42]
[13906 224 47]
[13906 224 50]
[16944 202 52]
[16944 202 26]
[16944 202 33]
[21786 180 52]
[21786 180 61]
[ 7890 262 52]
[ 7890 262 54]
[ 7890 262 24]
[ 7890 262 27]
[ 7890 262 32]
[ 7890 262 38]
[ 7890 262 42]
[21786 180 26]
[21786 180 28]
[21786 180 31]
[21786 180 34]
[21786 180 46]
[21786 180 49]
[20648 184 54]
[20648 184 22]
[20648 184 29]
[20648 184 33]
[20648 184 37]
[20648 184 45]
[20648 184 47]
[20648 184 50]
[12580 226 26]
[12580 226 28]
[12580 226 31]
[12580 226 34]
[12580 226 47]
[12580 226 50]
[12580 226 53]
[12580 226 58]
[19954 186 52]
[19954 186 58]
[19954 186 25]
[19954 186 28]
[19954 186 34]
[19954 186 49]
[14948 220 57]
[14948 220 26]
[14948 220 32]
[14948 220 35]
[14948 220 44]
[14948 220 46]
[14948 220 50]
[ 2890 272 24]
[ 2890 272 26]
[ 2890 272 30]
[ 2890 272 33]
[ 2890 272 35]
[ 2890 272 47]
[ 2890 272 50]
[ 2890 272 53]
[ 2890 272 55]
[17940 198 52]
[17940 198 55]
[17940 198 57]
[17940 198 24]
[17940 198 27]
[17940 198 32]
[17940 198 34]
[17940 198 47]
[17940 198 49]
[18994 192 53]
[18994 192 24]
[18994 192 30]
[18994 192 34]
[18994 192 36]
[18994 192 50]
[13260 226 22]
[13260 226 25]
[13260 226 27]
[13260 226 30]
[13260 226 37]
[13260 226 45]
[13260 226 48]
[13260 226 52]
[13260 226 54]
[13260 226 59]
[10274 238 54]
[10274 238 26]
[10274 238 33]
[10274 238 51]
[15962 216 23]
[15962 216 26]
[15962 216 29]
[15962 216 32]
[15962 216 36]
[15962 216 45]
[15962 216 50]
[15962 216 53]
[15962 216 55]
[ 6884 254 52]
[ 6884 254 55]
[ 6884 254 25]
[ 6884 254 27]
[ 6884 254 31]
[ 6884 254 33]
[ 6884 254 50]]
cottom
[[ 7924 2769 54]
[ 7924 2769 23]
[ 7924 2769 26]
[ 7924 2769 28]
[ 7924 2769 30]
[ 7924 2769 34]
[ 7924 2769 36]
[ 7924 2769 48]
[19032 2707 25]
[19032 2707 27]
[19032 2707 31]
[19032 2707 33]
[19032 2707 45]
[19032 2707 51]
[19032 2707 54]
[13272 2745 25]
[13272 2745 27]
[13272 2745 31]
[13272 2745 34]
[13272 2745 51]
[ 5914 2783 26]
[ 5914 2783 33]
[ 5914 2783 50]
[ 5914 2783 53]
[19994 2703 25]
[19994 2703 33]
[19994 2703 47]
[19994 2703 50]
[13272 2745 55]
[19994 2703 52]
[ 8944 2763 54]
[12594 2745 23]
[12594 2745 26]
[12594 2745 28]
[12594 2745 31]
[12594 2745 33]
[12594 2745 48]
[12594 2745 50]
[ 3938 2785 22]
[ 3938 2785 24]
[ 3938 2785 26]
[ 3938 2785 29]
[ 3938 2785 32]
[ 3938 2785 37]
[ 3938 2785 48]
[ 3938 2785 51]
[14944 2739 52]
[14944 2739 56]
[ 8944 2763 25]
[ 8944 2763 27]
[ 8944 2763 33]
[ 8944 2763 51]
[12594 2745 53]
[12594 2745 56]
[12594 2745 59]
[ 3938 2785 55]
[14944 2739 25]
[14944 2739 27]
[14944 2739 33]
[14944 2739 50]
[ 6886 2777 54]
[ 6886 2777 25]
[ 6886 2777 27]
[ 6886 2777 33]
[ 6886 2777 51]
[11432 2749 55]
[11432 2749 27]
[11432 2749 31]
[11432 2749 33]
[11432 2749 51]
[16964 2735 26]
[16964 2735 32]
[16964 2735 49]
[16964 2735 52]
[16964 2735 55]
[16964 2735 57]
[10274 2755 25]
[10274 2755 27]
[10274 2755 33]
[10274 2755 50]
[10274 2755 52]
[10274 2755 55]
[20650 2711 28]
[20650 2711 31]
[20650 2711 34]
[20650 2711 48]
[20650 2711 53]
[20650 2711 55]
[ 2924 2783 21]
[ 2924 2783 24]
[ 2924 2783 28]
[ 2924 2783 31]
[ 2924 2783 33]
[ 2924 2783 38]
[ 2924 2783 46]
[ 2924 2783 48]
[ 2924 2783 53]
[ 2924 2783 56]
[13940 2749 54]
[13940 2749 24]
[13940 2749 26]
[13940 2749 30]
[13940 2749 33]
[13940 2749 35]
[13940 2749 50]
[ 4926 2783 24]
[ 4926 2783 27]
[ 4926 2783 30]
[ 4926 2783 35]
[ 4926 2783 47]
[ 4926 2783 50]
[17978 2723 53]
[ 4926 2783 53]
[17978 2723 24]
[17978 2723 26]
[17978 2723 30]
[17978 2723 36]
[17978 2723 49]
[21824 2699 22]
[21824 2699 28]
[21824 2699 30]
[21824 2699 37]
[21824 2699 47]
[21824 2699 52]
[21824 2699 55]
[21824 2699 60]
[15962 2733 52]
[15962 2733 57]
[15962 2733 24]
[15962 2733 28]
[15962 2733 31]
[15962 2733 34]
[15962 2733 40]
[15962 2733 46]
[15962 2733 50]]
adding (12579,226)!
adding (6884,253)!
adding (13259,225)!
adding (14948,220)!
adding (19954,185)!
adding (15961,216)!
adding (10274,237)!
adding (8909,252)!
adding (5930,260)!
adding (4905,269)!
adding (13906,224)!
adding (16943,202)!
adding (11399,240)!
adding (18993,191)!
adding (7889,261)!
adding (20648,184)!
adding (2889,272)!
adding (17939,198)!
adding (3927,271)!
adding (21786,180)!
centersCount 4524 -> 20
circlesCount: 147
adding (10273,1230)!
adding (5914,1257)!
adding (11432,1224)!
adding (8943,1238)!
adding (16964,1210)!
adding (19994,1178)!
adding (13272,1220)!
adding (6886,1252)!
adding (12594,1220)!
adding (7924,1244)!
adding (19031,1182)!
adding (3938,1259)!
adding (14944,1214)!
adding (20649,1185)!
adding (17978,1197)!
adding (4926,1258)!
adding (15961,1208)!
adding (2923,1258)!
adding (13939,1223)!
adding (21824,1173)!
centersCount 4472 -> 20
circlesCount: 135
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment