Commit 29cdddd8 authored by Vladislav Vinogradov's avatar Vladislav Vinogradov

fixed bug in buildPointList

parent 3ae2244b
...@@ -64,23 +64,22 @@ namespace cv { namespace gpu { namespace device ...@@ -64,23 +64,22 @@ namespace cv { namespace gpu { namespace device
const int x = blockIdx.x * blockDim.x * PIXELS_PER_THREAD + threadIdx.x; const int x = blockIdx.x * blockDim.x * PIXELS_PER_THREAD + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y; const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (y >= src.rows)
return;
if (threadIdx.x == 0) if (threadIdx.x == 0)
s_qsize[threadIdx.y] = 0; s_qsize[threadIdx.y] = 0;
__syncthreads(); __syncthreads();
// fill the queue if (y < src.rows)
const uchar* srcRow = src.ptr(y);
for (int i = 0, xx = x; i < PIXELS_PER_THREAD && xx < src.cols; ++i, xx += blockDim.x)
{ {
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; if (srcRow[xx])
const int qidx = Emulation::smem::atomicAdd(&s_qsize[threadIdx.y], 1); {
s_queues[threadIdx.y][qidx] = val; const unsigned int val = (y << 16) | xx;
const int qidx = Emulation::smem::atomicAdd(&s_qsize[threadIdx.y], 1);
s_queues[threadIdx.y][qidx] = val;
}
} }
} }
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment