title: CUDA memory allocation issues with large images
body: <p>I have a function for making a histogram from an image (sequential version given (Homework))</p>
  
  <pre><code>CImg&lt; unsigned char &gt; histogramImage = CImg&lt; unsigned char &gt;(BAR_WIDTH * HISTOGRAM_SIZE, HISTOGRAM_SIZE, 1, 1);
  unsigned int *histogram;
  histogram = (unsigned int *)malloc(HISTOGRAM_SIZE * sizeof(unsigned int));
   memset(reinterpret_cast&lt; void * &gt;(histogram), 0, HISTOGRAM_SIZE * sizeof(unsigned int));
  
  cudaMemset(gpuImage, 0, grayImage.width() * grayImage.height() * sizeof(unsigned char));
  
  cuda_err = cudaMemcpy(gpuImage, grayImage, grayImage.width() * grayImage.height() * sizeof(unsigned char), cudaMemcpyHostToDevice);
  if (cuda_err != cudaSuccess)
  {
      std::cout &lt;&lt; "ERROR: Failed cudaMemcpy" &lt;&lt; std::endl;
     return -1;
  }
  
  unsigned int *gpuhistogram;
  cuda_err = cudaMalloc((void **)(&amp;gpuhistogram), HISTOGRAM_SIZE * sizeof(unsigned int));
  if (cuda_err != cudaSuccess)
  {
      std::cout &lt;&lt; "ERROR: Failed cudaMalloc" &lt;&lt; std::endl;
  }
  cudaMemset (gpuhistogram, 0, HISTOGRAM_SIZE * sizeof(unsigned int));
  
  histogram1D(gpuImage, histogramImage, grayImage.width(), grayImage.height(), gpuhistogram, HISTOGRAM_SIZE, BAR_WIDTH, total, gridSize, blockSize);
  
  cuda_err = cudaMemcpy(histogram, gpuhistogram, HISTOGRAM_SIZE * sizeof(unsigned int), cudaMemcpyDeviceToHost);
  if (cuda_err != cudaSuccess)
  {
      std::cout &lt;&lt; "ERROR: Failed cudaMemcpy" &lt;&lt; std::endl;
  }
  </code></pre>
  
  <p>That calls</p>
  
  <pre><code>void histogram1D(unsigned char *grayImage, unsigned char *histogramImage, const int width, const int height, unsigned int *histogram, const unsigned int HISTOGRAM_SIZE, const unsigned int BAR_WIDTH, NSTimer &amp;timer, dim3 grid_size, dim3 block_size) {
  
  NSTimer kernelTime = NSTimer("kernelTime", false, false);
  
  kernelTime.start();
  histo &lt;&lt;&lt; grid_size, block_size &gt;&gt;&gt; (grayImage, histogram,width);
  cudaDeviceSynchronize();
  kernelTime.stop();
  
  cout &lt;&lt; fixed &lt;&lt; setprecision(6);
  cout &lt;&lt; "histogram1D (kernel): \t\t" &lt;&lt; kernelTime.getElapsed() &lt;&lt; " seconds." &lt;&lt; endl;
  }
  </code></pre>
  
  <p>The kernel function is</p>
  
  <pre><code>__global__ void histo(unsigned char *inputImage, unsigned int *histogram, int width)
  {
  
  int x = threadIdx.x + (blockIdx.x * blockDim.x);
  int y = threadIdx.y + (blockIdx.y * blockDim.y);
  
  unsigned int index = static_cast&lt; unsigned int &gt;(inputImage[(y * width) + x]);
  atomicAdd(&amp;histogram[index],1);
  }
  </code></pre>
  
  <p>The problem I have is, when I call this with images ranging from 1024x1024 to 3543x2480, it works. However, I have one image which is 8192x8192 and when the function returns, the values in *histogram are all still 0. My trials seems to indicate it has to do with the memory allocation for *gpuhistogram (shouldn't unsigned int be large enough?)since the sequential version of this works. How do fix this? Any ideas?</p>