与 CUDA 共享内存互斥 - 添加到项目列表

Question

我的问题如下：我有一张图像，其中我使用 GPU 检测了一些兴趣点。该检测在处理方面是一个重量级的测试，但平均只有大约 25 个点中的 1 个通过测试。该算法的最后阶段是建立一个点列表。在 CPU 上，这将实现为：

forall pixels x,y
{
    if(test_this_pixel(x,y))
        vector_of_coordinates.push_back(Vec2(x,y));
}

在 GPU 上，我让每个 CUDA 块处理 16x16 像素。问题是我需要做一些特别的事情来最终在全局内存中有一个统一的点列表。目前我正在尝试在每个块的共享内存中生成一个本地点列表，最终将写入全局内存。我试图避免将任何东西发送回 CPU，因为在此之后还有更多的 CUDA 阶段。

我期待我可以使用原子操作来实现共享内存上的

push_back

功能。但是我无法使它正常工作。有两个问题。第一个烦人的问题是我经常遇到以下编译器崩溃：

nvcc error : 'ptxas' died with status 0xC0000005 (ACCESS_VIOLATION)

使用原子操作时。我是否可以编译一些东西是命中注定的。有谁知道是什么原因造成的？

以下内核将重现错误：

__global__ void gpu_kernel(int w, int h, RtmPoint *pPoints, int *pCounts)
{
    __shared__ unsigned int test;
    atomicInc(&test, 1000);
}

其次，我的代码包含共享内存上的互斥锁挂起 GPU，我不明白为什么：

__device__ void lock(unsigned int *pmutex)
{
    while(atomicCAS(pmutex, 0, 1) != 0);
}

__device__ void unlock(unsigned int *pmutex)
{
    atomicExch(pmutex, 0);
}

__global__ void gpu_kernel_non_max_suppress(int w, int h, RtmPoint *pPoints, int *pCounts)
{
    __shared__ RtmPoint localPoints[64];
    __shared__ int localCount;
    __shared__ unsigned int mutex;

    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    int threadid = threadIdx.y * blockDim.x + threadIdx.x;
    int blockid = blockIdx.y * gridDim.x + blockIdx.x;

    if(threadid==0)
    {
        localCount = 0;
        mutex = 0;
    }

    __syncthreads();

    if(x<w && y<h)
    {
        if(some_test_on_pixel(x,y))
        {
            RtmPoint point;
            point.x = x;
            point.y = y;

            // this is a local push_back operation
            lock(&mutex);
            if(localCount<64) // we should never get >64 points per block
                localPoints[localCount++] = point;
            unlock(&mutex);
        }
    }

    __syncthreads();

    if(threadid==0)
        pCounts[blockid] = localCount;
    if(threadid<localCount)
        pPoints[blockid * 64 + threadid] = localPoints[threadid];
}

在this site的示例代码中，作者设法成功地在共享内存上使用原子操作，所以我很困惑为什么我的案例不起作用。如果我注释掉锁定和解锁行，代码运行正常，但显然错误地添加到列表中。

如果有一些关于为什么会发生这个问题的建议，以及是否有更好的解决方案来实现这个目标，我将不胜感激，因为无论如何我都担心使用原子操作或互斥锁的性能问题。

Answer 1

我建议使用 prefix-sum 来实现该部分以增加并行性。为此，您需要使用共享数组。基本上 prefix-sum 会将数组 (1,1,0,1) 转换为 (0,1,2,2,3)，即，将计算一个就地运行的独占总和，以便您获得每个线程写索引。

__shared__ uint8_t vector[NUMTHREADS];

....

bool emit  = (x<w && y<h);
     emit  = emit && some_test_on_pixel(x,y);
__syncthreads();
scan(emit, vector);
if (emit) {
     pPoints[blockid * 64 + vector[TID]] = point;
}

前缀和示例：

    template <typename T>
__device__ uint32 scan(T mark, T *output) {
#define GET_OUT (pout?output:values)
#define GET_INP (pin?output:values)
  __shared__ T values[numWorkers];
  int pout=0, pin=1;
  int tid = threadIdx.x;

  values[tid] = mark;

  syncthreads();

  for( int offset=1; offset < numWorkers; offset *= 2) {
    pout = 1 - pout; pin = 1 - pout;
    syncthreads();
    if ( tid >= offset) {
      GET_OUT[tid] = (GET_INP[tid-offset]) +( GET_INP[tid]);
    }
    else {
      GET_OUT[tid] = GET_INP[tid];
    }
    syncthreads();
  }

  if(!pout)
    output[tid] =values[tid];

  __syncthreads();

  return output[numWorkers-1];

#undef GET_OUT
#undef GET_INP
}

Answer 2

根据此处的建议，我将最后使用的代码包含在内。它使用 16x16 像素块。请注意，我现在正在将数据写到一个全局数组中而不将其分解。我使用全局

atomicAdd

函数为每组结果计算基址。因为每个块只调用一次，所以我没有发现太多的减慢，同时我通过这样做获得了更多的便利。我还避免为

prefix_sum

的输入和输出共享缓冲区。

GlobalCount

在内核调用之前设置为零。

#define BLOCK_THREADS 256

__device__ int prefixsum(int threadid, int data)
{
    __shared__ int temp[BLOCK_THREADS*2];

    int pout = 0;
    int pin = 1;

    if(threadid==BLOCK_THREADS-1)
        temp[0] = 0;
    else
        temp[threadid+1] = data;

    __syncthreads();

    for(int offset = 1; offset<BLOCK_THREADS; offset<<=1)
    {
        pout = 1 - pout;
        pin = 1 - pin;

        if(threadid >= offset)
            temp[pout * BLOCK_THREADS + threadid] = temp[pin * BLOCK_THREADS + threadid] + temp[pin * BLOCK_THREADS + threadid - offset];
        else
            temp[pout * BLOCK_THREADS + threadid] = temp[pin * BLOCK_THREADS + threadid];

        __syncthreads();
    }

    return temp[pout * BLOCK_THREADS + threadid];
}

__global__ void gpu_kernel(int w, int h, RtmPoint *pPoints, int *pGlobalCount)
{
    __shared__ int write_base;

    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    int threadid = threadIdx.y * blockDim.x + threadIdx.x;
    int valid = 0;

    if(x<w && y<h)
    {
        if(test_pixel(x,y))
        {
            valid = 1;
        }
    }

    int index = prefixsum(threadid, valid);

    if(threadid==BLOCK_THREADS-1)
    {
        int total = index + valid;
        if(total>64)
            total = 64; // global output buffer is limited to 64 points per block
        write_base = atomicAdd(pGlobalCount, total); // get a location to write them out
    }

    __syncthreads(); // ensure write_base is valid for all threads

    if(valid)
    {
        RtmPoint point;
        point.x = x;
        point.y = y;
        if(index<64)
            pPoints[write_base + index] = point;
    }
}