Bit-Array in CUDA

Question

Implementiere Sieb von Eratosthenes in CUDA und habe eine sehr seltsame Ausgabe. Verwenden Sie unsigned char * als Datenstruktur und verwenden Sie die folgenden Makros, um die Bits zu manipulieren.

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

stelle ich das bisschen, es ist eine Primzahl zu bezeichnen, sonst ist es = 0. Hier ist, wo ich meine Kernel nennen

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

Hier ist mein Kernel

__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); 
    // } 

} 

} Der Ausgang sollte sein: 2, 3, 5, 7, 11, 13, 17, 19, 23, 29 während meine Ausgabe ist: 2, 3, 5, 9, 7, 11, 13, 17, 19, 23, 29

Wenn Sie sich den Kernel-Code ansehen, wenn ich diese Zeilen auskommentiere, bekomme ich die richtige Antwort, was bedeutet dass mit meinen Schleifen oder meiner Überprüfung nichts falsch ist!

Answer 1

Mehrere Threads greifen gleichzeitig auf das gleiche Wort (char) im globalen Speicher zu und das geschriebene Ergebnis wird dadurch beschädigt.

Sie könnten atomare Operationen verwenden, um dies zu verhindern, aber die bessere Lösung wäre, Ihren Algorithmus zu ändern: Anstatt jeden Thread mehrere, 2, 3, 4, 5 aussortieren zu lassen ... lässt jeder Thread einen Bereich wie [0..7], [8..15], ... so dass die Länge jedes Bereichs ein Vielfaches von 8 Bits ist und keine Kollisionen auftreten.

Answer 2

Ich würde vorschlagen, die Makros mit Methoden zu ersetzen, mit zu beginnen. Sie können Methoden verwenden, die von __host__ und __device__ eingeleitet werden, um cpp- und cu-spezifische Versionen zu generieren, falls erforderlich. Das wird die Möglichkeit ausschließen, dass der Vorprozessor etwas Unerwartetes tut.

Jetzt debuggen Sie nur den bestimmten Code-Zweig, der die falsche Ausgabe verursacht, und überprüfen Sie, ob jede Stufe der Reihe nach korrekt ist, und Sie werden das Problem finden.