当前位置： 首页 > news >正文

CUDA-GDB(8)——检查程序状态

news 2025/10/28 14:46:15

安全二次方(security²)博客目录导读

目录

8.1. 内存与变量

8.2. 变量存储与访问性

8.3. CUDA信息命令

8.3.1. info cuda 设备

8.3.2. cuda sms信息

8.3.3. cuda线程束信息

8.3.4. cuda通道信息

8.3.5. cuda内核信息

8.3.6. CUDA块信息

8.3.7. info cuda 线程

8.3.8. cuda启动追踪信息

8.3.9. cuda启动子进程信息

8.3.10. CUDA上下文信息

8.3.11. cuda托管内存信息

8.4. 反汇编

8.5. 寄存器

8.6. 常量存储体

8.1. 内存与变量

GDB的print命令已扩展为能够解析任何程序变量的位置，可用于显示包括以下在内的任何CUDA程序变量的内容：

  • 通过 cudaMalloc() 分配的数据

  • 驻留在各种GPU内存区域中的数据，例如共享内存、本地内存和全局内存

  • 特殊的CUDA运行时变量，例如 threadIdx

8.2. 变量存储与访问性

根据变量类型和使用情况，变量可以存储在寄存器中，也可以存储在local、shared、const或global内存中。您可以打印任何变量的地址来查找其存储位置，并直接访问关联的内存。

下面的示例展示了如何直接访问类型为shared int *的变量数组，以查看数组中存储的值。

(cuda-gdb) print &array
$1 = (@shared int (*)[0]) 0x20
(cuda-gdb) print array[0]@4
$2 = {0, 128, 64, 192}

你也可以访问共享内存中起始偏移量的索引，查看存储的值是什么：

(cuda-gdb) print *(@shared int*)0x20
$3 = 0
(cuda-gdb) print *(@shared int*)0x24
$4 = 128
(cuda-gdb) print *(@shared int*)0x28
$5 = 64

以下示例展示了如何访问内核输入参数的起始地址。

(cuda-gdb) print &data
$6 = (const @global void * const @parameter *) 0x10
(cuda-gdb) print *(@global void * const @parameter *) 0x10
$7 = (@global void * const @parameter) 0x110000</>

8.3. CUDA info命令

这些命令用于显示有关GPU和应用程序CUDA状态的信息。可用选项包括：

devices

关于所有设备的信息

sms

关于当前设备中所有活跃SM的信息

warps

关于当前SM中所有活跃warp的信息

lanes

关于当前warp中所有active通道的信息

kernels

关于所有active内核的信息

blocks

关于当前内核中所有active block的信息

threads

关于当前内核中所有acvtive线程的信息

launch trace

关于当前焦点内核的父内核信息

launch children

关于由焦点内核启动的内核信息

contexts

关于所有上下文的信息

可以对每个info cuda命令应用过滤器。过滤器会限制该命令的作用范围。一个过滤器由一个或多个限制条件组成。限制条件可以是以下任意一种：

  • device n

  • sm n

  • warp n

  • lane n

  • kernel n

  • grid n

  • block x[,y] 或 block (x[,y])

  • thread x[,y[,z]] 或 thread (x[,y[,z]])

  • breakpoint all 和 breakpoint n

其中n, x, y, z是整数，或者是以下特殊关键字之一：current, any和all。current表示应使用当前焦点中的对应值。any和all表示可以接受任何值。

注意

breakpoint all 和 breakpoint n 过滤器仅对 info cuda threads 命令有效。

8.3.1. info cuda 设备

该命令会枚举系统中所有按设备索引排序的GPU。*表示当前聚焦的设备。此命令支持过滤器，默认值为device all。如果未找到活动GPU，该命令将输出No CUDA Devices。设备在首次内核启动之前不会被判定为活动状态。

(cuda-gdb) info cuda devicesDev PCI Bus/Dev ID                Name Description SM Type SMs Warps/SM Lanes/Warp Max Regs/Lane Active SMs Mask0        06:00.0 GeForce GTX TITAN Z      GK110B   sm_35  15       64         32           256 0x000000001        07:00.0 GeForce GTX TITAN Z      GK110B   sm_35  15       64         32           256 0x00000000

