如何将元素（cv::Point）添加到共享数组中 - CUDA

Question

我是 Cuda 技术的新手。我需要帮助 CUDA 在二进制（单色）图像中找到只有像素，其值为白色（255）。然后需要像素来对输出数组进行排序。我的解决方案基于关键部分。但是，它给出的结果不正确。

//----- call kernel: -----
{
    const dim3 block(16,16);
    const dim3 grid(divUp(_binImg.cols, block.x), divUp(_binImg.rows, block.y));
    // others allocations, declarations ...
    cudaCalcWhitePixels<<<grid, block>>>(_binImg, _index, _pointsX, _pointsY);
}

__device__ int lock = 0;
__global__ void cudaCalcWhitePixels(cv::gpu::PtrStepSzb _binImg, int *_index, int *_pointsX, int *_pointsY)
{
    extern int lock;
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockIdx.y * blockDim.y + threadIdx.y;

    __syncthreads();

    if(x < _binImg.cols && y < _binImg.rows)
    {
        if(_binImg.ptr(y)[x] == 255)
        {
            do{} while(atomicCAS(&lock, 0, 1) != 0)

            //----- critical section ------

            _pointsX[*_index] = x;
            _pointsY[*_index] = y;
            (*_index)++;
            lock = 0;

            //----- end CS ------
        }
    }
}

在我看来，关键部分工作不正常。图像中的白色像素大约占 1%。

你能告诉我吗？谢谢你，祝你有美好的一天:)

编辑： 解决方案：

__global__ void cudaCalcWhitePixels(cv::gpu::PtrStepSzb _binImg, int *_index, int *_pointsX, int *_pointsY)
{
    int myIndex = 0;
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockIdx.y * blockDim.y + threadIdx.y;

    __syncthreads();

    if(x < _binImg.cols && y < _binImg.rows)
    {
        if(_binImg.ptr(y)[x] == 255)
        {
            //----- critical section ------

            myIndex = atomicAdd(_index, 1);
            _pointsX[myIndex] = x;
            _pointsY[myIndex] = y;

            //----- end CS ------
        }
    }
}

Answer 1

来自以下 URL 的这段代码可以帮助您了解如何使用 atomicCAS() 创建临界区。

https://github.com/ArchaeaSoftware/cudahandbook/blob/master/memory/spinlockReduction.cu

class cudaSpinlock {
public:
    cudaSpinlock( int *p );
    void acquire();
    void release();
private:
    int *m_p;
};

inline __device__
cudaSpinlock::cudaSpinlock( int *p )
{
    m_p = p;
}

inline __device__ void
cudaSpinlock::acquire( )
{
    while ( atomicCAS( m_p, 0, 1 ) );
}

inline __device__ void
cudaSpinlock::release( )
{
    atomicExch( m_p, 0 );
}

由于(*_index)++; 是您在 CS 中执行的唯一原子操作，您可以考虑改用atomicAdd()。

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd

另一方面，您可以尝试使用thrust::copy_if() 来简化编码。