Я пытаюсь уменьшить количество используемых регистров в своем ядре, поэтому решил попробовать встроенный 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 других переменных нет. Как я могу решить эту ситуацию?
nvcc
с помощью-maxrregcount
, см. stackoverflow.com/a/17554199/678093 - person m.s.   schedule 27.03.2015