8.3.2. cuda sms信息

该命令显示设备的所有SM（流式多处理器）及其上关联的活动warp（线程束）。此命令支持过滤器，默认值为device current sm all。标记为*的SM表示当前处于聚焦状态。结果按设备分组显示。

(cuda-gdb) info cuda smsSM Active Warps Mask
Device 0
* 0 0xffffffffffffffff1 0xffffffffffffffff2 0xffffffffffffffff3 0xffffffffffffffff4 0xffffffffffffffff5 0xffffffffffffffff6 0xffffffffffffffff7 0xffffffffffffffff8 0xffffffffffffffff
...

8.3.3. cuda线程束信息

该命令可让您深入一层，打印出当前关注的SM中所有线程束的信息。此命令支持过滤器，默认值为device current sm current warp all。该命令可用于显示每个线程束正在执行哪个块。

(cuda-gdb) info cuda warps
Wp /Active Lanes Mask/ Divergent Lanes Mask/Active Physical PC/Kernel/BlockIdx
Device 0 SM 0
* 0    0xffffffff    0x00000000 0x000000000000001c    0    (0,0,0)1    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)2    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)3    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)4    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)5    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)6    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)7    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)...

8.3.4. cuda通道信息

该命令显示当前关注的warp中的所有通道(线程)。此命令支持过滤器，默认值为device current sm current warp current lane all。在下面的示例中，您可以看到所有通道都位于相同的物理PC位置。该命令可用于显示每个通道执行的是哪个线程。

(cuda-gdb) info cuda lanesLn    State  Physical PC        ThreadIdx
Device 0 SM 0 Warp 0
*  0    active 0x000000000000008c   (0,0,0)1    active 0x000000000000008c   (1,0,0)2    active 0x000000000000008c   (2,0,0)3    active 0x000000000000008c   (3,0,0)4    active 0x000000000000008c   (4,0,0)5    active 0x000000000000008c   (5,0,0)6    active 0x000000000000008c   (6,0,0)7    active 0x000000000000008c   (7,0,0)8    active 0x000000000000008c   (8,0,0)9    active 0x000000000000008c   (9,0,0)10    active 0x000000000000008c  (10,0,0)11    active 0x000000000000008c  (11,0,0)12    active 0x000000000000008c  (12,0,0)13    active 0x000000000000008c  (13,0,0)14    active 0x000000000000008c  (14,0,0)15    active 0x000000000000008c  (15,0,0)16    active 0x000000000000008c  (16,0,0)...

8.3.5. cuda内核信息

该命令显示当前聚焦GPU上所有活跃的内核。它会打印每个内核的SM掩码、内核ID和网格ID，以及相关的维度和参数。内核ID在所有GPU中是唯一的，而网格ID在每个GPU内是唯一的。Parent列显示父级网格的内核ID。此命令支持过滤器，默认值为kernel all。

(cuda-gdb) info cuda kernelsKernel Parent Dev Grid Status   SMs Mask   GridDim  BlockDim      Name Args
*      1      -   0    2 Active 0x00ffffff (240,1,1) (128,1,1) acos_main parms=...

此命令还将显示通过动态并行在GPU上启动的网格。具有负网格ID的内核是从GPU启动的，而具有正网格ID的内核是从CPU启动的。

8.3.6. CUDA块信息

该命令显示当前焦点内核中所有活跃或正在运行的块。结果按内核分组显示。此命令支持过滤器，默认值为kernel current block all。默认情况下输出会被合并。

(cuda-gdb) info cuda blocksBlockIdx   To BlockIdx  Count  State
Kernel 1
*  (0,0,0)    (191,0,0)    192    running

可以按如下方式关闭合并功能，此时将显示更多关于设备和SM的信息：

(cuda-gdb) set cuda coalescing off

以下是关闭合并功能时相同命令的输出。

