第7章精讲：MUSA 程序调试（错误码 / musa-gdb / compute-sanitizer / 精度定位）

一、MUSA-GDB：GDB 的超集，不是另一套体系

MUSA-GDB 是摩尔线程公司基于 GNU GDB 扩展开发的专用调试器，同时支持主机和设备源代码调试。它兼容 MUSA Driver API、MUSA Runtime API 或两者组合开发的所有 MUSA 应用程序，能在异构计算环境中显著提升开发效率和问题定位精度。

理解它最重要的一点：MUSA-GDB 是 GDB 的超集——所有原生 GDB 命令仍然可用，额外新增了面向 GPU 线程/坐标的 musa 前缀命令族。这意味着你已有的 CPU 调试经验（break、run、continue、print）可以直接迁移过来，学习门槛主要集中在"如何用坐标定位到具体的 GPU 线程"这一新概念上。

二、调试环境搭建：两步走

调试环境搭建分两步：① 确保基础环境兼容；② 通过编译保留调试信息。

步骤 1：安装 MUSA 驱动

MUSA 驱动包安装方式与操作系统相关，典型命令（以 deb 包为例，路径与包名需按你的发行版与驱动版本调整，属"需按环境适配"）：

sudo dpkg -i musa_package.deb

步骤 2：编译应用（保留调试符号）

使用 mcc 编译时需增加编译选项 -g -gD，以便 MUSA-GDB 调试时获得程序调试信息（行号信息、变量地址信息等）：

mcc -g -gD test.mu -o test -lmusart

• -g：生成主机端调试信息；
• -gD：生成设备端调试信息；
• -lmusart：链接 MUSA 运行时库。

该命令生成的可执行文件能让 MUSA-GDB 精准关联源代码与设备端执行状态，为后续断点设置、变量查询提供基础。

调试命令规范

MUSA-GDB 在兼容 GNU GDB 命令体系的基础上，新增了专门针对 MUSA 应用的命令，例如 info musa（查看当前运行上下文信息）与 musa（切换当前调试焦点，后接坐标参数）：

(musa-gdb) info musa threads

与 GDB 命令一样，MUSA 命令的帮助信息也可以通过 help 获取。

三、内核调试技巧：线程定位、断点设置、状态查询

海量线程的执行差异、多层级内存的访问规则、核函数的逻辑流转，均需要具有针对性的调试方法。本节围绕"线程定位、断点设置、状态信息查询"三种方法拆解实操技巧。

1. 内核线程坐标与焦点管理

在 MUSA 应用中可能同时运行多个主机线程和大量的设备线程。为了简化应用状态的可视化与调试操作，调试命令通常会作用于当前聚焦实体：

• 在主机端，当焦点设置在某一主机线程上时，调试命令默认仅作用于该主机线程；
• 在设备端，调试焦点始终定位于粒度最小的单元——设备线程。

两套坐标体系需要区分清楚：

调试过程中，只要保持软件坐标与硬件坐标之间的一致性，二者就可同时使用、互换使用。此外，有时会使用另一个软件坐标——线程网格（grid）。网格与核函数的区别在于作用域不同：每个 GPU 上的网格 ID 是唯一的，而核函数 ID 在所有 GPU 范围内是唯一的，因此每个核函数等价于某一特定的"(线程网格, 设备)"对。

教学类比：可以把"软件坐标"理解成程序员写代码时用的逻辑编号（第几个线程、第几个块），"硬件坐标"理解成芯片物理执行单元的编号（第几个 lane、第几个 warp、第几个 MP）。调试器帮我们做的事情，就是把这两套编号系统实时对应起来。

查看当前焦点与切换焦点的命令：

(musa-gdb) musa device mpx warp lane block thread
block (0,0,0),thread (0,0,0) device 0,mpx 0,warp 0,lane 0
(musa-gdb) musa kernel thread
kernel 1, thread (0,0,0)
(musa-gdb) musa device 0 mpx 1 warp 2 lane 3
[Switching focus to MUSA kernel 1, grid 2, block (6,0,0), thread (20,0,0), device 0, mpx 1, warp 2, lane 3]

2. 断点设置与状态信息查询

MUSA-GDB 支持多种类型的断点设置，在语法上与 CPU 调试命令一致，且会自动分辨 CPU 断点和 GPU 断点。

Pending 断点机制：对于设置在设备代码上的断点，在对应内核的 ELF 文件尚未加载至设备前，该断点处于 pending 状态。待镜像装载完成，MUSA-GDB 解析后，即可进行断点调试。

