title: CUDA counting, reduction and thread warps
body: <p>I'm trying to create a cuda program that counts the number of true values (defined by non-zero values) in a long vector through a reduction algorithm.  I'm getting funny results.  I get either 0 or (ceil(N/threadsPerBlock)*threadsPerBlock), neither is correct.</p>
  
  <pre><code>__global__ void count_reduce_logical(int *  l, int * cntl, int N){
      // suml is assumed to blockDim.x long and hold the partial counts
      __shared__ int cache[threadsPerBlock];
      int cidx = threadIdx.x;
      int tid = threadIdx.x + blockIdx.x*blockDim.x;
  
      int cnt_tmp=0;
      while(tid&lt;N){
          if(l[tid]!=0)
                  cnt_tmp++;
          tid+=blockDim.x*gridDim.x;
      }
      cache[cidx]=cnt_tmp;
      __syncthreads();
      //reduce
      int k =blockDim.x/2;
      while(k!=0){
          if(threadIdx.x&lt;k)
              cache[cidx] += cache[cidx];
          __syncthreads();
          k/=2;
      }
      if(cidx==0)
          cntl[blockIdx.x] = cache[0];
  }
  </code></pre>
  
  <p>The host code then collects the cntl results and finishes summation.  This is going to be part of a larger project where the data is already on the GPU, so it makes sense to do the computations there, if they work correctly.</p>