Объем локальной памяти на поток CUDA

Question

Я прочитал в документации NVIDIA ( http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html, таблица #12), что объем локальной памяти на поток 512 Ko для моего 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 (в Linux с CUDA 6.5). Он работал без ошибки для MEMSIZE=200000, но "ошибка нехватки памяти" произошла во время выполнения для MEMSIZE=250000.

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

Answer 1

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

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

Кадр стека 65000 байт, хранилища разлива 0 байт, разливы 0 байт

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

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

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

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

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