与 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' dead 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

我建议使用前缀和来实现该部分以增加并行度。为此，您需要使用共享数组。基本上，前缀和会将数组 (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;
    }
}