Битовый массив в CUDA

Question

Я внедряю Sieve of Eratosthenes в CUDA и получаю очень странный вывод. Я использую unsigned char * в качестве структуры данных и использую следующие макросы для управления битами.

#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. Вот где я называю свое ядро ​​

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

Вот мое ядро ​​

__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 пока мой вывод: 2, 3, 5, 9, 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 версии, где это необходимо.Это исключит возможность того, что препроцессор сделает что-то неожиданное.

Теперь просто отладьте конкретную ветвь кода, которая вызывает неправильный вывод, проверяя, что каждый этап корректен по очереди, и вы найдете проблему.