我正在尝试了解我的每个 CUDA 线程的资源使用情况,以用于手写内核。
我使用 nvcc -arch=sm_20 -ptxas-options=-v
将我的 kernel.cu
文件编译为 kernel.o
文件
我得到了以下输出(通过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__
函数。 IS 72 字节总和__global__
和__device__
函数的堆栈帧的内存是多少?0 字节溢出存储
和0 字节溢出加载有什么区别
为什么
cmem
的信息(我假设是常量内存)会以不同的数字重复两次?在内核中我没有使用任何常量 内存。这是否意味着编译器在后台会告诉 GPU 使用一些常量内存?
最佳答案
- 每个 CUDA 线程使用 46 个寄存器? 是的,正确
- 没有寄存器溢出到本地内存? 是的,正确
__global__
和__device__
函数的堆栈帧的内存总和是 72 字节吗? 是的,正确- 0 字节溢出存储和 0 字节溢出加载有什么区别?
- 公平的问题,负载可能大于存储,因为您可能会溢出计算值,加载一次,丢弃它(即将其他内容存储到该寄存器中)然后再次加载(即重用它)。 更新: 另请注意,溢出加载/存储计数基于@njuffa 在下面的评论中描述的静态分析
- 为什么 cmem 的信息(我假设是常量内存)会以不同的数字重复两次?在内核中,我没有使用任何常量内存。这是否意味着编译器会在后台告诉 GPU 使用一些常量内存?
- 常量内存用于几个目的,包括
__constant__
变量和内核参数,使用不同的“银行”,开始有点详细,但只要您使用小于 64KB 的内存__constant__
变量和小于 4KB 的内核参数就可以了。
- 常量内存用于几个目的,包括
关于memory - 解释 ptxas 的详细输出,第一部分,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/12388207/