Предотвращение ненужных операций mov во встроенном PTX - PullRequest
Новая
       39

Предотвращение ненужных операций mov во встроенном PTX

2 голосов
/ 31 марта 2012

При записи PTX в отдельный файл параметр ядра может быть загружен в регистр с помощью: 

.reg .u32 test;
ld.param.u32 test, [test_param];

Однако, при использовании встроенного PTX, Использование встроенной сборки PTX в CUDA (версия 01) примечание к приложению описывает синтаксис, в котором загрузка параметра тесно связана с другой операцией.Он предоставляет такой пример: 

asm("add.s32 %0, %1, %2;" : "=r"(i) : "r"(j), "r"(k));

, который генерирует: 

ld.s32 r1, [j];
ld.s32 r2, [k];
add.s32 r3, r1, r2;
st.s32 [i], r3;

Во многих случаях необходимо разделить две операции.Например, может потребоваться сохранить параметр в регистре вне цикла, а затем повторно использовать и модифицировать регистр внутри цикла.Единственный способ сделать это - использовать дополнительную инструкцию mov, чтобы переместить параметр из регистра, в который он был загружен неявно, в другой регистр, который я смогу использовать позже.

Есть ли способизбежать этой дополнительной инструкции mov при переходе от PTX в отдельном файле к встроенному PTX?

1 Ответ

3 голосов
/ 03 апреля 2012

На вашем месте я бы не слишком беспокоился об этих операциях mov.

Имейте в виду, что PTX не является кодом окончательной сборки.PTX дополнительно компилируется в CUBIN перед запуском ядра.Среди прочего, этот последний шаг выполняет распределение регистров и удаляет все ненужные mov операции.

В частности, если вы переходите от %r1 к %r2 и затем никогда не используете %r1 вообщеалгоритм может присвоить %r1 и %r2 одному и тому же аппаратному регистру и удалить перемещение.

...