解释 ptxas 的详细输出，第一部分答案

【问题标题】：Interpreting the verbose output of ptxas, part I解释 ptxas 的详细输出，第一部分
【发布时间】：2012-09-05 11:24:24
【问题描述】：

我正在尝试了解我的每个 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 byte spill stores和0 bytes spill loads有什么区别
为什么cmem 的信息（我假设它是常量内存）用不同的数字重复了两次？在内核中我没有使用任何常量记忆。这是否意味着编译器在后台会告诉 GPU 使用一些常量内存？

_{这个问题“继续”在：Interpreting the verbose output of ptxas, part II}

【问题讨论】：

'Used 46 registers' 表示编译器为每个线程保留了 46 个寄存器用于编译的内核，其他寄存器被溢出。您可以通过从内核 PTX 中使用的寄存器总数中减去这个数字 (46) 来找到溢出寄存器的数量。
@Ahmad：你的第一句话是正确的，但第二句话不是。内核可以使用少于每个线程允许的最大寄存器，并且不会溢出到本地内存。
为了详细说明 talonmies 回复，PTX 是具有无限寄存器的高级抽象。那是因为它可以为多代GPU编译并且寄存器的数量可以不同。只有当您编译为特定于机器的代码时，您才能真正了解寄存器的使用情况。在任何情况下，ptxas（将 PTX 编译为特定于机器的代码）都会告诉您溢出的数量。
编译器还使用常量内存来存储数字常量（如果它们太大而无法在指令操作码中硬编码）。虽然我不确定如果你说你自己在程序中不使用常量内存，为什么 cmem 会重复两次
curiousexplorer：我已经为这个问题发布了“第二部分”。

标签： memory cuda gpu-constant-memory ptxas

【解决方案1】：

每个 CUDA 线程使用 46 个寄存器？ 是的，正确
没有寄存器溢出到本地内存？ 是的，正确
__global__ 和__device__ 函数的堆栈帧的内存总和是 72 字节吗？ 是的，正确
0 字节溢出存储和 0 字节溢出加载有什么区别？
- 公平的问题，负载可能大于存储，因为您可能会溢出计算值，加载一次，丢弃它（即将其他内容存储到该寄存器中）然后再次加载（即重用它）。 更新： 另请注意，溢出加载/存储计数基于@njuffa 在下面的 cmets 中描述的静态分析
为什么 cmem 的信息（我假设它是常量内存）用不同的数字重复了两次？在内核中，我没有使用任何常量内存。这是否意味着编译器会在后台告诉 GPU 使用一些常量内存？
- 常量内存用于几个目的，包括__constant__ 变量和内核参数，使用不同的“银行”，开始变得有点详细，但只要你为__constant__ 变量和内核参数小于 4KB 就可以了。

【讨论】：

请注意，溢出加载和存储是静态计算的，即本地加载和本地存储指令的数量乘以每个加载/存储的访问宽度。它们被标准化为字节，因为如果编译器有足够的关于对齐和寄存器分配的信息，它可能能够向量化溢出加载/存储。由于计数是静态的，因此这不能直接衡量溢出流量，因为溢出/填充可能在循环内。如果重复使用溢出的数据，溢出负载可能会超过溢出存储。这意味着溢出加载字节 >= 溢出存储字节。
谢谢@njuffa - 好点。编译器无法知道循环的行程计数（除非编译时常量）。真正分析溢出/填充成本的最佳方法是使用分析器，例如 Nsight（或独立的 NVVP），它会根据执行而不是编译为您提供数据。
同意分析。编译器的溢出统计信息作为一线指标有一定的用处。如果没有溢出，则无需担心。如果数字很小（例如