title: Bit array in CUDA
body: <p>Am implementing Sieve of Eratosthenes in CUDA and am having a very weird output. Am using unsigned char* as the data structure and using the following macros to manipulate the bits.</p>
  
  <pre><code>#define ISBITSET(x,i) ((x[i&gt;&gt;3] &amp; (1&lt;&lt;(i&amp;7)))!=0)
  #define SETBIT(x,i) x[i&gt;&gt;3]|=(1&lt;&lt;(i&amp;7));
  #define CLEARBIT(x,i) x[i&gt;&gt;3]&amp;=(1&lt;&lt;(i&amp;7))^0xFF;
  </code></pre>
  
  <p>I set the bit to denote it's a prime number, otherwise it's = 0. 
  Here is where i call my kernel</p>
  
  <pre><code>size_t p=3;
  size_t primeTill = 30;
  
  while(p*p&lt;=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&lt;&lt;&lt; dimGrid, dimBlock &gt;&gt;&gt;( 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 &lt; primeTill +1; i++)
          {
              printf("Bit %d is %d\n", i, ISBITSET(h_a, i));
          }
      }           
      p++;
  }
  </code></pre>
  
  <p>Here is my kernel</p>
  
  <pre><code>__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 &gt;= p &amp;&amp; r &lt;= size )
  {
      while(ISBITSET(d_out, r ) == 1 ){
          CLEARBIT(d_out, r);
      }
  
      // if(r == 9)
      // {
      //  /* code */
      //  CLEARBIT(d_out, 9);
      // }
  
  }
  </code></pre>
  
  <p>}
  The output should be:
  2, 3, 5, 7, 11, 13, 17, 19, 23, 29
  while my output is:
  2, 3, 5, 9, 7, 11, 13, 17, 19, 23, 29</p>
  
  <p>If you take a look at the kernel code, if i uncomment those lines i will get the correct answer, which means that there is nothing wrong with my loops or my checking!</p>