ROCm/HIP项目调试指南：从基础到高级技巧-优快云博客

本文链接：https://blog.youkuaiyun.com/gitblog_00829/article/details/148549834

ROCm/HIP项目调试指南：从基础到高级技巧

HIP HIP: C++ Heterogeneous-Compute Interface for Portability 项目地址: https://gitcode.com/gh_mirrors/hi/HIP

概述

在ROCm/HIP开发环境中，调试异构计算应用程序是一项关键技能。本文将详细介绍HIP应用程序的调试方法，包括基础工具使用、高级技巧和环境变量配置，帮助开发者快速定位和解决问题。

追踪工具ltrace的使用

ltrace是Linux系统中强大的动态库调用追踪工具，特别适合用于观察HIP应用程序的运行时行为。

基本用法

追踪HIP API调用的基本命令格式：

ltrace -C -e "hip*" ./your_hip_app

示例输出解析：

hipGetChanDesc->hipCreateChannelDesc(0x7ffdc4b66860, 32, 0, 0) = 0x7ffdc4b66860
hipGetChanDesc->hipMallocArray(0x7ffdc4b66840, 0x7ffdc4b66860, 8, 8) = 0

这显示了应用程序调用的HIP API序列及其参数和返回值。

追踪HSA运行时

要深入了解底层HSA运行时行为，可以使用：

ltrace -C -e "hsa*" ./your_hip_app

典型输出包括HSA内存管理、事件处理等底层操作：

libhsa-runtime64.so.1->hsaKmtAllocMemory(0, 4096, 64, 0x7fff325a6690) = 0
libhsa-runtime64.so.1->hsaKmtMapMemoryToGPUNodes(0x7f1202749000, 4096, 0x7fff325a6690, 0) = 0

使用ROCgdb进行调试

ROCgdb是基于GNU GDB的ROCm专用调试器，支持HIP应用程序的源码级调试。

基本调试流程

启动ROCgdb：

rocgdb ./your_hip_app

设置断点：

(gdb) break main

运行程序：

(gdb) run

继续执行：

(gdb) c

调试示例：内存错误分析

当遇到段错误时，ROCgdb可以提供详细的调用栈信息：

Thread 1 received signal SIGSEGV, Segmentation fault.
0x000000000020f78e in simpleTest2<float> (numElements=4194304, usePinnedHost=true)
    at /path/to/file.cpp:104
104             A_h1[i] = 3.14f + 1000 * i;

使用bt命令查看完整调用栈：

#0  0x000000000020f78e in simpleTest2<float> (...) at file.cpp:104
#1  0x000000000020e96c in main (...) at file.cpp:163

关键环境变量配置

HIP提供了一系列环境变量用于调试和性能分析。

内核序列化控制

内核排队序列化：
- AMD_SERIALIZE_KERNEL=1：内核排队前等待完成
- AMD_SERIALIZE_KERNEL=2：内核排队后等待完成
- AMD_SERIALIZE_KERNEL=3：两者都启用
内存拷贝序列化：
- AMD_SERIALIZE_COPY=1：拷贝操作前等待完成
- AMD_SERIALIZE_COPY=2：拷贝操作后等待完成
- AMD_SERIALIZE_COPY=3：两者都启用

设备可见性控制

在多设备系统中，可以限制HIP可见的设备：

HIP_VISIBLE_DEVICES=0,1  # 只显示设备0和1

或在代码中设置：

setenv("HIP_VISIBLE_DEVICES", "0,1,2", 1);

HSA相关变量

禁用DMA引擎：

HSA_ENABLE_SDMA=0  # 使用计算着色器进行内存拷贝

禁用中断：

HSA_ENABLE_INTERRUPT=0  # 使用内存轮询代替中断

高级调试技巧

同步执行模式：设置AMD_SERIALIZE_KERNEL=3和AMD_SERIALIZE_COPY=3可以强制内核同步执行，便于定位异步错误。
虚拟内存错误分析：GPU内核中的VM错误可能由以下原因引起：
- 数组越界访问
- 无效的内存指针
- 同步问题
- 编译器生成的错误代码
- 运行时问题
GDB环境变量设置：在GDB会话中直接设置环境变量：
```
(gdb) set env AMD_SERIALIZE_KERNEL 3
```
多线程调试：使用info thread查看所有线程，thread <id>切换线程上下文。