title: How to calculate histogram using shared memory
body: <p>I want to calculate a histogram  using shared memory, and I write the code like below,but I found the result of cuda and cpu is different,and the result of cuda is blockDim.x times of the cpu, why?
  And how can I fix it? 
  thank you!</p>
  
  <pre><code>__global__ void CalHistKernel(int*imgData,int*bins,int datasize)
  {
  __shared__ int _bins[3];
  
  int tx=threadIdx.x;
  int idx=blockIdx.x*blockDim.x+threadIdx.x;//blockDim.y=1
  if(tx&lt;3)
  {
      _bins[tx]=0;       
  }
  __syncthreads();
  
  if(idx&lt;datasize)
  {       
      atomicAdd((int*)&amp;_bins[imgData[idx]],1);     
  
  }
  __syncthreads();
  for(int i=0;i&lt;3;i++)   
      atomicAdd((int*)&amp;bins[i],_bins[i]);
  }  
  </code></pre>