Cuda PTX регистрирует объявление и использование

Я пытаюсь уменьшить количество используемых регистров в своем ядре, поэтому решил попробовать встроенный PTX.

Это ядро:

#define Feedback(a, b, c, d, e) d^e^(a&c)^(a&e)^(b&c)^(b&e)^(c&d)^(d&e)^(a&d&e)^(a&c&e)^(a&b&d)^(a&b&c)

__global__ void Test(unsigned long a, unsigned long b, unsigned long c, unsigned long d, unsigned long e, unsigned long f, unsigned long j, unsigned long h, unsigned long* res)
{
    res[0] = Feedback( a, b, c, d, e );  
    res[1] = Feedback( b, c, d, e, f );
    res[2] = Feedback( c, d, e, f, j );  
    res[3] = Feedback( d, e, f, j, h );
}  

Используя 14 регистров, я думаю, что это больше, чем нужно, поэтому я пишу Inline PTX:

    __global__ void Feedback_ASM(unsigned long a, unsigned long b, unsigned long c, unsigned long d, unsigned long e, unsigned long f, unsigned long j, unsigned long h, unsigned long* res)
{
asm(".reg .u32 %r<10>;\n");

// 1
asm("ld.param.u32   %r1, [__cudaparm__Z7Feedback_ASMmmmmmmmmPm_a];\n"
    "ld.param.u32   %r2, [__cudaparm__Z7Feedback_ASMmmmmmmmmPm_b];\n"
    "ld.param.u32   %r3, [__cudaparm__Z7Feedback_ASMmmmmmmmmPm_c];\n"
    "ld.param.u32   %r4, [__cudaparm__Z7Feedback_ASMmmmmmmmmPm_d];\n"
    "ld.param.u32   %r5, [__cudaparm__Z7Feedback_ASMmmmmmmmmPm_e];\n");

asm("and.b32 %r7, %r1, %r3;\n"
    "xor.b32 %r8, %r7, %r4;\n"
    "xor.b32 %r7, %r8, %r5;\n"
    "and.b32 %r8, %r1, %r5;\n"
    "xor.b32 %r9, %r7, %r8;\n"
    .............................
    "xor.b32 %r8, %r7, %r9;\n"
    "and.b32 %r6, %r1, %r2;\n"
    "and.b32 %r7, %r6, %r3;\n"
    "xor.b32 %r9, %r7, %r8;\n");

asm("ld.param.u32   %r8, [__cudaparm__Z7Feedback_ASMmmmmmmmmPm_res];\n"
    "st.global.u32  [%r8+0], %r9;");     
// 2
...
// 3
...
// 4
...
}     

Но это ядро ​​тоже использует 14 регистров! Я немного запутался. Я задекларировал всего 10 регистров, В файле ptx других переменных нет. Как я могу решить эту ситуацию?


person Alex Teos    schedule 27.03.2015    source источник
comment
PTX - это всего лишь виртуальная машина, промежуточное представление кода, который в конечном итоге запускается на оборудовании (и его SSA, который, я думаю, ваш код фактически нарушает, я удивлен, что он компилируется). Нет никакой связи между распределением регистров в PTX и окончательным распределением регистров в машинном коде, созданном ассемблером.   -  person talonmies    schedule 27.03.2015
comment
Он не только компилируется, но и работает корректно. Есть ли способ повлиять на количество регистров? На мой взгляд, они действительно использовали слишком много.   -  person Alex Teos    schedule 27.03.2015
comment
Вы передаете желаемые вычислительные возможности во время компиляции? Кроме того, почему для вас действительно так важно уменьшить количество используемых регистров? 14 вроде не много.   -  person jepio    schedule 27.03.2015
comment
Потому что в этом случае я вычисляю только 4 переменные, а в полной программе мне нужно 32 переменные, поэтому количество регистров увеличивается до 60. В результате у меня есть 17% занятости. Я пытаюсь найти способ контролировать этот процесс. P.S. Извините, но я не понимаю ваш первый вопрос :(   -  person Alex Teos    schedule 27.03.2015
comment
Какую версию CUDA вы используете и как компилируете? Вы можете влиять на nvcc с помощью -maxrregcount, см. stackoverflow.com/a/17554199/678093   -  person m.s.    schedule 27.03.2015


Ответы (1)


Как уже указывалось, РТХ является промежуточным кодом. «Регистры» PTX являются виртуальными регистрами и не обязательно отражают фактическое использование регистров устройства.

Чтобы получить представление о фактическом использовании регистра устройства, выполните компиляцию с подробной опцией ptxas:

nvcc -Xptxas -v ...

или воспользуйтесь одним из профилировщиков. Вы также можете проверить машинный код напрямую, используя:

cuobjdump -sass myexe

(где myexe заменяется именем вашего исполняемого файла).

Чтобы контролировать использование регистров, вы можете использовать параметр компиляции nvcc:

nvcc -maxrregcount 10 ...

(где 10 заменено количеством регистров на поток, которым вы хотите ограничить все ядра в вашем коде) или вы можете использовать запустить директиву bounds в вашем коде, которая может контролировать использование регистров для каждого ядра отдельно.

person Robert Crovella    schedule 01.04.2015