(cuda-gdb) info cuda blocksBlockIdx   State    Dev SM
Kernel 1
*   (0,0,0)   running   0   0(1,0,0)   running   0   3(2,0,0)   running   0   6(3,0,0)   running   0   9(4,0,0)   running   0  12(5,0,0)   running   0  15(6,0,0)   running   0  18(7,0,0)   running   0  21(8,0,0)   running   0   1...

8.3.7. info cuda 线程

该命令显示应用程序当前活跃的CUDA块和线程，以及这些块中的线程总数。同时还会显示虚拟程序计数器(PC)及相关的源文件和行号信息。结果按内核(kernel)分组显示。该命令支持过滤器，默认过滤器为kernel current block all thread all。默认情况下输出会进行如下合并：

(cuda-gdb) info cuda threadsBlockIdx ThreadIdx To BlockIdx ThreadIdx Count   Virtual PC    Filename   Line
Device 0 SM 0
* (0,0,0  (0,0,0)    (0,0,0)  (31,0,0)    32  0x000000000088f88c   acos.cu   376(0,0,0)(32,0,0)  (191,0,0) (127,0,0) 24544  0x000000000088f800   acos.cu   374...

可以按如下方式关闭合并功能，这样输出时会显示更多信息。

(cuda-gdb) set cuda coalescing off

(cuda-gdb) info cuda threadsBlockIdx  ThreadIdx  Virtual PC         Dev SM Wp Ln   Filename  Line
Kernel 1
*  (0,0,0)    (0,0,0)  0x000000000088f88c   0  0  0  0    acos.cu    376(0,0,0)    (1,0,0)  0x000000000088f88c   0  0  0  1    acos.cu    376(0,0,0)    (2,0,0)  0x000000000088f88c   0  0  0  2    acos.cu    376(0,0,0)    (3,0,0)  0x000000000088f88c   0  0  0  3    acos.cu    376(0,0,0)    (4,0,0)  0x000000000088f88c   0  0  0  4    acos.cu    376(0,0,0)    (5,0,0)  0x000000000088f88c   0  0  0  5    acos.cu    376(0,0,0)    (6,0,0)  0x000000000088f88c   0  0  0  6    acos.cu    376(0,0,0)    (7,0,0)  0x000000000088f88c   0  0  0  7    acos.cu    376(0,0,0)    (8,0,0)  0x000000000088f88c   0  0  0  8    acos.cu    376(0,0,0)    (9,0,0)  0x000000000088f88c   0  0  0  9    acos.cu    376...

注意

在合并形式中，线程必须连续才能被合并。如果某些线程当前未在硬件上运行，它们将在线程范围内产生空洞。例如，如果一个内核由2个16线程的块组成，且只有最低的8个线程处于活动状态，则将打印2个合并范围：一个范围对应块0的线程0到7，另一个范围对应块1的线程0到7。由于块0中线程8-15未运行，这两个范围无法被合并。

该命令还支持breakpoint all和breakpoint breakpoint_number作为筛选条件。前者显示命中用户设置的所有CUDA断点的线程，后者显示命中指定breakpoint_number号CUDA断点的线程。

(cuda-gdb) info cuda threads breakpoint allBlockIdx ThreadIdx         Virtual PC Dev SM Wp Ln        Filename  Line
Kernel 0(1,0,0)   (0,0,0) 0x0000000000948e58   0 11  0  0 infoCommands.cu    12(1,0,0)   (1,0,0) 0x0000000000948e58   0 11  0  1 infoCommands.cu    12(1,0,0)   (2,0,0) 0x0000000000948e58   0 11  0  2 infoCommands.cu    12(1,0,0)   (3,0,0) 0x0000000000948e58   0 11  0  3 infoCommands.cu    12(1,0,0)   (4,0,0) 0x0000000000948e58   0 11  0  4 infoCommands.cu    12(1,0,0)   (5,0,0) 0x0000000000948e58   0 11  0  5 infoCommands.cu    12(cuda-gdb) info cuda threads breakpoint 2 lane 1BlockIdx ThreadIdx         Virtual PC Dev SM Wp Ln        Filename  Line
Kernel 0(1,0,0)   (1,0,0) 0x0000000000948e58   0 11  0  1 infoCommands.cu    12

8.3.8. cuda启动追踪信息

此命令显示焦点内核的内核启动追踪。追踪中的第一个元素即为焦点内核。下一个元素是启动该内核的内核。追踪将持续进行，直到没有父内核为止。这种情况下，内核是由CPU启动的。

对于跟踪中的每个内核，该命令会打印内核在跟踪中的层级、内核ID、设备ID、网格ID、状态、内核维度、内核名称以及内核参数。

(cuda-gdb) info cuda launch traceLvl Kernel Dev Grid     Status   GridDim  BlockDim Invocation
*   0      3   0   -7     Active  (32,1,1)  (16,1,1) kernel3(c=5)1      2   0   -5 Terminated (240,1,1) (128,1,1) kernel2(b=3)2      1   0    2     Active (240,1,1) (128,1,1) kernel1(a=1)

一个已启动但未在GPU上运行的内核将显示为Pending状态。当前正在GPU上运行的内核会被标记为Active。等待再次激活的内核将显示为Sleeping。当内核终止时，它会被标记为Terminated。在少数情况下，当调试器无法确定内核是挂起还是终止时，状态会被设置为Undetermined。

该命令支持过滤器，默认值为kernel all。

注意

使用set cuda software_preemption on时，不会报告任何内核处于活动状态。

8.3.9. cuda启动子进程信息

该命令显示当前聚焦内核启动的所有未终止内核列表。对于每个内核，会显示内核ID、设备ID、网格ID、内核维度、内核名称以及内核参数。

(cuda-gdb) info cuda launch childrenKernel Dev Grid GridDim BlockDim Invocation
*      3   0   -7 (1,1,1)  (1,1,1) kernel5(a=3)18   0   -8 (1,1,1) (32,1,1) kernel4(b=5)

该命令支持过滤器，默认值为kernel all。

8.3.10. CUDA上下文信息

该命令会枚举所有GPU上正在运行的CUDA上下文。*标记表示当前处于焦点状态的上下文。此命令可显示某个上下文当前是否在设备上处于活动状态。

(cuda-gdb) info cuda contextsContext Dev    State0x080b9518   0 inactive
* 0x08067948   0   active

8.3.11. cuda托管内存信息

此命令根据当前焦点显示设备或主机上的所有静态托管变量。

(cuda-gdb) info cuda managed
Static managed variables on device 0 are:
managed_var = 3
managed_consts = {one = 1, e = 2.71000004, pi = 3.1400000000000001}

8.4. 反汇编

可以使用标准的GDB反汇编指令（如x/i和display/i）来反汇编设备的SASS代码。

(cuda-gdb) x/4i $pc-320xa689a8 <acos_main(acosParams)+824>: MOV R0, c[0x0][0x34]0xa689b8 <acos_main(acosParams)+840>: MOV R3, c[0x0][0x28]0xa689c0 <acos_main(acosParams)+848>: IMUL R2, R0, R3
=> 0xa689c8 <acos_main(acosParams)+856>: MOV R0, c[0x0][0x28]

注意

要使反汇编指令正常工作，必须安装cuobjdump并将其包含在您的$PATH环境变量中。

在反汇编视图中，当前程序计数器(pc)会以=>为前缀。对于Maxwell(SM 5.0)及更新的架构，如果指令触发异常，则会以*>为前缀。如果pc和errorpc是同一指令，则会以*=>为前缀。

例如，考虑以下异常：

CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x555555c08620 (memexceptions_kernel.cu:17)Thread 1 "memexceptions" received signal CUDA_EXCEPTION_14, Warp Illegal Address.
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
0x0000555555c08fb0 in exception_kernel<<<(1,1,1),(1,1,1)>>> (data=0x7fffccc00000, exception=MMU_FAULT) at memexceptions_kernel.cu:50
50  }
(cuda-gdb)

