pull/11/head
Vladislav Vinogradov 12 years ago
parent db79022b85
commit 97731c152d
  1. 2
      modules/gpu/include/opencv2/gpu/gpu.hpp
  2. 12
      modules/gpu/src/cuda/hough.cu
  3. 18
      modules/gpu/src/hough.cpp

@ -824,7 +824,7 @@ CV_EXPORTS void HoughLines(const GpuMat& src, GpuMat& lines, float rho, float th
CV_EXPORTS void HoughLines(const GpuMat& src, GpuMat& lines, GpuMat& accum, GpuMat& buf, float rho, float theta, int threshold, bool doSort = false, int maxLines = 4096);
CV_EXPORTS void HoughLinesTransform(const GpuMat& src, GpuMat& accum, GpuMat& buf, float rho, float theta);
CV_EXPORTS void HoughLinesGet(const GpuMat& accum, GpuMat& lines, float rho, float theta, int threshold, bool doSort = false, int maxLines = 4096);
CV_EXPORTS void HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines, OutputArray h_voices = noArray());
CV_EXPORTS void HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines, OutputArray h_votes = noArray());
////////////////////////////// Matrix reductions //////////////////////////////

@ -225,7 +225,7 @@ namespace cv { namespace gpu { namespace device
////////////////////////////////////////////////////////////////////////
// linesGetResult
__global__ void linesGetResult(const DevMem2Di accum, float2* out, int* voices, const int maxSize, const float threshold, const float theta, const float rho, const int numrho)
__global__ void linesGetResult(const DevMem2Di accum, float2* out, int* votes, const int maxSize, const float threshold, const float theta, const float rho, const int numrho)
{
__shared__ int smem[8][32];
@ -257,12 +257,12 @@ namespace cv { namespace gpu { namespace device
if (ind < maxSize)
{
out[ind] = make_float2(radius, angle);
voices[ind] = smem[threadIdx.y][threadIdx.x];
votes[ind] = smem[threadIdx.y][threadIdx.x];
}
}
}
int linesGetResult_gpu(DevMem2Di accum, float2* out, int* voices, int maxSize, float rho, float theta, float threshold, bool doSort)
int linesGetResult_gpu(DevMem2Di accum, float2* out, int* votes, int maxSize, float rho, float theta, float threshold, bool doSort)
{
void* counter_ptr;
cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) );
@ -272,7 +272,7 @@ namespace cv { namespace gpu { namespace device
const dim3 block(32, 8);
const dim3 grid(divUp(accum.cols, block.x - 2), divUp(accum.rows, block.y - 2));
linesGetResult<<<grid, block>>>(accum, out, voices, maxSize, threshold, theta, rho, accum.cols - 2);
linesGetResult<<<grid, block>>>(accum, out, votes, maxSize, threshold, theta, rho, accum.cols - 2);
cudaSafeCall( cudaGetLastError() );
cudaSafeCall( cudaDeviceSynchronize() );
@ -285,8 +285,8 @@ namespace cv { namespace gpu { namespace device
if (doSort && total_count > 0)
{
thrust::device_ptr<float2> out_ptr(out);
thrust::device_ptr<int> voices_ptr(voices);
thrust::sort_by_key(voices_ptr, voices_ptr + total_count, out_ptr, thrust::greater<int>());
thrust::device_ptr<int> votes_ptr(votes);
thrust::sort_by_key(votes_ptr, votes_ptr + total_count, out_ptr, thrust::greater<int>());
}
return total_count;

@ -58,7 +58,7 @@ namespace cv { namespace gpu { namespace device
{
int buildPointList_gpu(DevMem2Db src, unsigned int* list);
void linesAccum_gpu(const unsigned int* list, int count, DevMem2Di accum, float rho, float theta, size_t sharedMemPerBlock);
int linesGetResult_gpu(DevMem2Di accum, float2* out, int* voices, int maxSize, float rho, float theta, float threshold, bool doSort);
int linesGetResult_gpu(DevMem2Di accum, float2* out, int* votes, int maxSize, float rho, float theta, float threshold, bool doSort);
}
}}}
@ -116,13 +116,13 @@ void cv::gpu::HoughLines(const GpuMat& src, GpuMat& lines, GpuMat& accum, GpuMat
HoughLinesGet(accum, lines, rho, theta, threshold, doSort, maxLines);
}
void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, OutputArray h_voices_)
void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, OutputArray h_votes_)
{
if (d_lines.empty())
{
h_lines_.release();
if (h_voices_.needed())
h_voices_.release();
if (h_votes_.needed())
h_votes_.release();
return;
}
@ -132,12 +132,12 @@ void cv::gpu::HoughLinesDownload(const GpuMat& d_lines, OutputArray h_lines_, Ou
cv::Mat h_lines = h_lines_.getMat();
d_lines.row(0).download(h_lines);
if (h_voices_.needed())
if (h_votes_.needed())
{
h_voices_.create(1, d_lines.cols, CV_32SC1);
cv::Mat h_voices = h_voices_.getMat();
cv::gpu::GpuMat d_voices(1, d_lines.cols, CV_32SC1, const_cast<int*>(d_lines.ptr<int>(1)));
d_voices.download(h_voices);
h_votes_.create(1, d_lines.cols, CV_32SC1);
cv::Mat h_votes = h_votes_.getMat();
cv::gpu::GpuMat d_votes(1, d_lines.cols, CV_32SC1, const_cast<int*>(d_lines.ptr<int>(1)));
d_votes.download(h_votes);
}
}

Loading…
Cancel
Save