CUDA большие входные массивы

Question

Я новичок в CUDA, я работал над "алгоритмом сокращения".

Алгоритм работает для любого размера массива менее 1<<24.

Когда я использую массивы размером 1<<25, программа возвращает 0 в "общей сумме", что неверно. Сумма должна мне 2^25

РЕДАКТИРОВАТЬ cuda-memcheck compiled_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]
=========

Моя настройка:

• Элемент списка
• Nvidia Tesla K40
• CUDA 6.5
• Scientific Linux выпуск 6.4 (Carbon)

Программа состоит из ядра, оболочки ядра и основной для запуска оболочки ядра.

/* -------- 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>0; s>>=1)
  {
    if (tid < s)
    {
      if (pos+s < 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) && (pos < size))
  {
    d_out[blockIdx.x] = d_in[pos];
  }
}

Вот ядро ​​оболочки

/* -------- 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(&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<<<num_blocks, num_threads>>>(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(&d_intermediate, sizeof(int)*num_blocks);
  }
  while(size > num_threads); // if it is too small, compute rest.

  // computing rest
  reduce_kernel<<<1, size>>>(d_out, d_in, prev_num_blocks);
}

Вот главное:

/* -------- 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<<24;
  const int size_out = 1;
  int * d_in;
  int * d_out;
  cudaMalloc(&d_in, sizeof(int)*size);
  cudaMalloc(&d_out, sizeof(int)*size_out);

  int * h_in = (int *)malloc(size*sizeof(int));
  for (int i = 0; i <  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(&result, d_out, sizeof(int), cudaMemcpyDeviceToHost);
  printf("\nFINAL SUM IS: %d\n", result);
}

Answer 1

Этот метод компиляции вашего кода:

nvcc -o my_reduce my_reduce.cu

сборка для вычислительной архитектуры cc2.0 на CUDA 6.5

Эта архитектура ограничена 65535 блоками (в х-измерении, которое является единственным измерением, которое вы используете) в сетке.

на size из 1<<24с num_threads=512количество запущенных блоков:

  num_blocks = (size) / num_threads;

что составляет 1<<24/512 или 31250 блоков

при некотором числе чуть выше 1<<25 вы превысите предел блока устройства cc2.0.

Чтобы это исправить, скомпилируйте с

nvcc -o -arch=sm_35 my_reduce my_reduce.cu

которая является правильной архитектурой компиляции (т.е. возможностью вычислений) для вашего K40, и увеличит предел блока до 2^31-1

И, пожалуйста, используйте правильную проверку ошибок CUDA каждый раз, когда у вас возникают проблемы с кодом CUDA, прежде чем обращаться за помощью здесь. Даже если вы не понимаете результатов ошибки, это, вероятно, поможет тем, кто пытается вам помочь.