我正在尝试了解我的每个 CUDA 线程的资源使用情况,以用于手写内核。
我将我kernel.cu
的文件编译成一个kernel.o
文件nvcc -arch=sm_20 -ptxas-options=-v
我得到以下输出(通过c++filt
):
ptxas info : Compiling entry function 'searchkernel(octree, int*, double, int, double*, double*, double*)' for 'sm_20'
ptxas info : Function properties for searchkernel(octree, int*, double, int, double*, double*, double*)
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__
函数。__global__
和__device__
函数的堆栈帧的内存总和是 72 字节吗?0 byte spill stores
和有什么区别0 bytes spill loads
为什么信息
cmem
(我假设是恒定记忆)用不同的数字重复两次?在内核中,我没有使用任何常量内存。这是否意味着编译器会在后台告诉 GPU 使用一些常量内存?
这个问题在以下内容中“继续”:解释 ptxas 的详细输出,第二部分