Зачем тратить так много регистров? - PullRequest
Новая
       70

Зачем тратить так много регистров?

0 голосов
/ 22 сентября 2019

Я немного поиграл с кодом вроде 

__shared__ int temp[512];
...
temp[threadIdx.x] = pSource[i];

#pragma unroll
for (int j = 1; j < 16; j++)
    temp[threadIdx.x] += pSource[i + blockDim.x * j];

Компилятор развернул этот цикл (даже без прагмы) и потратил 26 регистров на поток для моей довольно маленькой процедуры (даже если она содержит какой-то другой кодостальное не требует даже близко к этому количеству регистров).Сгенерированный код выглядит так (смотрите сначала в конце) 

IADD R12, R3, R6;
MOV32I R13, 0x4;
MOV R14, c[0x0][0x8];
IADD R2, R12, c[0x0][0x8];
IMAD R4.CC, R12, R13, c[0x0][0x20];
ISCADD R3, R14, R12, 0x2;
IMAD R15, R14, 0x7, R12;
IMAD.HI.X R5, R12, R13, c[0x0][0x24];
IMAD.U32.U32 R8.CC, R2, R13, c[0x0][0x20];
IMAD.U32.U32 RZ, R1, RZ, RZ;
LD.E R11, [R4];
IMAD.U32.U32.HI.X R9, R2, R13, c[0x0][0x24];
ISCADD R2, R14, R12, 0x1;
LD.E R10, [R8];
IMAD.U32.U32 R18.CC, R2, R13, c[0x0][0x20];
IMAD R4, R14, 0x5, R12;
IMAD R5, R14, 0x6, R12;
IMAD.U32.U32.HI.X R19, R2, R13, c[0x0][0x24];
IMAD R2, R14, 0x3, R12;
LD.E R9, [R18];
IMAD.U32.U32 R16.CC, R2, R13, c[0x0][0x20];
IMAD.U32.U32.HI.X R17, R2, R13, c[0x0][0x24];
...
LD.E R19, [R22];                              
IADD R10, R11, R10;                           
IMAD R11, R14, 0xe, R12;                      
IMAD R14, R14, 0xf, R12;                      
IADD R9, R10, R9;                             
IMAD.U32.U32 R10.CC, R11, R13, c[0x0][0x20];  
IMAD.U32.U32.HI.X R11, R11, R13, c[0x0][0x24];
IMAD.U32.U32 R12.CC, R14, R13, c[0x0][0x20];  
LD.E R10, [R10];                              
IMAD.U32.U32.HI.X R13, R14, R13, c[0x0][0x24];
IADD R8, R9, R8;                              
LD.E R12, [R12];                              
IADD R7, R8, R7;                              
IADD R2, R7, R2;                              
IADD R2, R2, R18;                             
IADD R2, R2, R17;                             
IADD R2, R2, R16;                             
IADD R2, R2, R15;                             
IADD R2, R2, R5;                              
IADD R2, R2, R4;                              
IADD R2, R2, R3;                              
IADD R2, R2, R19;                             
IADD R3, R2, R10;

Таким образом, компилятор в основном откладывает окончательное суммирование до самого конца.Для чего?

У меня достаточно потоков для загрузки конвейеров ALU, поэтому это не важно, когда доступны результаты вычислений ... Проблема только в доступе к памяти устройства - похоже, было бы необходимо иметь около 100 деформацийчтобы скрыть его задержку ... Ограничение использования регистров до 20 ускорилось примерно на 10%, так как 6 блоков по 256 потоков каждый мог уместиться в многопроцессорную систему вместо 4 ... Мне интересно, есть ли другие нюансы архитектуры, которые пытался компиляториспользовать здесь?

NVidia GT520M (возможность 2.1), пробовал компилировать до 2.0 и 3.0 с тем же эффектом.

...