Количество локальной памяти в потоке CUDA

Question

Я прочитал в документации NVIDIA (http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications, таблица № 12), что количество локальной памяти на поток составляет 512 кодов для моего GPU (GTX 580, вычислительная способность 2.0).

Я безуспешно пытался проверить этот лимит на Linux с помощью CUDA 6.5.

Вот код, который я использовал (его единственная цель состоит в том, чтобы проверить локальный предел памяти, он не делает никаких вычислений полезную):

#include <iostream> 
#include <stdio.h> 

#define MEMSIZE 65000 // 65000 -> out of memory, 60000 -> ok 

inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=false) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if(abort) 
      exit(code); 
    } 
} 

inline void gpuCheckKernelExecutionError(const char *file, int line) 
{ 
    gpuAssert(cudaPeekAtLastError(), file, line); 
    gpuAssert(cudaDeviceSynchronize(), file, line);  
} 


__global__ void kernel_test_private(char *output) 
{ 
    int c = blockIdx.x*blockDim.x + threadIdx.x; // absolute col 
    int r = blockIdx.y*blockDim.y + threadIdx.y; // absolute row 

    char tmp[MEMSIZE]; 
    for(int i = 0; i < MEMSIZE; i++) 
     tmp[i] = 4*r + c; // dummy computation in local mem 
    for(int i = 0; i < MEMSIZE; i++) 
     output[i] = tmp[i]; 
} 

int main(void) 
{ 
    printf("MEMSIZE=%d bytes.\n", MEMSIZE); 

    // allocate memory 
    char output[MEMSIZE]; 
    char *gpuOutput; 
    cudaMalloc((void**) &gpuOutput, MEMSIZE); 

    // run kernel 
    dim3 dimBlock(1, 1); 
    dim3 dimGrid(1, 1); 
    kernel_test_private<<<dimGrid, dimBlock>>>(gpuOutput); 
    gpuCheckKernelExecutionError(__FILE__, __LINE__); 

    // transfer data from GPU memory to CPU memory 
    cudaMemcpy(output, gpuOutput, MEMSIZE, cudaMemcpyDeviceToHost); 

    // release resources 
    cudaFree(gpuOutput); 
    cudaDeviceReset(); 

    return 0; 
} 

И команду компиляции строка:

nvcc -o cuda_test_private_memory -Xptxas -v -O2 --compiler-options -Wall cuda_test_private_memory.cu 

компиляция нормально, и отчеты:

ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z19kernel_test_privatePc' for 'sm_20' 
ptxas info : Function properties for _Z19kernel_test_privatePc 
    65000 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 21 registers, 40 bytes cmem[0] 

я получил «из памяти» ошибка во время выполнения на GTX 580, когда я достиг 65000 байт на поток. Вот точный вывод программы в консоли:

MEMSIZE=65000 bytes. 
GPUassert: out of memory cuda_test_private_memory.cu 48 

я также сделал тест с GTX 770 GPU (на Linux с помощью CUDA 6.5). Он работал без ошибок для MEMSIZE = 200000, но ошибка «вне памяти» произошла во время выполнения для MEMSIZE = 250000.

Как это объяснить? Я делаю что-то неправильно ?

Answer 1

Кажется, вы работаете в не локальное ограничение памяти, но ограничение на размер стека:

ptxas информация: Функциональные свойства для _Z19kernel_test_privatePc

65000 байт стека кадра, 0 байт разливами магазины, 0 bytes spill load

Переменная, которую вы намеревались быть локальной, находится в стеке (GPU thread) в этом случае.

На основе информации, предоставленной @njuffa here, доступный предельный размер стека является меньшим из:

1. Максимальный размер локальной памяти (512 Кбайт для cc2.x и выше) памяти
2. GPU/(# OF SMs)/(максимальные потоки на SM)

Очевидно, что первое ограничение не является проблемой. Я предполагаю, что у вас есть «стандартный» GTX580, который имеет 1,5 ГБ памяти и 16 SM. Устройство cc2.x имеет максимум 1536 резидентных потоков на мультипроцессор. Это означает, что мы имеем 1536MB/16/1536 = 1MB/16 = 65536 байт. Существует некоторое накладное и другое использование памяти, которое вычитается из общей доступной памяти, поэтому ограничение размера стека составляет некоторую сумму ниже 65536, где-то между 60000 и 65000 в вашем случае, по-видимому.

Я подозреваю, подобный расчет на вашем GTX770 даст аналогичный результат, т.е. максимальный размер стека между 200000 и 250000.