title: Accessing global memory in CUDA is slow
body: <p>I have a CUDA kernel doing some computation on a local variable (in register), and after it gets computed, its value gets written into a global array <code>p</code>:</p>
  
  <pre><code>__global__ void dd( float* p, int dimX, int dimY, int dimZ  )
  {
      int 
          i = blockIdx.x*blockDim.x + threadIdx.x,
          j = blockIdx.y*blockDim.y + threadIdx.y,
          k = blockIdx.z*blockDim.z + threadIdx.z,
          idx = j*dimX*dimY + j*dimX +i;   
  
      if (i &gt;= dimX || j &gt;= dimY || k &gt;= dimZ)
      {
          return;
      }   
  
      float val = 0;
  
      val = SomeComputationOnVal();
  
       p[idx ]=  val;
      __syncthreads();    
  
  } 
  </code></pre>
  
  <p>Unfortunately, this function executes very slow.</p>
  
  <p>However, it runs very fast if I do this:</p>
  
  <pre><code>  __global__ void dd(   float* p, int dimX, int dimY, int dimZ  )
      {
          int 
              i = blockIdx.x*blockDim.x + threadIdx.x,
              j = blockIdx.y*blockDim.y + threadIdx.y,
              k = blockIdx.z*blockDim.z + threadIdx.z,
              idx = j*dimX*dimY + j*dimX +i;   
  
          if (i &gt;= dimX || j &gt;= dimY || k &gt;= dimZ)
          {
              return;
          }   
  
          float val = 0;
  
          //val = SomeComputationOnVal();
  
           p[idx ]=  val;
          __syncthreads();    
  
      } 
  </code></pre>
  
  <p>It also runs very fast if I do this:</p>
  
  <pre><code>__global__ void dd( float* p, int dimX, int dimY, int dimZ  )
  {
      int 
          i = blockIdx.x*blockDim.x + threadIdx.x,
          j = blockIdx.y*blockDim.y + threadIdx.y,
          k = blockIdx.z*blockDim.z + threadIdx.z,
          idx = j*dimX*dimY + j*dimX +i;   
  
      if (i &gt;= dimX || j &gt;= dimY || k &gt;= dimZ)
      {
          return;
      }   
  
      float val = 0;
  
      val = SomeComputationOnVal();
  
    //   p[idx ]=  val;
      __syncthreads();    
  
  } 
  </code></pre>
  
  <p>So I am confused, and have no idea how to solve this problem. I have used NSight step in, and did not find access violations. </p>
  
  <p>Here is how I launch the kernel (dimX:924; dimY: 16: dimZ: 1120):</p>
  
  <pre><code>dim3 
        blockSize(8,16,2),
        gridSize(dimX/blockSize.x+1,dimY/blockSize.y, dimZ/blockSize.z);
  float* dev_p;       cudaMalloc((void**)&amp;dev_p, dimX*dimY*dimZ*sizeof(float));
  
  dd&lt;&lt;&lt;gridSize, blockSize&gt;&gt;&gt;(     dev_p,dimX,dimY,dimZ);
  </code></pre>
  
  <p>Could anyone please gives some pointers? Because it does not make much sense to me. All computation of val is fast, and the final step is to move <code>val</code> into <code>p</code>. <code>p</code> never gets involved in the computation, and it only shows up once. So why is it so slow?</p>
  
  <p>The computations are basically a loop over a 512  X 512 matrix. It is pretty fair amount of computation I'd say.</p>