disas 命令可用于查看触发异常的PC和错误PC。

(cuda-gdb) disas $pc,+16
Dump of assembler code from 0x555555c08fb0 to 0x555555c08fc0:
=> 0x0000555555c08fb0 <_Z16exception_kernelPv11exception_t+3504>:  ERRBAR
End of assembler dump.

(cuda-gdb) disas $errorpc,+16
Dump of assembler code from 0x555555c08620 to 0x555555c08630:
*> 0x0000555555c08620 <_Z16exception_kernelPv11exception_t+1056>:  ST.E.U8.STRONG.SYS [R6.64], R5
End of assembler dump.

8.5. 寄存器

可以使用标准的GDB命令（如info registers）来检查/修改设备寄存器代码。

(cuda-gdb) info registers $R0 $R1 $R2 $R3
R0             0xf0 240
R1             0xfffc48 16776264
R2             0x7800   30720
R3             0x80 128

寄存器也可以通过内置变量$R访问，例如：

(cuda-gdb) printf "%d %d\n", $R0*$R3, $R2
30720 30720

谓词寄存器和CC寄存器的值可以通过打印系统寄存器组或使用它们各自的伪名称来检查：$P0..$P6 和 $CC。

(cuda-gdb) info registers system
P0             0x1  1
P1             0x1  1
P2             0x0  0
P3             0x0  0
P4             0x0  0
P5             0x0  0
P6             0x1  1
CC             0x0  0

