CUDA: ошибочная статистика lmem отображается для sm_20?

Question

Ядро CUDA, скомпилированное с опцией --ptxas-options=-v кажется, отображает ошибочную статистику lmem (локальная память), когда sm_20 Архитектура GPU уточняется. То же самое дает значимую статистику sm_10 / sm_11 / sm_12 / sm_13 архитектуры.

Может кто-нибудь уточнить, нужно ли считать статистику sm_20 lmem по-разному или она ошибочна?

Вот ядро:

__global__ void fooKernel( int* dResult )
{
        const int num = 1000;
        int val[num]; 

        for ( int i = 0; i < num; ++i )
        val[i] = i * i; 

        int result = 0; 

        for ( int i = 0; i < num; ++i )
        result += val[i]; 

        *dResult = result;

        return;
}

--ptxas-options=-v а также sm_20 доклад:

1>ptxas info    : Compiling entry function '_Z9fooKernelPi' for 'sm_20'
1>ptxas info    : Used 5 registers, 4+0 bytes lmem, 36 bytes cmem[0]

--ptxas-options=-v а также sm_10 / sm_11 / sm_12 / sm_13 доклад:

1>ptxas info    : Compiling entry function '_Z9fooKernelPi' for 'sm_10'
1>ptxas info    : Used 3 registers, 4000+0 bytes lmem, 4+16 bytes smem, 4 bytes cmem[1]

sm_20 сообщает lmem 4 байта, что просто невозможно, если вы видите массив байтов 4x1000, используемый в ядре. Более старые архитектуры GPU сообщают правильную статистику 4000 байт.

Это было опробовано в CUDA 3.2. Я ссылался на раздел " Статистика генерации кода" в руководстве NVCC (v3.2), но это не помогает объяснить эту аномалию.

Answer 1

Компилятор правильный. Благодаря умной оптимизации массив не нужно хранить. То, что вы делаете, по сути, расчет result += i * i без хранения временных val,

Взгляд на сгенерированный код ptx не покажет различий между sm_10 и sm_20. Декомпиляция сгенерированных кубинов с помощью decuda покажет оптимизацию.

Кстати: старайтесь избегать локальной памяти! Это так же медленно, как глобальная память.