我想了解手写内核的每个CUDA线程的资源使用情况。解释--ptxas选项的输出= -v
我编译kernel.cu
文件到kernel.o
文件,nvcc -arch=sm_20 -ptxas-options=-v
,我得到了下面的输出
ptxas info : Compiling entry function '_Z12searchkernel6octreePidiPdS1_S1_' for 'sm_20'
ptxas info : Function properties for _Z12searchkernel6octreePidiPdS1_S1_
72 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 46 registers, 176 bytes cmem[0], 16 bytes cmem[14]
看着上面的输出,是正确地说,
- 每个CUDA线程使用46个寄存器?
- 有没有注册泄漏到本地内存?
我对理解输出也有一些问题。
我的内核调用了一大堆
__device__
函数。 IS为72个字节的__global__
和__device__
函数的堆栈帧的内存总和为 ?是什么
0 byte spill stores
和0 bytes spill loads
之间的差异为什么是
cmem
的信息(我假设是恒定的内存),重复两次用不同的数字?在内核中,我没有使用任何常量的内存。这是否意味着编译器正在引擎盖下,告诉GPU使用一些常量内存?
“用于46个寄存器”表示编译器已经保留了每个线程46个寄存器进行编译的内核和其它寄存器溢出。您可以通过从内核PTX中使用的寄存器总数中减去此数字(46)来查找溢出寄存器的数量。 – ahmad
@Ahmad:你的第一句话是正确的,但第二句话是不正确的。内核可以使用少于每个线程的最大允许寄存器数量,并且没有溢出到本地内存。 – talonmies
为了详细说明talonmies的回复,PTX是一个具有无限寄存器的高级抽象。这是因为它可以编译多代GPU,并且寄存器的数量可以不同。只有当你编译到机器特定的代码时,你才能真正看到注册使用。无论如何,ptxas(将PTX编译为机器特定的代码)会告诉您溢出的数量。 – Tom