断点触发后的并行特性：断点触发后，调试器会强制当前活跃的所有 GPU 线程暂停执行。需注意的是，GPU 执行模型具有高度并行性，不同线程的执行路径与调度时间存在差异，可能导致部分线程尚未执行至断点处。

断点类型：

• 符号断点：使用 break + 函数名称 在函数入口处设置断点，例如 break target_function。
• 行断点：使用 break + 行号 的方法在特定文件的指定行上设置断点，例如 break target_file.mu:100。如果断点指定的行位于模板代码中，将会创建多个断点，每个断点对应一个模板实例。

变量查询（print 命令扩展）：MUSA-GDB 扩展了 print 命令，以支持解析任意程序变量的存储位置，并显示任何 MUSA 程序变量的内容。该扩展功能覆盖三类变量和内存区域：① 通过 musaMalloc() 等接口分配的设备全局内存中的数据；② 位于 GPU 各内存区域中的变量，包括线程间共享存储器、线程私有存储器以及全局内存；③ 特殊的 MUSA 运行时内置变量，如 threadIdx、blockIdx 等线程或线程块索引变量。

(musa-gdb) print &array
$1 = (int (*)[0])0x10
(musa-gdb) print array[0]@4
$2 = {0,1,2,3}

array[0]@4 语法表示"从 array[0] 起连续打印 4 个元素"，这是 GDB 原生语法在 MUSA-GDB 中的直接复用。

Attach 功能：MUSA-GDB 同样具有 Attach 功能，允许调试器附加到一个正在运行的主机与设备程序上，无须停止该程序即可动态分析其状态：

musa-gdb -p <pid>

四、Step-by-Step 实验：axpby 核函数的完整十步调试

本节通过一个简单的 MUSA 核函数计算示例（axpby：y = tx*x + tmpg1 + tmpg2，随后 y *= y），结合断点设置、线程管理、变量查询等核心调试技巧，完整演示从编译到调试结束的全流程。以下代码结构完整，可直接编译运行（需按你的实际环境调整头文件路径与库路径）：

#include <iostream>
// 设备端全局变量
__device__ float tmpg1 = 1.0f;
__device__ float tmpg2 = 2.0f;

// 核函数：实现目标计算逻辑
__global__ void axpby(float *x, float *y, float tx) {
    float val = tx + tmpg1 + tmpg2;
    y[threadIdx.x] = val + x[threadIdx.x];
    y[threadIdx.x] *= y[threadIdx.x];
}

int main(int argc, char *argv[]) {
    const int kDataLen = 4;
    float a = 2.0f;
    float host_x[kDataLen] = {1.0, 2.0, 3.0, 4.0}; // 主机端输入数据
    float host_y[kDataLen];

    // 设备内存分配，主机到设备数据复制
    float *device_x;
    float *device_y;
    musaMalloc(&device_x, kDataLen * sizeof(float));
    musaMalloc(&device_y, kDataLen * sizeof(float));
    musaMemcpy(device_x, host_x, kDataLen * sizeof(float),
               musaMemcpyHostToDevice);

    // 启动内核函数
    axpby<<<1, kDataLen>>>(device_x, device_y, a);

    // 设备同步并复制结果
    musaDeviceSynchronize();
    musaMemcpy(host_y, device_y, kDataLen * sizeof(float),
               musaMemcpyDeviceToHost);

    // 输出结果
    for (int i = 0; i < kDataLen; ++i) {
        std::cout << "y[" << i << "] = " << host_y[i] << "\n";
    }
    // 资源释放
    musaFree(device_x);
    musaFree(device_y);

    musaDeviceReset();
    return 0;
}

手工验算（tx=a=2.0，可作为课堂"先手算后验证"的练习）：val = tx + tmpg1 + tmpg2 = 2.0 + 1.0 + 2.0 = 5.0，然后 y[i] = val + x[i]，再 y[i] *= y[i]。逐项计算：i=0 时 y=5+1=6→36；i=1 时 y=5+2=7→49；i=2 时 y=5+3=8→64；i=3 时 y=5+4=9→81。这与调试输出的 y[0]=36, y[1]=49, y[2]=64, y[3]=81 完全一致。

Step-by-Step：十步调试流程

1. 编译：添加调试符号选项

   mcc -g -gD axpby.mu -o axpby -lmusart

   编译成功后生成可执行文件 axpby。
2. 启动调试器

   musa-gdb axpby

