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

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

1 голос
/ 24 февраля 2011

Ядро CUDA, скомпилированное с параметром --ptxas-options=-v, похоже, отображает ошибочную статистику lmem (локальная память) , если указана sm_20 архитектура GPU.То же самое дает значимую статистику lmem с 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 сообщает значение 4 байта , что просто невозможно, если вы видите массив байтов 4x1000, используемый в ядре.Более старые архитектуры GPU сообщают правильную статистику 4000 байт lmem.

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

1 Ответ

1 голос
/ 24 февраля 2011

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

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

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

...