Bitmap in CUDA

Question

I introduce a sieve of eratosthenes in CUDA and get a very strange result. I use unsigned char * as a data structure and using the following macros to control bits.

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

I set the bit to denote this prime number, otherwise it = 0. Here I call my kernel

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

Here is my core

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

}

} The output should be: 2, 3, 5, 7, 11, 13, 17, 19, 23, 29 while my result: 2, 3, 5, 9, 7, 11, 13, 17, 19, 23 , 29

If you look at the kernel code, if I uncomment these lines, I will get the correct answer, which means that there is nothing wrong with my loops or my check!

Answer 1

(char) , , .

, , : , 2, 3, 4, 5,... , [ 0..7], [8..15],... 8 .

Answer 2

. , __host__ __device__, cpp cu. - .

, , , , .