title: cudaDeviceSynchronize() error code 77: cudaErrorIllegalAddress
body: <p>Thank you very much for reading my threads. </p>
  
  <p>I am doing CUDA work, but keep getting cudaDeviceSynchronize() error code 77: cudaErrorIllegalAddress, without any idea why. I did the search for both the code and the function, surprisingly , only a few records showed up. Very strange.</p>
  
  <p>I basically sum up all pixels of images. To make my questions have as much reference as it can, I am showing all my CUDA code here:</p>
  
  <pre><code>#include "cuda_runtime.h"
  #include "device_launch_parameters.h"
  #include "thorcalgpu.h"
  #include &lt;stdio.h&gt;
  #include "math.h"
  #include &lt;vector&gt;
  #include &lt;algorithm&gt;
  #include &lt;stdlib.h&gt;
  #include &lt;stdio.h&gt;
  #include &lt;vector&gt;
  #include &lt;numeric&gt;
  #include &lt;iostream&gt;
  
  using namespace std;
  
  float random_float(void)
  {
    return static_cast&lt;float&gt;(rand()) / RAND_MAX;
  }
  
  
  __global__ void reduceSum(unsigned short *input,
                            unsigned long long *per_block_results,
                            const int n)
  {
      extern __shared__ unsigned long long sdata[];
  
      unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
  
      // load input into __shared__ memory
      unsigned short x = 0;
      if(i &lt; n)
      {
          x = input[i];
      }
      sdata[threadIdx.x] = x;
      __syncthreads();
  
      // contiguous range pattern
      for(int offset = blockDim.x / 2; offset &gt; 0; offset &gt;&gt;= 1)
      {
          if(threadIdx.x &lt; offset)
          {
              // add a partial sum upstream to our own
              sdata[threadIdx.x] += sdata[threadIdx.x + offset];
          }
  
          // wait until all threads in the block have
          // updated their partial sums
          __syncthreads();
      }
  
      // thread 0 writes the final result
      if(threadIdx.x == 0)
      {
          per_block_results[blockIdx.x] = sdata[0];
      }
  }
  
  // Helper function for using CUDA to add vectors in parallel.
  //template &lt;class T&gt;
  cudaError_t gpuWrapper(float *mean,  int N,  vector&lt;string&gt; filelist)
  {
      int size = N*N;
      unsigned long long* dev_sum = 0;
      unsigned short* dev_img = 0;
      cudaError_t cudaStatus;
      const int block_size = 512;
      const int num_blocks = (size/block_size) + ((size%block_size) ? 1 : 0);
      int L = filelist.size();
  
      // Choose which GPU to run on, change this on a multi-GPU system.
  
      double totalgpuinittime = 0;
      StartCounter(7);
  
      cudaStatus = cudaSetDevice(0);
      if (cudaStatus != cudaSuccess) 
      {
          fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
          goto Error;
      }
  
      // Allocate GPU buffers for three vectors (two input, one output)    .
      cudaStatus = cudaMalloc((void**)&amp;dev_img, size * sizeof(unsigned short));
      if (cudaStatus != cudaSuccess) 
      {
          fprintf(stderr, "cudaMalloc failed!");
          goto Error;
      }
  
      cudaStatus = cudaMalloc((void**)&amp;dev_sum, num_blocks*sizeof(unsigned long long));
      if (cudaStatus != cudaSuccess) 
      {
          fprintf(stderr, "cudaMalloc failed!");
          goto Error;
      }
  
      totalgpuinittime = GetCounter(7);
  
      unsigned short* img;
      unsigned short* pimg;
      unsigned long long* sum = new unsigned long long[num_blocks];
      unsigned long long* psum = sum;
  
      cout&lt;&lt;endl;
      cout &lt;&lt; "gpu looping starts, and in progress ..." &lt;&lt; endl;
      StartCounter(6);
  
      double totalfileiotime = 0;
      double totalh2dcpytime = 0;
      double totalkerneltime = 0;
      double totald2hcpytime = 0;
      double totalcpusumtime = 0;
      double totalloopingtime = 0;
  
      for (int k = 0; k &lt; L; k++)
      {
          StartCounter(1);
          img = (unsigned short*)LoadTIFF(filelist[k].c_str());
          totalfileiotime += GetCounter(1);
  
          psum = sum;
          pimg = img;
  
          float gpumean = 0;
  
          memset(psum, 0, sizeof(unsigned long long)*num_blocks);
  
          StartCounter(2);
          // Copy input vectors from host memory to GPU buffers.
          cudaStatus = cudaMemcpy(dev_img, pimg, size * sizeof(unsigned short), cudaMemcpyHostToDevice);
          if (cudaStatus != cudaSuccess) 
          {
              fprintf(stderr, "cudaMemcpy failed!");
              goto Error;
          }
  
          cudaStatus = cudaMemcpy(dev_sum, psum, num_blocks*sizeof(unsigned long long), cudaMemcpyHostToDevice);
          if (cudaStatus != cudaSuccess) 
          {
              fprintf(stderr, "cudaMemcpy failed!");
              goto Error;
          }   
  
          totalh2dcpytime += GetCounter(2);
  
          StartCounter(3);
          //reduceSum&lt;&lt;&lt;num_blocks,block_size,num_blocks * sizeof(unsigned long long)&gt;&gt;&gt;(dev_img, dev_sum, size);
           //reduceSum&lt;&lt;&lt;num_blocks,block_size,block_size * sizeof(unsigned short)&gt;&gt;&gt;(dev_img, dev_sum, size);
            reduceSum&lt;&lt;&lt;num_blocks,block_size&gt;&gt;&gt;(dev_img, dev_sum, size);
          totalkerneltime += GetCounter(3);
  
        // Check for any errors launching the kernel
          cudaStatus = cudaGetLastError();
          if (cudaStatus != cudaSuccess) 
          {
              fprintf(stderr, "reduction Kernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
              goto Error;
          }
  
          // cudaDeviceSynchronize waits for the kernel to finish, and returns
          // any errors encountered during the launch.
  
                  // !!!!!! following is where the code 77 error occurs!!!!!!!
          cudaStatus = cudaDeviceSynchronize();
          if (cudaStatus != cudaSuccess) 
          {
              fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
              goto Error;
          }
  
          // Copy output vector from GPU buffer to host memory.
          StartCounter(4);
          cudaStatus = cudaMemcpy(psum, dev_sum, num_blocks * sizeof(unsigned long long ), cudaMemcpyDeviceToHost);
          if (cudaStatus != cudaSuccess) 
          {
              fprintf(stderr, "cudaMemcpy failed!");
              goto Error;
          }
          totald2hcpytime += GetCounter(4);
  
          StartCounter(5);
          for (int i = 0; i &lt; num_blocks; i++)
          {
              gpumean += *psum;
              psum++;
          }
  
          gpumean /= N*N;
          totalcpusumtime += GetCounter(5);
  
          delete img; 
          img = NULL;
  
          cout&lt;&lt;gpumean&lt;&lt;endl;
  
      }
  
      int S = 1e+6;
      int F = filelist.size();
      float R = S/F;
  
      totalloopingtime = GetCounter(6);
      cout&lt;&lt;"gpu looping ends."&lt;&lt;endl&lt;&lt;endl;
      cout&lt;&lt; "analysis:"&lt;&lt;endl;
      cout&lt;&lt;"gpu initialization time: "&lt;&lt;totalgpuinittime&lt;&lt;" sec"&lt;&lt;endl&lt;&lt;endl;
      cout&lt;&lt;"file I/O time: "&lt;&lt;endl;
      cout&lt;&lt;" total "&lt;&lt;totalfileiotime&lt;&lt;" sec | average "&lt;&lt;totalfileiotime*R&lt;&lt;" usec/frame"&lt;&lt;endl&lt;&lt;endl;
      cout&lt;&lt;"host-to-device copy time: "&lt;&lt;endl;
      cout&lt;&lt;" total "&lt;&lt;totalh2dcpytime&lt;&lt;" sec | average "&lt;&lt;totalh2dcpytime*R&lt;&lt;" usec/frame"&lt;&lt;endl&lt;&lt;endl;
      cout&lt;&lt;"pure gpu kerneling time: "&lt;&lt;endl;
      cout&lt;&lt;" total "&lt;&lt;totalkerneltime&lt;&lt;" sec | average "&lt;&lt;totalkerneltime*R&lt;&lt;" usec/frame"&lt;&lt;endl&lt;&lt;endl;
      cout&lt;&lt;"device-to-host copy time: "&lt;&lt;endl;
      cout&lt;&lt;" total "&lt;&lt;totald2hcpytime&lt;&lt;" sec | average "&lt;&lt;totald2hcpytime*R&lt;&lt;" usec/frame"&lt;&lt;endl&lt;&lt;endl;
      /*cout&lt;&lt;"cpu summing time: "&lt;&lt;endl;
      cout&lt;&lt;" total: "&lt;&lt;totalcpusumtime&lt;&lt;" sec | average: "&lt;&lt;totalcpusumtime*R&lt;&lt;" usec/frame"&lt;&lt;endl&lt;&lt;endl;;*/
  
      /*cout &lt;&lt;"gpu looping time: " &lt;&lt; endl;
      cout&lt;&lt;" total: "&lt;&lt;totalloopingtime&lt;&lt;" sec | average: "&lt;&lt;totalloopingtime*R&lt;&lt;" usec/frame"&lt;&lt;endl;*/
  
  
  Error:
      cudaFree(dev_sum);
      cudaFree(dev_img);
  
      delete sum;
      sum = NULL;
  
      return cudaStatus;
  }
  
  void kernel(float* &amp;mean, int N, vector&lt;string&gt; filelist)
  {
      // wrapper and kernel
      cudaError_t cudaStatus = gpuWrapper(mean, N,  filelist);
  
      if (cudaStatus != cudaSuccess) 
      {
          fprintf(stderr, "gpuWapper failed!");
  
      }
  
     // printf("mean is: %f\n", mean);
  
      // cudaDeviceReset must be called before exiting in order for profiling and
      // tracing tools such as Nsight and Visual Profiler to show complete traces.
  
      StartCounter(8);
      cudaStatus = cudaDeviceReset();
      if (cudaStatus != cudaSuccess) 
      {
          fprintf(stderr, "cudaDeviceReset failed!");
  
      }
      cout&lt;&lt;"gpu reset time: "&lt;&lt;GetCounter(8)&lt;&lt;" sec"&lt;&lt;endl&lt;&lt;endl;
      //return *mean;
  }
  </code></pre>
  
  <p>I have assigned enough and equivalent memory space for both host and device memory. Any comments is appreciated.</p>