对 cuda 中的循环全局缓冲区的原子操作

Question

我正在实现一个循环全局内存，以使所有线程同时将数据读/写到同一缓冲区。它是CPU中非常简单的生产者/消费者算法。但是我发现我的 cuda 代码有问题。循环缓冲区定义如下：

#define BLOCK_NUM 1024
#define THREAD_NUM 64
#define BUFFER_SIZE BLOCK_NUM*THREAD_NUM*10
struct Stack {
    bool bDirty[BUFFER_SIZE];
    unsigned int index;
    unsigned int iStackSize;
}

读取设备实现为

__device__ void read(Stack *pStack) {
    unsigned int index = atomicDec(&pStack->index, BUFFER_SIZE-1);
    if(- -index >= BUFFER_SIZE)
        index = BUFFER_SIZE - 1;
    // check
    if(pStack->bDirty[index] == false) {
        printf(“no datan”);
        return;
    }
    //set read flag
    pStack->bDirty[index] = false;
    atomicSub(&pStack->iStackSize, 1);
}

写入设备功能为：

__device__ void write(Stack *pStack) {
    unsigned int index = atomicInc(&pStack->index, BUFFER_SIZE - 1);
    //check
    if(pStack->bDirty[index] == true) {
        printf(“why dirtyn”);
        return;
    }
    pStack->bDirty[index] = true;
    atomicAdd(&pStack->iStackSize, 1);
}

为了以更健壮的方式测试读/写函数，我编写了以下内核：

__global__ void kernelWrite(Stack *pStack) {
    if(threadIdx.x != 0) //make write less than thread number for testing purpose
        write(pStack);
}
__global__ void kernelRead(Stack *pStack) {
    read(pStack);
    __syncthreads();
    if(threadIdx.x % 3 != 0) // make write less than read
        write(pStack);
    __syncthreads();
}

在主函数中，我使用了一个死循环来测试读/写是否是原子的。

int main() {
    Stack *pHostStack = (Stack*)malloc(sizeof(Stack));
    Stack *pStack;
    cudaMalloc(&pStack, sizeof(Stack));
    cudaMemset(pStack, 0, sizeof(Stack));
    while(true) { //dead loop
        kernelWrite<<<BLOCK_NUM, THREAD_NUM>>>(pStack);
        cudaDeviceSynchonize();
        cudaMemcpy(pHostStack, pStack, sizeof(Stack), cudaMemcpyDeviceToHost);
        while(pHost->iStackSize >= BLOCK_NUM*THREAD_NUM) {
             kernelRead<<<BLOCK_NUM, THREAD_NUM>>>(pStack);
                   cudaDeviceSynchonize();
                   cudaMemcpy(pHostStack, pStack, sizeof(Stack), cudaMemcpyDeviceToHost);
         }
    return 0;
}

当我执行上面的代码时，我收到错误msg"为什么脏"和"没有数据"。读/写逻辑有什么问题？

顺便说一下，我没有将线程 ID 映射到线性缓冲区地址，因为在我的应用程序中可能只有 10% 的线程写入缓冲区，它是不可预测/随机的。

Answer 1

关键问题是原子操作不是真正的原子操作，因为读取和写入同一个缓冲区。奇怪的是，当总线程数小于 4096 时，不会显示任何错误消息。