From 29cdddd84540990b1140987facae68fac9b55eb6 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Thu, 30 Aug 2012 09:48:16 +0400 Subject: [PATCH] fixed bug in buildPointList --- modules/gpu/src/cuda/hough.cu | 21 ++++++++++----------- 1 file changed, 10 insertions(+), 11 deletions(-) diff --git a/modules/gpu/src/cuda/hough.cu b/modules/gpu/src/cuda/hough.cu index 580b9de..9ee7621 100644 --- a/modules/gpu/src/cuda/hough.cu +++ b/modules/gpu/src/cuda/hough.cu @@ -64,23 +64,22 @@ namespace cv { namespace gpu { namespace device const int x = blockIdx.x * blockDim.x * PIXELS_PER_THREAD + threadIdx.x; const int y = blockIdx.y * blockDim.y + threadIdx.y; - if (y >= src.rows) - return; - if (threadIdx.x == 0) s_qsize[threadIdx.y] = 0; - __syncthreads(); - // fill the queue - const uchar* srcRow = src.ptr(y); - for (int i = 0, xx = x; i < PIXELS_PER_THREAD && xx < src.cols; ++i, xx += blockDim.x) + if (y < src.rows) { - if (srcRow[xx]) + // fill the queue + const uchar* srcRow = src.ptr(y); + for (int i = 0, xx = x; i < PIXELS_PER_THREAD && xx < src.cols; ++i, xx += blockDim.x) { - const unsigned int val = (y << 16) | xx; - const int qidx = Emulation::smem::atomicAdd(&s_qsize[threadIdx.y], 1); - s_queues[threadIdx.y][qidx] = val; + if (srcRow[xx]) + { + const unsigned int val = (y << 16) | xx; + const int qidx = Emulation::smem::atomicAdd(&s_qsize[threadIdx.y], 1); + s_queues[threadIdx.y][qidx] = val; + } } } -- 2.7.4