CUDAのビット配列

Question

CUDAでEratosthenesのSieveを実装しており、非常に奇妙な出力を持っています。 unsigned char *をデータ構造体として使用し、以下のマクロを使用してビットを操作します。私はここ

size_t p=3; 
size_t primeTill = 30; 

while(p*p<=primeTill) 
{ 
    if(ISBITSET(h_a, p) == 1){ 
     int dimA = 30; 
     int numBlocks = 1; 
     int numThreadsPerBlock = dimA; 
     dim3 dimGrid(numBlocks); 
     dim3 dimBlock(numThreadsPerBlock); 
     cudaMemcpy(d_a, h_a, memSize, cudaMemcpyHostToDevice);   
     cudaThreadSynchronize();  
     reverseArrayBlock<<< dimGrid, dimBlock >>>(d_a, primeTill, p); 
     cudaThreadSynchronize();  
     cudaMemcpy(h_a, d_a, memSize, cudaMemcpyDeviceToHost); 
     cudaThreadSynchronize();  
     printf("This is after removing multiples of %d\n", p); 
     //Loop 
     for(size_t i = 0; i < primeTill +1; i++) 
     { 
      printf("Bit %d is %d\n", i, ISBITSET(h_a, i)); 
     } 
    }   
    p++; 
} 

私のカーネルを呼び出すところ

#define ISBITSET(x,i) ((x[i>>3] & (1<<(i&7)))!=0) 
#define SETBIT(x,i) x[i>>3]|=(1<<(i&7)); 
#define CLEARBIT(x,i) x[i>>3]&=(1<<(i&7))^0xFF; 

私はそれ以外の場合は、ここで= 0 だ、それは素数だ示すためにビットをセットである私のカーネル

__global__ void reverseArrayBlock(unsigned char *d_out, int size, size_t p) 
{ 
int id = blockIdx.x*blockDim.x + threadIdx.x; 
int r = id*p; 
if(id >= p && r <= size) 
{ 
    while(ISBITSET(d_out, r) == 1){ 
     CLEARBIT(d_out, r); 
    } 

    // if(r == 9) 
    // { 
    // /* code */ 
    // CLEARBIT(d_out, 9); 
    // } 

} 

}です 出力は ,2,3,5,7,11,13,17,19,23,29,です。出力はです。カーネルコードを調べると、これらの行のコメントを外すと正しい答えが得られます。つまり、私のループや私のチェックに間違いがないこと！

Answer 1

複数のスレッドが同時にグローバルメモリ内の同じワード（char）にアクセスしているため、結果が壊れてしまいます。

これを防ぐために原子操作を使用できますが、より良い解決策は、アルゴリズムを変更することです：すべてのスレッドが2,3,4,5の倍数を篩に通すのではなく、すべてのスレッドが[0..7]、[8..15]、...のように、すべての範囲の長さは8ビットの倍数であり、衝突は発生しません。

Answer 2

マクロを最初から始める方法に置き換えることをお勧めします。 __host__と__device__の前にあるメソッドを使用して、必要に応じてcppおよびcu固有のバージョンを生成できます。それは、予期せぬことをするプリプロセッサの可能性を根絶するでしょう。

誤った出力を引き起こしている特定のコードブランチをデバッグし、各ステージが順番に正しいことを確認するだけで、問題が見つかります。