8.6. 常量存储体

分配在GPU内存常量地址空间中的内存驻留在称为常量存储体的二维数组中。 常量存储体标记为c[X][Y]，其中X是存储体编号，Y是偏移量。 特定存储体/偏移量对的内存地址可通过便捷函数$_cuda_const_bank(bank, offset)获取。

(cuda-gdb) disass $pc,+16
Dump of assembler code from 0x7fffd5043d40 to 0x7fffd5043d50:
=> 0x00007fffd5043d40 <_Z9acos_main10acosParams+1856>:  MOV R0, c[0x0][0xc]
End of assembler dump.
(cuda-gdb) p *$_cuda_const_bank(0x0,0xc)
$1 = 8
http://www.dtcms.com/a/537982.html

相关文章：

  • 青海网站建设公明网站建设怎么做
  • 学院网站建设策划书村镇建设年度报表登录网站
  • 西安网站seo技术外贸企业网站模板建设可以吗
  • Rust：函数栈帧 Box智能指针
  • 如何实现大模型 “边生成边显示“
  • 网站排版教程程序员 做网站 微信公众号 赚钱
  • 无人机数据 → 三维模型与光谱指数 → 多源融合特征 → 机器学习模型与机理解释 → 生态应用案例与科研论文
  • 做性的网站有哪些内容科技股有哪些股票龙头2021
  • 深圳网站建设找智恒网络网站做竞价优化
  • 计算机视觉：基于YOLOv8/YOLOv7/YOLOv6/YOLOv5的零售柜商品检测识别系统（Python+PySide6界面+训练代码）（源码+文档）✅
  • 重庆网站设计公司推荐永久免费虚拟主机
  • 软件自学网站房地产设计公司
  • 网络科普：自治系统编号
  • 网站不显示index.html北京最大的广告制作公司
  • TCP 消息分段与粘包问题的完整解决方案
  • 网站怎么运营推广电话销售管理系统
  • 邢台公司网站建设南漳网站制作
  • 度假村网站模板关键词研究工具
  • 【算法】day13 链表
  • 可以做网站的语言济南泰安网站建设公司
  • 超级工程网站建设上海中心大厦wordpress 登陆后台
  • 淮安网站定制wordpress多用户图库
  • 顺企网萍乡网站建设网站排名优化怎么弄
  • switch宝可梦传说Z-A金手指1.0.1免通信进化和持物通信进化修改
  • 校园电子商务网站建设网站制作教程网站
  • 长沙第三方网站建设公司辽宁工程建设信息网诚信库怎么填
  • 建网站需要什么软件义乌跨境电商公司前十名
  • 网站建设费 科目百度推广怎么注册账号
  • FP8013降压恒流单芯片切五路调光调色方案
  • 知识就是力量——数据库