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

Question

При записи 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?

Answer 1

Если бы я был тобой, я бы не слишком беспокоился об этих операциях.

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

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