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

Question

我的问题如下:我有一张图像，其中我使用 GPU 检测到一些兴趣点。就处理而言，该检测是一项重量级测试，但平均而言，只有约二十五分之一的点通过测试。该算法的最后阶段是建立点列表。在 CPU 上，这将实现为:

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

在 GPU 上，我让每个 CUDA block 处理 16x16 像素。问题是我需要做一些特殊的事情才能最终在全局内存中拥有一个统一的点列表。目前，我正在尝试在每个 block 的共享内存中生成一个本地点列表，这些点最终将被写入全局内存。我试图避免将任何内容发送回 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
}