title: CUDA large input arrays
body: <p>I am new to CUDA, I have been working on a "Reduce algorithm".</p>
  
  <p>The algorithm works for any array size less than 1&lt;&lt;24.</p>
  
  <p>When I use arrays of size 1&lt;&lt;25 the program returns 0 in "total sum" which is wrong. The sum should me 2^25</p>
  
  <p><strong>EDIT</strong> cuda-memcheck compiled_code</p>
  
  <pre><code>========= CUDA-MEMCHECK
  @@STARTING@@ 
  ========= Program hit cudaErrorInvalidValue (error 11) due to "invalid argument" on CUDA API call to cudaLaunch. 
  =========     Saved host backtrace up to driver entry point at error
  =========     Host Frame:/usr/lib64/libcuda.so.1 [0x2f2d83]
  =========     Host Frame:test [0x3b37e]
  =========     Host Frame:test [0x2b71]
  =========     Host Frame:test [0x2a18]
  =========     Host Frame:test [0x2a4c]
  =========     Host Frame:test [0x2600]
  =========     Host Frame:test [0x2904]
  =========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1ed5d]
  =========     Host Frame:test [0x23e9]
  =========
  </code></pre>
  
  <p>My setup is:</p>
  
  <ul>
  <li>List item</li>
  <li>Nvidia Tesla K40</li>
  <li>CUDA 6.5</li>
  <li>Scientific Linux release 6.4 (Carbon)</li>
  </ul>
  
  <p>The program consists of a kernel, a kernel wrapper and a main to execute the kernel wrapper.</p>
  
  <pre><code>/* -------- KERNEL -------- */
  __global__ void reduce_kernel(int * d_out, int * d_in, int size)
  {
    // position and threadId
    int pos = blockIdx.x * blockDim.x + threadIdx.x;
    int tid = threadIdx.x;
  
    // do reduction in global memory
    for (unsigned int s = blockDim.x / 2; s&gt;0; s&gt;&gt;=1)
    {
      if (tid &lt; s)
      {
        if (pos+s &lt; size) // Handling out of bounds
        {
          d_in[pos] = d_in[pos] + d_in[pos+s];
        }
      }
      __syncthreads();
    }
  
    // only thread 0 writes result, as thread
    if ((tid==0) &amp;&amp; (pos &lt; size))
    {
      d_out[blockIdx.x] = d_in[pos];
    }
  }
  </code></pre>
  
  <p>Here is the Kernel wrapper</p>
  
  <pre><code>/* -------- KERNEL WRAPPER -------- */
  void reduce(int * d_out, int * d_in, int size, int num_threads)
  {
    // setting up blocks and intermediate result holder
  
    int num_blocks;
    if(((size) % num_threads))
      {
        num_blocks = ((size) / num_threads) + 1;    
      }
      else
      {
        num_blocks = (size) / num_threads;
      }
    int * d_intermediate;
    cudaMalloc(&amp;d_intermediate, sizeof(int)*num_blocks);
    cudaMemset(d_intermediate, 0, sizeof(int)*num_blocks);
    int prev_num_blocks;
    int i = 1;
    int size_rest = 0;
    // recursively solving, will run approximately log base num_threads times.
    do
    {
      printf("Round:%.d\n", i);
      printf("NumBlocks:%.d\n", num_blocks);
      printf("NumThreads:%.d\n", num_threads);
      printf("size of array:%.d\n", size);
      i++;
      reduce_kernel&lt;&lt;&lt;num_blocks, num_threads&gt;&gt;&gt;(d_intermediate, d_in, size);
      size_rest = size % num_threads;
      size = size / num_threads + size_rest;
  
      // updating input to intermediate
      cudaMemcpy(d_in, d_intermediate, sizeof(int)*num_blocks, cudaMemcpyDeviceToDevice);
  
      // Updating num_blocks to reflect how many blocks we now want to compute on
      prev_num_blocks = num_blocks;
      if(size % num_threads)
      {
        num_blocks = size / num_threads + 1;      
      }
      else
      {
        num_blocks = size / num_threads;
      }
      // updating intermediate
      cudaFree(d_intermediate);
      cudaMalloc(&amp;d_intermediate, sizeof(int)*num_blocks);
    }
    while(size &gt; num_threads); // if it is too small, compute rest.
  
    // computing rest
    reduce_kernel&lt;&lt;&lt;1, size&gt;&gt;&gt;(d_out, d_in, prev_num_blocks);
  }
  </code></pre>
  
  <p>Here is the main:</p>
  
  <pre><code>/* -------- MAIN -------- */
  int main(int argc, char **argv)
  {
    printf("@@STARTING@@ \n");
    // Setting num_threads
    int num_threads = 512;
    // Making non-bogus data and setting it on the GPU
    const int size = 1&lt;&lt;24;
    const int size_out = 1;
    int * d_in;
    int * d_out;
    cudaMalloc(&amp;d_in, sizeof(int)*size);
    cudaMalloc(&amp;d_out, sizeof(int)*size_out);
  
    int * h_in = (int *)malloc(size*sizeof(int));
    for (int i = 0; i &lt;  size; i++) h_in[i] = 1;
    cudaMemcpy(d_in, h_in, sizeof(int)*size, cudaMemcpyHostToDevice);
  
    // Running kernel wrapper
    reduce(d_out, d_in, size, num_threads);
    int result;
    cudaMemcpy(&amp;result, d_out, sizeof(int), cudaMemcpyDeviceToHost);
    printf("\nFINAL SUM IS: %d\n", result);
  }
  </code></pre>