title: 0xC0000005: Access violation reading location 0x0000000000000018 at cudaEventDestroy()
body: <p>When I am running a CUDA program, I met &quot;0xC0000005: Access violation reading location 0x0000000000000018.&quot; at cudaEventDestroy().</p>
  <p><strong>It is confusing because this exception sometimes appears and sometimes not</strong> (if not, the program runs normally without any error).</p>
  <p><code>start</code> &amp; <code>stop</code> is defined as <code>cudaEvent_t start, stop;</code> to compute time elapsed on GPU.</p>
  <p><a href="https://i.stack.imgur.com/Xvc0i.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/Xvc0i.png" alt="Exception Thrown" /></a></p>
  <p>Sometimes the Exception looks like: (at a different location)</p>
  <p><a href="https://i.stack.imgur.com/9H0yl.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/9H0yl.png" alt="enter image description here" /></a></p>
  <p><a href="https://i.stack.imgur.com/dWaWs.png" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/dWaWs.png" alt="When runs normally" /></a></p>
  <p>Here I attach the whole program code, in which 2 matrices A &amp; B add together to get S on both GPU and CPU, and compare the results.</p>
  <p>I am new to CUDA, I will appreciate it a lot if somebody may explain this to me.</p>
  <p>Note: the original code is provided by book &quot;professional CUDA C programming&quot;.</p>
  <pre><code>#include &lt;cmath&gt;
  #include &lt;iostream&gt;
  #include &quot;common.h&quot;
  #include &quot;cuda_runtime.h&quot;
  #include &quot;device_launch_parameters.h&quot;
  
  
  void initialData(float *ip, const int size)
  {
      for(int i = 0; i &lt; size; i++)
      {
          ip[i] = (float)(rand() &amp; 0xFF) / 10.0f;
      }
  }
  
  void sumMatrixOnHost(float *A, float *B, float *C, const int nx, const int ny)
  {
      float* ia = A;
      float* ib = B;
      float* ic = C;
  
      for (int iy = 0; iy &lt; ny; iy++)
      {
          for (int ix = 0; ix &lt; nx; ix++)
          {
              ic[ix] = ia[ix] + ib[ix];
          }
          ia += nx;
          ib += nx;
          ic += nx;
      }
  }
  
  void checkResult(float *hostRef, float *gpuRef, const int N)
  {
      double epsilon = 1.0E-8;
      bool match = 1;
      for (int i = 0; i &lt; N; i++)
      {
          if (abs(hostRef[i] - gpuRef[i]) &gt; epsilon)
          {
              match = 0;
              std::cout &lt;&lt; &quot;host&quot; &lt;&lt; hostRef[i] &lt;&lt; &quot;gpu&quot; &lt;&lt; gpuRef[i] &lt;&lt; std::endl;
              break;
          }
      }
      if (match)
          std::cout&lt;&lt;&quot;Arrays match.&quot;&lt;&lt;std::endl;
      else
          std::cout&lt;&lt;&quot;Arrays do not match.&quot;&lt;&lt;std::endl;
  }
  
  // grid 2D block 2D
  __global__ void sumMatrixOnGPU2D(float *MatA, float *MatB, float *MatC, int nx, int ny)
  {
      unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
      unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
      unsigned int idx = iy * nx + ix;
      if (ix &lt; nx &amp;&amp; iy &lt; ny)
          MatC[idx] = MatA[idx] + MatB[idx];
  }
  
  int main(int argc, char **argv)
  {
      clock_t startc, stopc;
      double cpu_time_used;
      cudaEvent_t start, stop;
      float gpu_time_used = 0;
      cudaEventCreate(&amp;start);
      cudaEventCreate(&amp;stop);
  
      // set up device
      int dev = 0;//
      cudaDeviceProp deviceProp;//
      CHECK(cudaGetDeviceProperties(&amp;deviceProp, dev));//
      std::cout &lt;&lt; &quot;Using Device &quot; &lt;&lt; dev &lt;&lt; &quot; &quot; &lt;&lt; deviceProp.name &lt;&lt; std::endl;// 0 NVIDIA GeForce RTX 3060 Laptop GPU
      CHECK(cudaSetDevice(dev));//
      cudaSetDevice(dev);//
  
      // set up data size of matrix
      int nx = 1 &lt;&lt; 14;// 16384
      int ny = 1 &lt;&lt; 14;//
      int nxy = nx * ny;//
      int nBytes = nxy * sizeof(float);//
      std::cout &lt;&lt; &quot;Matrix size: nx=&quot; &lt;&lt; nx &lt;&lt; &quot; ny=&quot; &lt;&lt; ny &lt;&lt; std::endl;
  
      // malloc host memory
      startc = clock();
      cudaEventRecord(start, 0);
      float *h_A, *h_B, *h_S, *d_S;//
      h_A = (float *)malloc(nBytes);//
      h_B = (float *)malloc(nBytes);//
      h_S = (float *)malloc(nBytes);//
      d_S = (float *)malloc(nBytes);//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;malloc() elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;malloc() elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // initialize matrix at host side
      startc = clock();
      cudaEventRecord(start, 0);
      initialData(h_A, nxy);//
      initialData(h_B, nxy);//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;matrix initialization elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;matrix initialization elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      startc = clock();
      cudaEventRecord(start, 0);
      memset(h_S, 0, nBytes);//
      memset(d_S, 0, nBytes);//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;memset() elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;memset() elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // add matrix at host side for result checks
      startc = clock();
      cudaEventRecord(start, 0);
      sumMatrixOnHost(h_A, h_B, h_S, nx, ny);//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;sumMatrixOnHost() elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;sumMatrixOnHost() elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // malloc device global memory
      float *d_MatA, *d_MatB, *d_MatC;
      startc = clock();
      cudaEventRecord(start, 0);
      CHECK(cudaMalloc((void **)&amp;d_MatA, nBytes));//
      CHECK(cudaMalloc((void **)&amp;d_MatB, nBytes));//
      CHECK(cudaMalloc((void **)&amp;d_MatC, nBytes));//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;cudaMalloc() elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;cudaMalloc() elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // transfer data from host to device
      startc = clock();
      cudaEventRecord(start, 0);
      CHECK(cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice));//
      CHECK(cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice));//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;cudaMemcpyHostToDevice elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;cudaMemcpyHostToDevice elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
      CHECK(cudaGetLastError());//
  
      // invoke kernel at host side
      int dimx = 32;
      int dimy = 32;
      dim3 block(dimx, dimy);
      dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);
  
      startc = clock();
      cudaEventRecord(start, 0);
      sumMatrixOnGPU2D &lt;&lt; &lt;grid, block &gt;&gt; &gt; (d_MatA, d_MatB, d_MatC, nx, ny);//
      CHECK(cudaDeviceSynchronize());// shoule be called after kernel called?
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;sumMatrixOnGPU2D&lt;&lt;&lt;(&quot; &lt;&lt; grid.x &lt;&lt; grid.y &lt;&lt; &quot;), (&quot; &lt;&lt; block.x &lt;&lt; block.y &lt;&lt; &quot;)&gt;&gt;&gt; elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;sumMatrixOnGPU2D&lt;&lt;&lt;(&quot; &lt;&lt; grid.x &lt;&lt; grid.y &lt;&lt; &quot;), (&quot; &lt;&lt; block.x &lt;&lt; block.y &lt;&lt; &quot;)&gt;&gt;&gt; elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
  
      // check kernel error
      startc = clock();
      cudaEventRecord(start, 0);
      CHECK(cudaGetLastError());//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;cudaGetLastError() elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;cudaGetLastError() elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // copy kernel result back to host side
      startc = clock();
      cudaEventRecord(start, 0);
      CHECK(cudaMemcpy(d_S, d_MatC, nBytes, cudaMemcpyDeviceToHost));//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;cudaMemcpyDeviceToHost elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;cudaMemcpyDeviceToHost elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // check device results
      startc = clock();
      cudaEventRecord(start, 0);
      checkResult(h_S, d_S, nxy);//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;check results elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;check results elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // free device global memory
      startc = clock();
      cudaEventRecord(start, 0);
      CHECK(cudaFree(d_MatA));//
      CHECK(cudaFree(d_MatB));//
      CHECK(cudaFree(d_MatC));//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;cudaFree() elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;cudaFree() elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // free host memory
      startc = clock();
      cudaEventRecord(start, 0);
      free(h_A);//
      free(h_B);//
      free(h_S);//
      free(d_S);//
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&amp;gpu_time_used, start, stop);
      std::cout &lt;&lt; &quot;free() elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
      std::cout &lt;&lt; &quot;free() elapsed &quot; &lt;&lt; 1e-3 * gpu_time_used &lt;&lt; &quot; sec on GPU&quot; &lt;&lt; std::endl;
  
      // reset device
      startc = clock();
      CHECK(cudaDeviceReset());// cuda reset
      stopc = clock(); cpu_time_used = ((double)(stopc - startc)) / CLOCKS_PER_SEC;
      std::cout &lt;&lt; &quot;cudaDeviceReset() elapsed &quot; &lt;&lt; cpu_time_used &lt;&lt; &quot; sec on CPU&quot; &lt;&lt; std::endl;
  
      cudaEventDestroy(start);
      cudaEventDestroy(stop);
  
      return 0;//
  }
  </code></pre>
  <p>And the &quot;common.h&quot; simply contains the #define of CHECK as below.</p>
  <pre><code>#define CHECK(call)                                                                         \
  {                                                                                           \
      const cudaError_t error = call;                                                         \
      if (error != cudaSuccess)                                                               \
      {                                                                                       \
          fprintf(stderr, &quot;Error: %s:%d, &quot;, __FILE__, __LINE__);                              \
          fprintf(stderr, &quot;code: %d, reason: %s\n&quot;, error,                                    \
              cudaGetErrorString(error));                                                     \
          exit(1);                                                                            \
      }                                                                                       \
  }                                                                                           \
  </code></pre>