3. 设置断点：覆盖主机端入口、设备端核函数入口、核函数关键行

   # 断点 1: 主机端 main 函数入口
   (musa-gdb) b main
   Breakpoint 1 at 0x401349: file axpby.mu, line 13.
   
   # 断点 2: 设备端内核函数 axpby 入口
   (musa-gdb) b axpby
   Function "axpby" not defined.
   Make breakpoint pending on future shared library load? (y or [n]) y
   Breakpoint 2 (axpby) pending.
   
   # 断点 3: 核函数第 8 行
   (musa-gdb) b axpby.mu:8
   Breakpoint 3 at 0x401324: file axpby.mu, line 10.
   
   # 断点 4: 核函数第 9 行
   (musa-gdb) b axpby.mu:9
   Note: breakpoint 3 also set at pc 0x401324.
   Breakpoint 4 at 0x401324: file axpby.mu, line 10.

   断点 2 提示 Function "axpby" not defined 时需确认 pending——这正好呼应"设备代码断点在 ELF 未加载前处于 pending 状态"的知识点。
4. 运行程序，命中 CPU 断点

   (musa-gdb) run
   Breakpoint 1, main (argc=1, argv=0x7fffffffe098) at axpby.mu:13
   13          const int kDataLen = 4;

5. 继续运行，命中设备断点：焦点自动切换到 GPU 线程

   (musa-gdb) continue
   [Switching focus to MUSA kernel 0 block (0,0,0), thread(0,0,0)  device 0 mp 0 warp 0 lane 0]
   MUSA thread hit Breakpoint 2, axpby (x=<optimized out>, y=<optimized out>, tx=2) at axpby.mu:7
   7           float val = tx + tmpg1 + tmpg2;

6. 查询线程信息

   (musa-gdb) info musa threads
   Kernel From-BlockIdx From-ThreadIdx To-BlockIdx To-ThreadIdx Count PC Dev MP Wp Ln Filename Line
   Kernel 0
   *    0 (0,0,0)   (0,0,0)   (0,0,0)   (4,0,0)   4   e000004008  0  0  0  0  axpby.mu 7

7. 切换线程：验证所有线程的执行一致性

   (musa-gdb) musa thread (1,0,0)
   [Switching focus to MUSA kernel 0 block (0,0,0), thread(1,0,0)  device 0 mp 0 warp 0 lane 1]

8. 运行到下一个设备断点

   (musa-gdb) continue
   MUSA thread hit Breakpoint 3, axpby (x=<optimized out>, y=<optimized out>, tx=2) at axpby.mu:8
   8           y[threadIdx.x] = val + x[threadIdx.x];

9. 输出变量值：验证基础计算是否符合预期（tx=2、tmpg1=1.0、tmpg2=2.0，预期 val=5.0）

   (musa-gdb) p val
   $1 = 5

10. 删除断点并继续运行：确认中间结果无误后放行程序完整执行

    (musa-gdb) delete breakpoints
    Delete all breakpoints? (y or n) y
    (musa-gdb) continue
    y[0] = 36
    y[1] = 49
    y[2] = 64
    y[3] = 81
    [Inferior 1 (process 479195) exited normally]

五、常见坑点与练习

1. 改错练习：将示例代码中 y[threadIdx.x] *= y[threadIdx.x] 故意改为越界写 y[threadIdx.x + 10]，使用 MUSA-GDB 设置断点观察行为差异。预期现象：程序可能不会立即崩溃，但会破坏其他内存区域——这正是 GPU 内存越界最隐蔽也最危险的地方，值得在实验课上专门演示。
2. 多线程断点练习：设置断点在核函数第 8 行，运行后依次切换 thread (0,0,0) 到 thread (3,0,0)，记录每个线程对应的 threadIdx.x 和执行到该行时 x[threadIdx.x] 的值。
3. 思考题：为什么 GPU 断点命中后需要显式识别"是哪个线程触发的"？如果不这样做会有什么风险？参考答案：不同线程可能异步先后到达同一断点地址，若不加区分地 continue，可能导致后续线程的现场被覆盖或误判为同一次命中，从而漏过某些线程的中间状态。

六、版本与兼容性

七、本课与相邻课程

• m11-l8（第 6 章）：本章调试的对象正是第 6 章教你写出来的核函数与内存管理代码
• m11-l3（第 3 章）：本章 musaMalloc / musaMemcpy / 核函数启动语法均来自第 3 章的编程模型
• m11-l5（第 8 章）：调试解决"对不对"的问题，第 8 章解决"快不快"的问题，二者是工程闭环的两端
• m11-l7（第 5 章）：调试环境搭建依赖第 5 章已完成的驱动与 SDK 安装