cuda - минимальный пример, высокое использование регистра
Рассмотрим эти три тривиальных минимальных ядра. Их использование регистра намного выше, чем я ожидаю. Зачем?
A:
__global__ void Kernel_A()
{
//empty
}
соответствующий PTX:
ptxas info : Compiling entry function '_Z8Kernel_Av' for 'sm_20'
ptxas info : Function properties for _Z8Kernel_Av
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 2 registers, 32 bytes cmem[0]
B:
template<uchar effective_bank_width>
__global__ void Kernel_B()
{
//empty
}
template
__global__ void Kernel_B<1>();
соответствующий PTX:
ptxas info : Compiling entry function '_Z8Kernel_BILh1EEvv' for 'sm_20'
ptxas info : Function properties for _Z8Kernel_BILh1EEvv
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 2 registers, 32 bytes cmem[0]
C:
template<uchar my_val>
__global__ void Kernel_C
(uchar *const device_prt_in,
uchar *const device_prt_out)
{
//empty
}
соответствующий PTX:
ptxas info : Compiling entry function '_Z35 Kernel_CILh1EEvPhS0_' for 'sm_20'
ptxas info : Function properties for _Z35 Kernel_CILh1EEvPhS0_
16 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 10 registers, 48 bytes cmem[0]
Вопрос:
Почему пустые ядра A и B использовали 2 регистра? CUDA всегда использует один неявный регистр, но почему используются 2 дополнительных явных регистра?
Ядро С еще более расстраивает. 10 регистров? Но есть только 2 указателя. Это дает 2*2 = 4 регистра для указателей. Даже если есть дополнительно 2 загадочных регистра (предложенных ядром A и ядром B), это даст всего 6. Все еще намного меньше чем 10!
В случае, если вы заинтересованы, вот ptx код для ядра А. ptx код для ядра B точно такой же, по модулю целочисленных значений и имен переменных.
.visible .entry _Z8Kernel_Av(
)
{
.loc 5 19 1
func_begin0:
.loc 5 19 0
.loc 5 19 1
func_exec_begin0:
.loc 5 22 2
ret;
tmp0:
func_end0:
}
И для ядра C...
.weak .entry _Z35Kernel_CILh1EEvPhS0_(
.param .u64 _Z35Kernel_CILh1EEvPhS0__param_0,
.param .u64 _Z35Kernel_CILh1EEvPhS0__param_1
)
{
.local .align 8 .b8 __local_depot2[16];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .s64 %rd<3>;
.loc 5 38 1
func_begin2:
.loc 5 38 0
.loc 5 38 1
mov.u64 %SPL, __local_depot2;
cvta.local.u64 %SP, %SPL;
ld.param.u64 %rd1, [_Z35Kernel_CILh1EEvPhS0__param_0];
ld.param.u64 %rd2, [_Z35Kernel_CILh1EEvPhS0__param_1];
st.u64 [%SP+0], %rd1;
st.u64 [%SP+8], %rd2;
func_exec_begin2:
.loc 5 836 2
tmp2:
ret;
tmp3:
func_end2:
}
- Почему он сначала объявляет переменную локальной памяти (
.local) - Почему два указателя (заданные в качестве аргументов функции) хранятся в регистрах? Разве для них не предусмотрено специальное пространство для параметров?
- Возможно, два указателя аргумента функции принадлежат регистрам - это объясняет два
.reg .b64линий. Но что это.reg .s64линия? Почему это там?
Становится еще хуже:
D:
template<uchar my_val>
__global__ void Kernel_D
(uchar * device_prt_in,
uchar *const device_prt_out)
{
device_prt_in = device_prt_in + blockIdx.x*blockDim.x + threadIdx.x;
}
дает
ptxas info : Used 6 registers, 48 bytes cmem[0]
То есть манипулирование аргументом (указателем) уменьшается с 10 до 6 регистров?
1 ответ
Первое, на что нужно обратить внимание: если вы беспокоитесь о регистрах, не смотрите на код PTX, потому что он ничего вам не скажет. PTX использует статическую форму одиночного назначения, а код, выдаваемый компилятором, не включает в себя "декорацию", необходимую для создания работающей точки входа машинного кода.
С этим из пути, давайте посмотрим на ядро A:
$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z8Kernel_Av' for 'sm_20'
ptxas info : Function properties for _Z8Kernel_Av
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 2 registers, 32 bytes cmem[0]
$ cuobjdump -sass null.cubin
code for sm_20
Function : _Z8Kernel_Av
/*0000*/ /*0x00005de428004404*/ MOV R1, c [0x1] [0x100];
/*0008*/ /*0x00001de780000000*/ EXIT;
.............................
Есть два ваших регистра. Пустые ядра не дают нулевых инструкций.
Кроме того, я не могу воспроизвести то, что вы показали. Если я посмотрю на ваше ядро C как написанное, я получу следующее (компилятор релиза CUDA 5):
$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z8Kernel_CILh1EEvPhS0_' for 'sm_20'
ptxas info : Function properties for _Z8Kernel_CILh1EEvPhS0_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 2 registers, 48 bytes cmem[0]
$ cuobjdump -sass null.cubin
code for sm_20
Function : _Z8Kernel_CILh1EEvPhS0_
/*0000*/ /*0x00005de428004404*/ MOV R1, c [0x1] [0x100];
/*0008*/ /*0x00001de780000000*/ EXIT;
........................................
то есть. 2 идентичных регистра кода для первых двух ядер.
и то же самое для ядра D:
$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z8Kernel_DILh1EEvPhS0_' for 'sm_20'
ptxas info : Function properties for _Z8Kernel_DILh1EEvPhS0_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 2 registers, 48 bytes cmem[0]
$ cuobjdump -sass null.cubin
code for sm_20
Function : _Z8Kernel_DILh1EEvPhS0_
/*0000*/ /*0x00005de428004404*/ MOV R1, c [0x1] [0x100];
/*0008*/ /*0x00001de780000000*/ EXIT;
........................................
Опять 2 регистра.
Для записи я использую версию nvcc:
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2012 NVIDIA Corporation
Built on Fri_Sep_28_16:10:16_PDT_2012
Cuda compilation tools, release 5.0, V0.2.1221