title: Cuda GPUassert: an illegal memory access was encountered
body: <p>I was trying to make a game program using __<strong>device</strong> __ variables instead of declaring it dynamically using cudaMalloc, but it keeps telling me that GPUassert: illegal memory access was encountered at the third last line where the cudaDeviceSynchronization() is called. I have tried the version using cudaMalloc and it worked out fine.</p>
  <pre><code>#include &quot;cuda_runtime.h&quot;
  #include &quot;device_launch_parameters.h&quot;
  
  #include &lt;iostream&gt;
  #include &lt;cmath&gt;
  #include &lt;stdio.h&gt;
  #include &lt;stdlib.h&gt;
  #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
  inline void gpuAssert(cudaError_t code, const char* file, int line, bool abort = true)
  {
      if (code != cudaSuccess)
      {
          fprintf(stderr, &quot;GPUassert: %s %s %d\n&quot;, cudaGetErrorString(code), file, line);
          if (abort) exit(code);
      }
  }
  
  #define M 3
  #define N 3
  #define K 3
  
  using namespace std;
  
  __device__ double A_dev[M * K];
  __device__ double B_dev[K * N];
  __device__ double C_dev[M * N];
  
  __global__ void gemm(double* A, double* B, double* C, int m, int n, int k)
  {
      int x = blockDim.x * blockIdx.x + threadIdx.x;
      int y = blockDim.y * blockIdx.y + threadIdx.y;
  
      int i = x * n + y;
      
      double sum = 0.0;
      for (int j = 0; j &lt; k; j++)
      {
          sum += A[x * k + j] * B[n * j + y];
      }
      C[i] = sum;
      printf(&quot;The value is %f&quot;, C[i]);
  
  }
  
  int main(void)
  {
      double A_h[M * K];
      double B_h[K * N];
      double C_h[M * N];
      
      for (int i = 0; i &lt; M*K; i++)
      {
          A_h[i] = (double)i;
          B_h[i] = (double)i;
          C_h[i] = 0.0;
      }
  
      gpuErrchk(cudaMemcpyToSymbol(A_dev, A_h, M * K * sizeof(double), 0, cudaMemcpyHostToDevice));
      gpuErrchk(cudaMemcpyToSymbol(B_dev, B_h, K * N * sizeof(double), 0, cudaMemcpyHostToDevice));
      gpuErrchk(cudaMemcpyToSymbol(C_dev, C_h, M * N * sizeof(double), 0, cudaMemcpyHostToDevice));
  
      gpuErrchk(cudaPeekAtLastError());
      gpuErrchk(cudaDeviceSynchronize());
  
      dim3 dimGrid(1, 1, 1);
      dim3 dimBlock(3, 3, 1);
      gemm &lt;&lt;&lt;dimGrid, dimBlock &gt;&gt;&gt; (A_dev, B_dev, C_dev, 3, 3, 3);
      gpuErrchk(cudaPeekAtLastError());
      gpuErrchk(cudaDeviceSynchronize());
  
      gpuErrchk(cudaMemcpyFromSymbol(C_h, C_dev, M * N * sizeof(double), 0, cudaMemcpyDeviceToHost));
  
      return 0;
  }
  </code></pre>