NVCC寄存器使用情况报告

Question

我试图让使用NVCC选项
--ptxas-options=v我CUDA内核有关的寄存器使用情况的信息，并同时与全球功能一切正常，我有一些困难由于

ptxas info : Used N registers

线设备那些缺少在输出中。我试图使用noinline关键字，并将它们保存在另一个文件中，与调用全局函数有关，因为我认为NVCC报告了全局函数的全部注册使用情况，包括内联后的被调用设备的全局函数，但没有任何变化。我可以获得关于设备功能的寄存器使用情况的信息，只将它们定义为全局的。

你有什么建议吗？

谢谢！

Answer 1

据我所知，ptxas（设备汇编器）只输出一个寄存器计数，它链接的代码。独立的__device__函数没有被汇编器链接，它们只被编译。因此，汇编器不会为器件功能发出寄存器计数值。我不相信这有一个解决方法。

但是，仍然可以通过使用cuobjdump从汇编程序输出中转储精灵数据来获得__device__函数的寄存器占用空间。你可以这样做如下：

$ 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寄存器。这是我知道检查器件功能寄存器占位情况的唯一方法。