キュー内の循環グローバルバッファに対するアトミック操作

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 data\n”); 
     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 dirty\n”); 
     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(); 
} 

main関数では、読み取り/書き込みがアトミックであるかどうかをテストするためにデッドループを使用しました。

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; 
} 

上記のコードを実行すると、エラーメッセージ「なぜ汚い」と「データなし」が表示されます。読み取り/書き込みロジックに何が間違っていますか？

ちなみに私のアプリケーションでは10％のスレッドしかバッファに書き込めないのでスレッドIDをリニアバッファアドレスにマップしません。これは予測不能/ランダムです。

Answer 1

重要な問題は、同じバッファに読み書きするため、アトミックな操作が実際の原子ではないことです。奇妙なことは、総スレッド数が4096未満の場合、エラーメッセージは表示されないということです。