Какие переменные потребляют регистры в CUDA? 

__global__ void add( int *c, const int* a, const int* b )
{
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x;
    c[offset] = a[offset] + b[offset];
}

В приведенном выше примере, я думаю, x, y, offset сохраняются в регистрах, а

  • nvcc -Xptxas -v дает 4 registers, 24+16 bytes smem

  • профайлер показывает 4 регистра

  • и глава файла ptx:

    .reg .u16 %rh<4>;
.reg .u32 %r<9>;    
.reg .u64 %rd<10>;  
.loc    15  21  0   

$LDWbegin__Z3addPiPKiS1_:   
.loc    15  26  0

Может ли кто-нибудь уточнить использование регистров? В Ферми максимальное количество регистров составляет 63 для каждого потока. В моей программе я хочу проверить случай, когда ядро ​​потребляет слишком много регистров (так что переменные могут быть автоматически сохранены в локальной памяти и, следовательно, приводят к снижению производительности). Затем в этот момент я могу разделить одно ядро ​​на два, чтобы каждый поток имел достаточно регистров. Предположим, что ресурсов SM достаточно для одновременных ядер.

Я не уверен, прав ли я.

Ответы

Ответ 1

Распределение регистров в PTX совершенно не имеет отношения к конечному потреблению ядра. PTX является лишь промежуточным представлением конечного машинного кода и использует статическую единую форму назначения, что означает, что каждый регистр в PTX используется только один раз. Часть PTX с сотнями регистров может скомпилироваться в ядро ​​с несколькими регистрами.

Назначение регистров выполняется с помощью ptxas как полностью автономный прогон компиляции (как статически, так и точно вовремя драйвером, либо и то, и другое), и он может выполнять много переупорядочения кода и оптимизации на входном PTX для улучшения пропускная способность и сохранение регистров, что означает, что между переменными исходного C или регистрами в PTX и окончательным количеством регистров собранного ядра существует незначительная или вообще отсутствует связь.

nvcc предоставляет некоторые способы влияния на поведение распределения регистров ассемблера. У вас есть __launch_bounds__, чтобы предоставить эвристические подсказки компилятору, который может влиять на распределение регистров, а компилятор/ассемблер принимает аргумент -maxrregcount (при потенциальном расходе переноса регистров в локальную память, что может снизить производительность). Ключевое слово volatile используется, чтобы иметь значение для более старых версий компилятора на основе nvopen64 и может влиять на поведение разлива локальной памяти. Но вы не можете произвольно контролировать или управлять распределением регистров в исходном коде кода C или в языке сборки PTX.