mirror of
https://github.com/opencv/opencv.git
synced 2026-07-30 15:53:03 +04:00
refactored generalized hough (both CPU and GPU):
removed set/get methods from Algorithm (implement owns) removed GHT_* enumeration
This commit is contained in:
@@ -213,98 +213,19 @@ Creates implementation for :ocv:class:`gpu::HoughCirclesDetector` .
|
||||
|
||||
|
||||
|
||||
gpu::GeneralizedHough
|
||||
---------------------
|
||||
.. ocv:class:: gpu::GeneralizedHough : public Algorithm
|
||||
|
||||
Base class for generalized hough transform. ::
|
||||
|
||||
class CV_EXPORTS GeneralizedHough : public Algorithm
|
||||
{
|
||||
public:
|
||||
static Ptr<GeneralizedHough> create(int method);
|
||||
|
||||
virtual void setTemplate(InputArray templ, int cannyThreshold = 100, Point templCenter = Point(-1, -1)) = 0;
|
||||
virtual void setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter = Point(-1, -1)) = 0;
|
||||
|
||||
virtual void detect(InputArray image, OutputArray positions, int cannyThreshold = 100) = 0;
|
||||
virtual void detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions) = 0;
|
||||
|
||||
virtual void downloadResults(InputArray d_positions, OutputArray h_positions, OutputArray h_votes = noArray()) = 0;
|
||||
};
|
||||
|
||||
|
||||
Finds arbitrary template in the grayscale image using Generalized Hough Transform.
|
||||
|
||||
|
||||
|
||||
gpu::GeneralizedHough::create
|
||||
-----------------------------
|
||||
Creates implementation for :ocv:class:`gpu::GeneralizedHough` .
|
||||
|
||||
.. ocv:function:: Ptr<GeneralizedHough> gpu::GeneralizedHough::create(int method)
|
||||
|
||||
:param method: Combination of flags ( ``cv::GeneralizedHough::GHT_POSITION`` , ``cv::GeneralizedHough::GHT_SCALE`` , ``cv::GeneralizedHough::GHT_ROTATION`` ) specifying transformation to find.
|
||||
|
||||
For full affine transformations (move + scale + rotation) [Guil1999]_ algorithm is used, otherwise [Ballard1981]_ algorithm is used.
|
||||
|
||||
|
||||
|
||||
gpu::GeneralizedHough::setTemplate
|
||||
gpu::createGeneralizedHoughBallard
|
||||
----------------------------------
|
||||
Set template to search.
|
||||
Creates implementation for generalized hough transform from [Ballard1981]_ .
|
||||
|
||||
.. ocv:function:: void gpu::GeneralizedHough::setTemplate(InputArray templ, int cannyThreshold = 100, Point templCenter = Point(-1, -1))
|
||||
|
||||
.. ocv:function:: void gpu::GeneralizedHough::setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter = Point(-1, -1))
|
||||
|
||||
:param templ: Template image. Canny edge detector will be applied to extract template edges.
|
||||
|
||||
:param cannyThreshold: Threshold value for Canny edge detector.
|
||||
|
||||
:param templCenter: Center for rotation. By default image center will be used.
|
||||
|
||||
:param edges: Edge map for template image.
|
||||
|
||||
:param dx: First derivative of template image in the vertical direction. Support only ``CV_32S`` type.
|
||||
|
||||
:param dy: First derivative of template image in the horizontal direction. Support only ``CV_32S`` type.
|
||||
.. ocv:function:: Ptr<GeneralizedHoughBallard> gpu::createGeneralizedHoughBallard()
|
||||
|
||||
|
||||
|
||||
gpu::GeneralizedHough::detect
|
||||
-----------------------------
|
||||
Finds template (set by :ocv:func:`gpu::GeneralizedHough::setTemplate` ) in the grayscale image.
|
||||
gpu::createGeneralizedHoughGuil
|
||||
-------------------------------
|
||||
Creates implementation for generalized hough transform from [Guil1999]_ .
|
||||
|
||||
.. ocv:function:: void gpu::GeneralizedHough::detect(InputArray image, OutputArray positions, int cannyThreshold = 100)
|
||||
|
||||
.. ocv:function:: void gpu::GeneralizedHough::detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions)
|
||||
|
||||
:param templ: Input image. Canny edge detector will be applied to extract template edges.
|
||||
|
||||
:param positions: Output vector of found objects. Each vector is encoded as a 4-element floating-point vector :math:`(x, y, scale, angle)` .
|
||||
|
||||
:param cannyThreshold: Threshold value for Canny edge detector.
|
||||
|
||||
:param edges: Edge map for input image.
|
||||
|
||||
:param dx: First derivative of input image in the vertical direction. Support only ``CV_32S`` type.
|
||||
|
||||
:param dy: First derivative of input image in the horizontal direction. Support only ``CV_32S`` type.
|
||||
|
||||
|
||||
|
||||
gpu::GeneralizedHough::downloadResults
|
||||
--------------------------------------
|
||||
Downloads results from :ocv:func:`gpu::GeneralizedHough::detect` to host memory.
|
||||
|
||||
.. ocv:function:: void gpu::GeneralizedHough::downloadResult(InputArray d_positions, OutputArray h_positions, OutputArray h_votes = noArray())
|
||||
|
||||
:param d_lines: Result of :ocv:func:`gpu::GeneralizedHough::detect` .
|
||||
|
||||
:param h_lines: Output host array.
|
||||
|
||||
:param h_votes: Optional output array for votes. Each vector is encoded as a 3-element integer-point vector :math:`(position_votes, scale_votes, angle_votes)` .
|
||||
.. ocv:function:: Ptr<GeneralizedHoughGuil> gpu::createGeneralizedHoughGuil()
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -283,24 +283,13 @@ CV_EXPORTS Ptr<HoughCirclesDetector> createHoughCirclesDetector(float dp, float
|
||||
//////////////////////////////////////
|
||||
// GeneralizedHough
|
||||
|
||||
//! finds arbitrary template in the grayscale image using Generalized Hough Transform
|
||||
//! Ballard, D.H. (1981). Generalizing the Hough transform to detect arbitrary shapes. Pattern Recognition 13 (2): 111-122.
|
||||
//! Detects position only without traslation and rotation
|
||||
CV_EXPORTS Ptr<GeneralizedHoughBallard> createGeneralizedHoughBallard();
|
||||
|
||||
//! Guil, N., González-Linares, J.M. and Zapata, E.L. (1999). Bidimensional shape detection using an invariant approach. Pattern Recognition 32 (6): 1025-1038.
|
||||
class CV_EXPORTS GeneralizedHough : public Algorithm
|
||||
{
|
||||
public:
|
||||
static Ptr<GeneralizedHough> create(int method);
|
||||
|
||||
//! set template to search
|
||||
virtual void setTemplate(InputArray templ, int cannyThreshold = 100, Point templCenter = Point(-1, -1)) = 0;
|
||||
virtual void setTemplate(InputArray edges, InputArray dx, InputArray dy, Point templCenter = Point(-1, -1)) = 0;
|
||||
|
||||
//! find template on image
|
||||
virtual void detect(InputArray image, OutputArray positions, int cannyThreshold = 100) = 0;
|
||||
virtual void detect(InputArray edges, InputArray dx, InputArray dy, OutputArray positions) = 0;
|
||||
|
||||
virtual void downloadResults(InputArray d_positions, OutputArray h_positions, OutputArray h_votes = noArray()) = 0;
|
||||
};
|
||||
//! Detects position, traslation and rotation
|
||||
CV_EXPORTS Ptr<GeneralizedHoughGuil> createGeneralizedHoughGuil();
|
||||
|
||||
////////////////////////// Corners Detection ///////////////////////////
|
||||
|
||||
|
||||
@@ -227,23 +227,59 @@ PERF_TEST_P(Sz_Dp_MinDist, HoughCircles,
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// GeneralizedHough
|
||||
|
||||
enum { GHT_POSITION = cv::GeneralizedHough::GHT_POSITION,
|
||||
GHT_SCALE = cv::GeneralizedHough::GHT_SCALE,
|
||||
GHT_ROTATION = cv::GeneralizedHough::GHT_ROTATION
|
||||
};
|
||||
|
||||
CV_FLAGS(GHMethod, GHT_POSITION, GHT_SCALE, GHT_ROTATION);
|
||||
|
||||
DEF_PARAM_TEST(Method_Sz, GHMethod, cv::Size);
|
||||
|
||||
PERF_TEST_P(Method_Sz, GeneralizedHough,
|
||||
Combine(Values(GHMethod(GHT_POSITION), GHMethod(GHT_POSITION | GHT_SCALE), GHMethod(GHT_POSITION | GHT_ROTATION), GHMethod(GHT_POSITION | GHT_SCALE | GHT_ROTATION)),
|
||||
GPU_TYPICAL_MAT_SIZES))
|
||||
PERF_TEST_P(Sz, GeneralizedHoughBallard, GPU_TYPICAL_MAT_SIZES)
|
||||
{
|
||||
declare.time(10);
|
||||
|
||||
const int method = GET_PARAM(0);
|
||||
const cv::Size imageSize = GET_PARAM(1);
|
||||
const cv::Size imageSize = GetParam();
|
||||
|
||||
const cv::Mat templ = readImage("cv/shared/templ.png", cv::IMREAD_GRAYSCALE);
|
||||
ASSERT_FALSE(templ.empty());
|
||||
|
||||
cv::Mat image(imageSize, CV_8UC1, cv::Scalar::all(0));
|
||||
templ.copyTo(image(cv::Rect(50, 50, templ.cols, templ.rows)));
|
||||
|
||||
cv::Mat edges;
|
||||
cv::Canny(image, edges, 50, 100);
|
||||
|
||||
cv::Mat dx, dy;
|
||||
cv::Sobel(image, dx, CV_32F, 1, 0);
|
||||
cv::Sobel(image, dy, CV_32F, 0, 1);
|
||||
|
||||
if (PERF_RUN_GPU())
|
||||
{
|
||||
cv::Ptr<cv::GeneralizedHoughBallard> alg = cv::gpu::createGeneralizedHoughBallard();
|
||||
|
||||
const cv::gpu::GpuMat d_edges(edges);
|
||||
const cv::gpu::GpuMat d_dx(dx);
|
||||
const cv::gpu::GpuMat d_dy(dy);
|
||||
cv::gpu::GpuMat positions;
|
||||
|
||||
alg->setTemplate(cv::gpu::GpuMat(templ));
|
||||
|
||||
TEST_CYCLE() alg->detect(d_edges, d_dx, d_dy, positions);
|
||||
|
||||
GPU_SANITY_CHECK(positions);
|
||||
}
|
||||
else
|
||||
{
|
||||
cv::Ptr<cv::GeneralizedHoughBallard> alg = cv::createGeneralizedHoughBallard();
|
||||
|
||||
cv::Mat positions;
|
||||
|
||||
alg->setTemplate(templ);
|
||||
|
||||
TEST_CYCLE() alg->detect(edges, dx, dy, positions);
|
||||
|
||||
CPU_SANITY_CHECK(positions);
|
||||
}
|
||||
}
|
||||
|
||||
PERF_TEST_P(Sz, GeneralizedHoughGuil, GPU_TYPICAL_MAT_SIZES)
|
||||
{
|
||||
declare.time(10);
|
||||
|
||||
const cv::Size imageSize = GetParam();
|
||||
|
||||
const cv::Mat templ = readImage("cv/shared/templ.png", cv::IMREAD_GRAYSCALE);
|
||||
ASSERT_FALSE(templ.empty());
|
||||
@@ -281,39 +317,32 @@ PERF_TEST_P(Method_Sz, GeneralizedHough,
|
||||
|
||||
if (PERF_RUN_GPU())
|
||||
{
|
||||
cv::Ptr<cv::GeneralizedHoughGuil> alg = cv::gpu::createGeneralizedHoughGuil();
|
||||
alg->setMaxAngle(90.0);
|
||||
alg->setAngleStep(2.0);
|
||||
|
||||
const cv::gpu::GpuMat d_edges(edges);
|
||||
const cv::gpu::GpuMat d_dx(dx);
|
||||
const cv::gpu::GpuMat d_dy(dy);
|
||||
cv::gpu::GpuMat posAndVotes;
|
||||
cv::gpu::GpuMat positions;
|
||||
|
||||
cv::Ptr<cv::gpu::GeneralizedHough> d_hough = cv::gpu::GeneralizedHough::create(method);
|
||||
if (method & GHT_ROTATION)
|
||||
{
|
||||
d_hough->set("maxAngle", 90.0);
|
||||
d_hough->set("angleStep", 2.0);
|
||||
}
|
||||
alg->setTemplate(cv::gpu::GpuMat(templ));
|
||||
|
||||
d_hough->setTemplate(cv::gpu::GpuMat(templ));
|
||||
TEST_CYCLE() alg->detect(d_edges, d_dx, d_dy, positions);
|
||||
|
||||
TEST_CYCLE() d_hough->detect(d_edges, d_dx, d_dy, posAndVotes);
|
||||
|
||||
const cv::gpu::GpuMat positions(1, posAndVotes.cols, CV_32FC4, posAndVotes.data);
|
||||
GPU_SANITY_CHECK(positions);
|
||||
}
|
||||
else
|
||||
{
|
||||
cv::Ptr<cv::GeneralizedHoughGuil> alg = cv::createGeneralizedHoughGuil();
|
||||
alg->setMaxAngle(90.0);
|
||||
alg->setAngleStep(2.0);
|
||||
|
||||
cv::Mat positions;
|
||||
|
||||
cv::Ptr<cv::GeneralizedHough> hough = cv::GeneralizedHough::create(method);
|
||||
if (method & GHT_ROTATION)
|
||||
{
|
||||
hough->set("maxAngle", 90.0);
|
||||
hough->set("angleStep", 2.0);
|
||||
}
|
||||
alg->setTemplate(templ);
|
||||
|
||||
hough->setTemplate(templ);
|
||||
|
||||
TEST_CYCLE() hough->detect(edges, dx, dy, positions);
|
||||
TEST_CYCLE() alg->detect(edges, dx, dy, positions);
|
||||
|
||||
CPU_SANITY_CHECK(positions);
|
||||
}
|
||||
|
||||
@@ -307,268 +307,6 @@ namespace cv { namespace gpu { namespace cudev
|
||||
return totalCount;
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
// Ballard_PosScale
|
||||
|
||||
__global__ void Ballard_PosScale_calcHist(const unsigned int* coordList, const float* thetaList,
|
||||
PtrStep<short2> r_table, const int* r_sizes,
|
||||
PtrStepi hist, const int rows, const int cols,
|
||||
const float minScale, const float scaleStep, const int scaleRange,
|
||||
const float idp, const float thetaScale)
|
||||
{
|
||||
const unsigned int coord = coordList[blockIdx.x];
|
||||
float2 p;
|
||||
p.x = (coord & 0xFFFF);
|
||||
p.y = (coord >> 16) & 0xFFFF;
|
||||
|
||||
const float theta = thetaList[blockIdx.x];
|
||||
const int n = __float2int_rn(theta * thetaScale);
|
||||
|
||||
const short2* r_row = r_table.ptr(n);
|
||||
const int r_row_size = r_sizes[n];
|
||||
|
||||
for (int j = 0; j < r_row_size; ++j)
|
||||
{
|
||||
const float2 d = saturate_cast<float2>(r_row[j]);
|
||||
|
||||
for (int s = threadIdx.x; s < scaleRange; s += blockDim.x)
|
||||
{
|
||||
const float scale = minScale + s * scaleStep;
|
||||
|
||||
float2 c = p - scale * d;
|
||||
|
||||
c.x *= idp;
|
||||
c.y *= idp;
|
||||
|
||||
if (c.x >= 0 && c.x < cols && c.y >= 0 && c.y < rows)
|
||||
::atomicAdd(hist.ptr((s + 1) * (rows + 2) + __float2int_rn(c.y + 1)) + __float2int_rn(c.x + 1), 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void Ballard_PosScale_calcHist_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
|
||||
PtrStepSz<short2> r_table, const int* r_sizes,
|
||||
PtrStepi hist, int rows, int cols,
|
||||
float minScale, float scaleStep, int scaleRange,
|
||||
float dp, int levels)
|
||||
{
|
||||
const dim3 block(256);
|
||||
const dim3 grid(pointsCount);
|
||||
|
||||
const float idp = 1.0f / dp;
|
||||
const float thetaScale = levels / (2.0f * CV_PI_F);
|
||||
|
||||
Ballard_PosScale_calcHist<<<grid, block>>>(coordList, thetaList,
|
||||
r_table, r_sizes,
|
||||
hist, rows, cols,
|
||||
minScale, scaleStep, scaleRange,
|
||||
idp, thetaScale);
|
||||
cudaSafeCall( cudaGetLastError() );
|
||||
|
||||
cudaSafeCall( cudaDeviceSynchronize() );
|
||||
}
|
||||
|
||||
__global__ void Ballard_PosScale_findPosInHist(const PtrStepi hist, const int rows, const int cols, const int scaleRange,
|
||||
float4* out, int3* votes, const int maxSize,
|
||||
const float minScale, const float scaleStep, const float dp, const int threshold)
|
||||
{
|
||||
const int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
const int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
|
||||
if (x >= cols || y >= rows)
|
||||
return;
|
||||
|
||||
for (int s = 0; s < scaleRange; ++s)
|
||||
{
|
||||
const float scale = minScale + s * scaleStep;
|
||||
|
||||
const int prevScaleIdx = (s) * (rows + 2);
|
||||
const int curScaleIdx = (s + 1) * (rows + 2);
|
||||
const int nextScaleIdx = (s + 2) * (rows + 2);
|
||||
|
||||
const int curVotes = hist(curScaleIdx + y + 1, x + 1);
|
||||
|
||||
if (curVotes > threshold &&
|
||||
curVotes > hist(curScaleIdx + y + 1, x) &&
|
||||
curVotes >= hist(curScaleIdx + y + 1, x + 2) &&
|
||||
curVotes > hist(curScaleIdx + y, x + 1) &&
|
||||
curVotes >= hist(curScaleIdx + y + 2, x + 1) &&
|
||||
curVotes > hist(prevScaleIdx + y + 1, x + 1) &&
|
||||
curVotes >= hist(nextScaleIdx + y + 1, x + 1))
|
||||
{
|
||||
const int ind = ::atomicAdd(&g_counter, 1);
|
||||
|
||||
if (ind < maxSize)
|
||||
{
|
||||
out[ind] = make_float4(x * dp, y * dp, scale, 0.0f);
|
||||
votes[ind] = make_int3(curVotes, curVotes, 0);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int Ballard_PosScale_findPosInHist_gpu(PtrStepi hist, int rows, int cols, int scaleRange, float4* out, int3* votes, int maxSize,
|
||||
float minScale, float scaleStep, float dp, int threshold)
|
||||
{
|
||||
void* counterPtr;
|
||||
cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
|
||||
|
||||
cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
|
||||
|
||||
const dim3 block(32, 8);
|
||||
const dim3 grid(divUp(cols, block.x), divUp(rows, block.y));
|
||||
|
||||
cudaSafeCall( cudaFuncSetCacheConfig(Ballard_PosScale_findPosInHist, cudaFuncCachePreferL1) );
|
||||
|
||||
Ballard_PosScale_findPosInHist<<<grid, block>>>(hist, rows, cols, scaleRange, out, votes,
|
||||
maxSize, minScale, scaleStep, dp, threshold);
|
||||
cudaSafeCall( cudaGetLastError() );
|
||||
|
||||
cudaSafeCall( cudaDeviceSynchronize() );
|
||||
|
||||
int totalCount;
|
||||
cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
|
||||
|
||||
totalCount = ::min(totalCount, maxSize);
|
||||
|
||||
return totalCount;
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
// Ballard_PosRotation
|
||||
|
||||
__global__ void Ballard_PosRotation_calcHist(const unsigned int* coordList, const float* thetaList,
|
||||
PtrStep<short2> r_table, const int* r_sizes,
|
||||
PtrStepi hist, const int rows, const int cols,
|
||||
const float minAngle, const float angleStep, const int angleRange,
|
||||
const float idp, const float thetaScale)
|
||||
{
|
||||
const unsigned int coord = coordList[blockIdx.x];
|
||||
float2 p;
|
||||
p.x = (coord & 0xFFFF);
|
||||
p.y = (coord >> 16) & 0xFFFF;
|
||||
|
||||
const float thetaVal = thetaList[blockIdx.x];
|
||||
|
||||
for (int a = threadIdx.x; a < angleRange; a += blockDim.x)
|
||||
{
|
||||
const float angle = (minAngle + a * angleStep) * (CV_PI_F / 180.0f);
|
||||
float sinA, cosA;
|
||||
sincosf(angle, &sinA, &cosA);
|
||||
|
||||
float theta = thetaVal - angle;
|
||||
if (theta < 0)
|
||||
theta += 2.0f * CV_PI_F;
|
||||
|
||||
const int n = __float2int_rn(theta * thetaScale);
|
||||
|
||||
const short2* r_row = r_table.ptr(n);
|
||||
const int r_row_size = r_sizes[n];
|
||||
|
||||
for (int j = 0; j < r_row_size; ++j)
|
||||
{
|
||||
const float2 d = saturate_cast<float2>(r_row[j]);
|
||||
|
||||
const float2 dr = make_float2(d.x * cosA - d.y * sinA, d.x * sinA + d.y * cosA);
|
||||
|
||||
float2 c = make_float2(p.x - dr.x, p.y - dr.y);
|
||||
c.x *= idp;
|
||||
c.y *= idp;
|
||||
|
||||
if (c.x >= 0 && c.x < cols && c.y >= 0 && c.y < rows)
|
||||
::atomicAdd(hist.ptr((a + 1) * (rows + 2) + __float2int_rn(c.y + 1)) + __float2int_rn(c.x + 1), 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
void Ballard_PosRotation_calcHist_gpu(const unsigned int* coordList, const float* thetaList, int pointsCount,
|
||||
PtrStepSz<short2> r_table, const int* r_sizes,
|
||||
PtrStepi hist, int rows, int cols,
|
||||
float minAngle, float angleStep, int angleRange,
|
||||
float dp, int levels)
|
||||
{
|
||||
const dim3 block(256);
|
||||
const dim3 grid(pointsCount);
|
||||
|
||||
const float idp = 1.0f / dp;
|
||||
const float thetaScale = levels / (2.0f * CV_PI_F);
|
||||
|
||||
Ballard_PosRotation_calcHist<<<grid, block>>>(coordList, thetaList,
|
||||
r_table, r_sizes,
|
||||
hist, rows, cols,
|
||||
minAngle, angleStep, angleRange,
|
||||
idp, thetaScale);
|
||||
cudaSafeCall( cudaGetLastError() );
|
||||
|
||||
cudaSafeCall( cudaDeviceSynchronize() );
|
||||
}
|
||||
|
||||
__global__ void Ballard_PosRotation_findPosInHist(const PtrStepi hist, const int rows, const int cols, const int angleRange,
|
||||
float4* out, int3* votes, const int maxSize,
|
||||
const float minAngle, const float angleStep, const float dp, const int threshold)
|
||||
{
|
||||
const int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
const int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
|
||||
if (x >= cols || y >= rows)
|
||||
return;
|
||||
|
||||
for (int a = 0; a < angleRange; ++a)
|
||||
{
|
||||
const float angle = minAngle + a * angleStep;
|
||||
|
||||
const int prevAngleIdx = (a) * (rows + 2);
|
||||
const int curAngleIdx = (a + 1) * (rows + 2);
|
||||
const int nextAngleIdx = (a + 2) * (rows + 2);
|
||||
|
||||
const int curVotes = hist(curAngleIdx + y + 1, x + 1);
|
||||
|
||||
if (curVotes > threshold &&
|
||||
curVotes > hist(curAngleIdx + y + 1, x) &&
|
||||
curVotes >= hist(curAngleIdx + y + 1, x + 2) &&
|
||||
curVotes > hist(curAngleIdx + y, x + 1) &&
|
||||
curVotes >= hist(curAngleIdx + y + 2, x + 1) &&
|
||||
curVotes > hist(prevAngleIdx + y + 1, x + 1) &&
|
||||
curVotes >= hist(nextAngleIdx + y + 1, x + 1))
|
||||
{
|
||||
const int ind = ::atomicAdd(&g_counter, 1);
|
||||
|
||||
if (ind < maxSize)
|
||||
{
|
||||
out[ind] = make_float4(x * dp, y * dp, 1.0f, angle);
|
||||
votes[ind] = make_int3(curVotes, 0, curVotes);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int Ballard_PosRotation_findPosInHist_gpu(PtrStepi hist, int rows, int cols, int angleRange, float4* out, int3* votes, int maxSize,
|
||||
float minAngle, float angleStep, float dp, int threshold)
|
||||
{
|
||||
void* counterPtr;
|
||||
cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
|
||||
|
||||
cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
|
||||
|
||||
const dim3 block(32, 8);
|
||||
const dim3 grid(divUp(cols, block.x), divUp(rows, block.y));
|
||||
|
||||
cudaSafeCall( cudaFuncSetCacheConfig(Ballard_PosRotation_findPosInHist, cudaFuncCachePreferL1) );
|
||||
|
||||
Ballard_PosRotation_findPosInHist<<<grid, block>>>(hist, rows, cols, angleRange, out, votes,
|
||||
maxSize, minAngle, angleStep, dp, threshold);
|
||||
cudaSafeCall( cudaGetLastError() );
|
||||
|
||||
cudaSafeCall( cudaDeviceSynchronize() );
|
||||
|
||||
int totalCount;
|
||||
cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
|
||||
|
||||
totalCount = ::min(totalCount, maxSize);
|
||||
|
||||
return totalCount;
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
// Guil_Full
|
||||
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -193,7 +193,7 @@ PARAM_TEST_CASE(GeneralizedHough, cv::gpu::DeviceInfo, UseRoi)
|
||||
{
|
||||
};
|
||||
|
||||
GPU_TEST_P(GeneralizedHough, POSITION)
|
||||
GPU_TEST_P(GeneralizedHough, Ballard)
|
||||
{
|
||||
const cv::gpu::DeviceInfo devInfo = GET_PARAM(0);
|
||||
cv::gpu::setDevice(devInfo.deviceID());
|
||||
@@ -218,16 +218,16 @@ GPU_TEST_P(GeneralizedHough, POSITION)
|
||||
templ.copyTo(imageROI);
|
||||
}
|
||||
|
||||
cv::Ptr<cv::gpu::GeneralizedHough> hough = cv::gpu::GeneralizedHough::create(cv::GeneralizedHough::GHT_POSITION);
|
||||
hough->set("votesThreshold", 200);
|
||||
cv::Ptr<cv::GeneralizedHoughBallard> alg = cv::gpu::createGeneralizedHoughBallard();
|
||||
alg->setVotesThreshold(200);
|
||||
|
||||
hough->setTemplate(loadMat(templ, useRoi));
|
||||
alg->setTemplate(loadMat(templ, useRoi));
|
||||
|
||||
cv::gpu::GpuMat d_pos;
|
||||
hough->detect(loadMat(image, useRoi), d_pos);
|
||||
alg->detect(loadMat(image, useRoi), d_pos);
|
||||
|
||||
std::vector<cv::Vec4f> pos;
|
||||
hough->downloadResults(d_pos, pos);
|
||||
d_pos.download(pos);
|
||||
|
||||
ASSERT_EQ(gold_count, pos.size());
|
||||
|
||||
|
||||
Reference in New Issue
Block a user