__device__ 函数中的 NVCC 寄存器使用报告

Question

我正在尝试使用NVCC选项获取有关CUDA内核中寄存器使用情况的一些信息，虽然使用全局函数一切正常，但我在设备函数方面遇到了一些困难，因为
--ptxas-options=v

ptxas info : Used N registers

输出中缺少行。我尝试使用 noinline 关键字并将它们保存在另一个文件中，相对于调用全局函数，因为我认为NVCC正在报告全局函数的完整寄存器使用情况，包括内联后调用的设备寄存器，但没有任何变化。我可以获得有关设备功能的寄存器使用情况的信息，只需将它们定义为全局即可。

你有什么建议吗？

谢谢！

Answer 1

据我了解，ptxas（设备汇编器）仅输出其链接的代码的寄存器计数。独立__device__函数不由汇编器链接，它们仅被编译。因此，汇编器不会发出设备函数的寄存器计数值。我不相信有解决方法。

但是，仍然可以 __device__通过使用从汇编器输出转储 elf 数据来获取函数 的寄存器占用空间cuobjdump。您可以按如下方式执行此操作：

$ cat vdot.cu
__device__  __noinline__ float vdot(float v1, float v2) {
    return (v1 * v2);
}

__device__ __noinline__  float vdot(float2 v1, float2 v2) {
    return (v1.x * v2.x) + (v1.y * v2.y);
}

__device__ __noinline__ float vdot(float4 v1, float4 v2) {
    return (v1.x * v2.x) + (v1.y * v2.y) + (v1.z * v2.z) + (v1.w * v2.w);
}

$ nvcc -std=c++11 -arch=sm_52 -dc -Xptxas="-v" vdot.cu
ptxas info    : 0 bytes gmem
ptxas info    : Function properties for cudaDeviceGetAttribute
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z4vdotff
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessor
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z4vdot6float4S_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaMalloc
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaGetDevice
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for _Z4vdot6float2S_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Function properties for cudaFuncGetAttributes
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads

__device__这里，我们在设备目标文件中有一组单独编译的三个函数。运行cuobjdump它会发出大量输出，但在其中您将获得每个函数的寄存器计数：

$ cuobjdump -elf ./vdot.o

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
compressed

<---Snipped--->


.text._Z4vdotff
bar = 0 reg = 6 lmem=0  smem=0
0xfec007f1  0x001fc000  0x00570003  0x5c980780  
0x00470000  0x5c980780  0x00370004  0x5c680000  
0xffe007ff  0x001f8000  0x0007000f  0xe3200000  
0xff87000f  0xe2400fff  0x00070f00  0x50b00000

在设备函数输出的第二行中，dot(float, float) 您可以看到该函数使用 6 个寄存器。这是我所知道的检查设备功能寄存器足迹的唯一方法。