在device函数中，NVCC寄存器使用情况报告

Question

在device函数中，NVCC寄存器使用情况报告

3

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

ptxas info : Used N registers

我尝试使用noinline关键字，并将它们保留在与调用全局函数相对应的另一个文件中，因为我认为NVCC在内联后会报告包括调用的设备函数在内的全局函数的完整寄存器使用情况，但是什么都没有改变。只有将设备函数定义为全局函数，我才能获得有关设备函数寄存器使用情况的信息。

您有什么建议吗？

谢谢！

- Christopher23

2个回答

4

我不知道是什么时候加入的，但我的CUDA 10 cuobjdump有一个-res-usage标志，显示如下：

$ cuobjdump -res-usage .../cuda_compile_1_generated_VisualOdometry.cu.o

Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
identifier = /home/mad/automy-system/vision/src/VisualOdometry.cu

Resource usage:
 Common:
  GLOBAL:0 CONSTANT[3]:24
 Function _Z17vo_compute_systemPfS_P6float4S_jS0_S0_f:
  REG:39 STACK:32 SHARED:168 LOCAL:0 CONSTANT[0]:404 CONSTANT[2]:80 TEXTURE:0 SURFACE:0 SAMPLER:0
 Function _Z13vo_pre_filterP6float4PfPjPK5uint2iijff:
  REG:16 STACK:0 SHARED:8 LOCAL:0 CONSTANT[0]:372 TEXTURE:0 SURFACE:0 SAMPLER:0

- Hugo Maxwell

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- talonmies · Accepted Answer

据我理解，ptxas（设备汇编器）仅在链接代码时输出寄存器计数。独立的__device__函数不会被汇编器链接，它们只是被编译。因此，汇编器不会为设备函数发出寄存器计数值。我不认为有任何解决方法。

然而，仍然可以通过使用cuobjdump从汇编器输出中转储elf数据来获取__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个寄存器。这是我所知道的检查设备函数寄存器占用情况的唯一方法。

在__device__函数中，NVCC寄存器使用情况报告

在device函数中，NVCC寄存器使用情况报告