gpt4 book ai didi

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

转载 作者:行者123 更新时间:2023-12-02 06:58:30 24 4
gpt4 key购买 nike

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

ptxas信息:使用了N个寄存器

输出中缺少

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

你有什么建议吗?

谢谢!

最佳答案

据我了解,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 个寄存器。这是我所知道的检查设备功能寄存器足迹的唯一方法。

关于cuda - __device__ 函数中的 NVCC 寄存器使用报告,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/45957072/

24 4 0
Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
广告合作:1813099741@qq.com 6ren.com