title: CUDA Address Out of Bounds
body: <p>I have been playing with a simple CUDA program that just zeros out global memory.  Below is the device code as well as the host code:</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  
  __global__ void kernel(float *data, int width) {
      int x = blockDim.x * blockIdx.x + threadIdx.x;
      int y = blockDim.y * blockIdx.y + threadIdx.y;
  
      if (x &gt; (width-1)) {
          printf("x = %d\n", x);
          printf("blockDim.x = %d\n", blockDim.x);
          printf("blockIdx.x = %d\n", blockIdx.x);
          printf("threadIdx.x = %d\n", threadIdx.x);
      }   
  
      if (y &gt; (width-1)) {
          printf("y = %d\n", y);
          printf("blockDim.y = %d\n", blockDim.y);
          printf("blockIdx.y = %d\n", blockIdx.y);
          printf("threadIdx.y = %d\n", threadIdx.y);
      }   
  
      data[y * width + x] = 0.0;
  }
  
  int main(void) {
      const int MATRIX_SIZE = 256;
      float *data, *dataGPU;
      int sizeOfMem;
      int x = MATRIX_SIZE;
      int y = MATRIX_SIZE;
  
      cudaDeviceReset();
      cudaDeviceSynchronize();
  
      sizeOfMem = sizeof(float) * x * y;
  
      data = (float *)malloc(sizeOfMem);
      cudaMalloc((void **)&amp;dataGPU, sizeOfMem);
  
      cudaMemcpy(dataGPU, data, sizeOfMem, cudaMemcpyHostToDevice);
  
      //int threads = 256;
      //int blocks = ((x * y) + threads - 1) / threads;
  
      dim3 threads(16, 16);
      dim3 blocks(x / 16, y / 16);
  
      kernel&lt;&lt;&lt;blocks, threads&gt;&gt;&gt;(dataGPU, MATRIX_SIZE);
      cudaThreadSynchronize();
  
      cudaMemcpy(data, dataGPU, sizeOfMem, cudaMemcpyDeviceToHost);
  
      cudaFree(dataGPU);
  
      free(data);
  
      return 0;
  }
  </code></pre>
  
  <p>I continue to receive address out of bounds error messages when running my code with cuda-memcheck.  But this is only if the dimensions of the matrix I have created are 128 or greater.  If I have dimensions that are less than 128, the errors are less frequent (I almost never receive errors).  You may notice that I included print statements in my kernel function.  These statements are only printed when I receive the error messages, because x and y should never be greater than width-1, or in this case 255.  This statement is true, if I have done my math correctly, which I believe I have.  Below is an error message I received from cuda-memcheck:</p>
  
  <pre><code>  ========= CUDA-MEMCHECK
    ========= Invalid __global__ write of size 4
    =========     at 0x00000298 in kernel(float*, int)
    =========     by thread (3,10,0) in block (15,1,0)
    =========     Address 0x2300da6bcc is out of bounds
    =========     Saved host backtrace up to driver entry point at kernel launch time
    =========     Host Frame:/usr/lib64/nvidia/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x472225]
    =========     Host Frame:./test_reg_memory [0x16c41]
    =========     Host Frame:./test_reg_memory [0x31453]
    =========     Host Frame:./test_reg_memory [0x276d]
    =========     Host Frame:./test_reg_memory [0x24f0]
    =========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x21b15]
    =========     Host Frame:./test_reg_memory [0x25cd]
    =========
    y = 2074
    blockDim.y = 16
    blockIdx.y = 1
    threadIdx.y = 10
  </code></pre>
  
  <p>This output does not make sense to me, because if I do the math,</p>
  
  <pre><code>y = blockDim.y * blockIdx.y + threadIdx.y = 16 * 1 + 10 = 26 (not 2074)
  </code></pre>
  
  <p>I have spent some time looking at CUDA programming forums, and nothing has seemed to helped.  There was one thread that I read indicating that I may have corrupted register memory.  However, the one that began the thread has this problem with a different GPU.  The thread is a bit unrelated, but I included the link anyway.</p>
  
  <p><a href="https://devtalk.nvidia.com/default/topic/498784/memory-corruption-on-a-fermi-class-gpu-error-only-on-fermis-program-works-on-non-fermis-/?offset=6" rel="nofollow">https://devtalk.nvidia.com/default/topic/498784/memory-corruption-on-a-fermi-class-gpu-error-only-on-fermis-program-works-on-non-fermis-/?offset=6</a></p>
  
  <p>Below I have included the nvcc version.</p>
  
  <pre><code> nvcc: NVIDIA (R) Cuda compiler driver
   Copyright (c) 2005-2015 NVIDIA Corporation
   Built on Tue_Aug_11_14:27:32_CDT_2015
   Cuda compilation tools, release 7.5, V7.5.17
  </code></pre>
  
  <p>Also, here is the GPU that I am using.</p>
  
  <pre><code> Device 0: "GeForce GT 640"
   CUDA Driver Version / Runtime Version 8.0 / 7.5
   CUDA Capability Major/Minor version number: 3.0
  </code></pre>
  
  <p>Could anyone with CUDA experience point something out that I may be doing wrong?</p>