MUSA Compute Sanitizer 用户指南

什么是 MUSA Compute Sanitizer？

MUSA Compute Sanitizer（简称 mt-compute-sanitizer）是摩尔线程自主研发的 GPU 运行时调试工具，专为 MUSA GPU 应用程序设计。该工具可帮助开发者在程序运行时检测和定位内存错误，包括越界访问、内存泄漏和 API 调用错误等问题。MUSA Compute Sanitizer 通过详细的错误报告和调用栈信息，帮助开发者快速定位并修复问题，提高代码质量和开发效率。

主要功能

功能	默认状态	说明
global memory 越界检查	默认启用	检测 kernel 访问不属于自己的 global memory 地址。
MUSA API 错误检查	默认启用	报告 MUSA Runtime / Driver API 返回的错误。
显存泄漏检查	默认关闭	检测 musaMalloc() 后未正确 musaFree() 的显存，需要 --leak-check=full。

备注

显存泄漏检查不是默认项。排查泄漏或长稳显存增长问题时，必须显式添加 --leak-check=full。

推荐运行方式

最常用的方式是在原始测试命令前加 mt-compute-sanitizer，并使用 --log-file 保存 sanitizer 日志：

/usr/local/musa/bin/mt-compute-sanitizer \
  --log-file san.log \
  ./your_test_program --case-args

建议同时保存应用输出：

/usr/local/musa/bin/mt-compute-sanitizer \
  --log-file san.log \
  ./your_test_program --case-args \
  > app.out 2>&1

文件	内容
san.log	sanitizer 检查到的 OOB、leak、API error、错误汇总。
app.out	被测程序自己的输出、报错和测试框架日志。

环境检查

排查问题前建议记录环境信息：

mthreads-gmi
/usr/local/musa/bin/mcc --version
/usr/local/musa/bin/mt-compute-sanitizer --version
/usr/local/musa/bin/mt-compute-sanitizer -h

提报问题时至少保留 Driver、Toolkit、mt-compute-sanitizer 版本和完整命令。

内存错误检测

越界检查

global memory 越界检查默认启用。典型命令如下：

/usr/local/musa/bin/mt-compute-sanitizer \
  --log-file oob.log \
  ./your_test_program --case-args \
  > oob.out 2>&1

查看结果：

grep "ERROR SUMMARY" oob.log
grep -n "Out of bounds memory access\|ERROR SUMMARY" oob.log

代码示例：越界访问

有 Bug 的代码

#include <musa_runtime_api.h>

__global__ void oob_load_kernel(const int *input, int *output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // BUG：只保护了写 output 的范围，但读 input 时使用 idx + 1。
    // 当 idx == n - 1 时，会加载 input[n]，发生 global memory 越界读。
    if (idx < n) {
        output[idx] = input[idx + 1];
    }
}

int main() {
    const int n = 64;
    int h_input[n] = {0};
    int h_output[n] = {0};
    int *d_input = nullptr;
    int *d_output = nullptr;

    musaMalloc((void **)&d_input, n * sizeof(int));
    musaMalloc((void **)&d_output, n * sizeof(int));

    // 加载输入数据到设备侧。
    musaMemcpy(d_input, h_input, n * sizeof(int), musaMemcpyHostToDevice);

    oob_load_kernel<<<1, 64>>>(d_input, d_output, n);
    musaDeviceSynchronize();

    musaMemcpy(h_output, d_output, n * sizeof(int), musaMemcpyDeviceToHost);
    musaFree(d_input);
    musaFree(d_output);
    return 0;
}

修复后的代码

#include <musa_runtime_api.h>

__global__ void oob_load_kernel(const int *input, int *output, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // 修复：读 input[idx + 1] 时，idx 必须小于 n - 1。
    if (idx + 1 < n) {
        output[idx] = input[idx + 1];
    }
}

int main() {
    const int n = 64;
    int h_input[n] = {0};
    int h_output[n] = {0};
    int *d_input = nullptr;
    int *d_output = nullptr;

    musaMalloc((void **)&d_input, n * sizeof(int));
    musaMalloc((void **)&d_output, n * sizeof(int));
    musaMemcpy(d_input, h_input, n * sizeof(int), musaMemcpyHostToDevice);

    oob_load_kernel<<<1, 64>>>(d_input, d_output, n);
    musaDeviceSynchronize();

    musaMemcpy(h_output, d_output, n * sizeof(int), musaMemcpyDeviceToHost);
    musaFree(d_input);
    musaFree(d_output);
    return 0;
}

典型 OOB 日志片段：

========= COMPUTE-SANITIZER
========= ERROR SUMMARY: 448 errors
========= ERROR SUMMARY: 448 errors were not printed. Use --print-limit option to adjust the number of printed errors

该示例说明：错误数量很多时，summary 仍会统计总错误数；如果详细错误被截断，可通过 --print-limit 调整打印数量。

常见修复方式：

常见原因	修复建议
kernel 中没有边界检查	添加 if (idx < n) 等边界保护。
launch 配置产生多余线程	确认多余线程不会访问非法地址。
offset / stride / shape 计算错误	检查索引和指针计算。

泄漏检查

显存泄漏检查默认关闭，必须显式开启：

/usr/local/musa/bin/mt-compute-sanitizer \
  --leak-check=full \
  --log-file leak.log \
  ./your_test_program --case-args \
  > leak.out 2>&1

查看结果：

grep "ERROR SUMMARY" leak.log
grep -n "Leaked\|musaMalloc\|ERROR SUMMARY" leak.log

代码示例：显存泄漏

有 Bug 的代码

#include <musa_runtime_api.h>

int main() {
    int *data = nullptr;

    // BUG：申请了 MUSA device memory，但程序退出前没有释放。
    musaMalloc((void **)&data, 4096);

    musaDeviceSynchronize();
    return 0;
}

修复后的代码

#include <musa_runtime_api.h>

int main() {
    int *data = nullptr;

    musaMalloc((void **)&data, 4096);
    musaDeviceSynchronize();

    // 修复：每个 musaMalloc 都应有对应的 musaFree。
    musaFree(data);
    return 0;
}

典型 leak 日志片段：

========= COMPUTE-SANITIZER
========= Leaked 4096 bytes at 0x1006b600000
=========     Saved host backtrace up to driver entry point at allocation time
=========         Host Frame: main [0x1155] in your_program
========= ERROR SUMMARY: 1 error

适用场景：

• 长稳后显存持续增长。
• 循环跑 case 后显存没有回落。
• 怀疑异常路径没有释放 MUSA memory。

常见修复方式：

常见原因	修复建议
musaMalloc() 后没有 musaFree()	为每个设备内存分配添加对应释放。
提前 return	使用统一清理路径，确保所有分支释放资源。
异常路径未释放	使用 RAII 或封装资源生命周期。

MUSA API 错误检查

--report-api-errors 默认值为 all，即默认报告 MUSA API error。

代码示例：API 错误

有 Bug 的代码

#include <musa_runtime_api.h>

int main() {
    // BUG：选择不存在的 device id，会触发 musaErrorInvalidDevice。
    musaSetDevice(9999);
    return 0;
}

修复后的代码

#include <musa_runtime_api.h>

int main() {
    int count = 0;
    musaGetDeviceCount(&count);
    if (count <= 0) {
        return 1;
    }

    // 修复：只选择合法 device id。
    musaSetDevice(0);
    return 0;
}

典型 API error 日志片段：

========= COMPUTE-SANITIZER
========= Program hit musaErrorInvalidDevice (error 101) due to "invalid device ordinal" on MUSA API call to musaSetDevice.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: main [0x1149] in your_program
========= ERROR SUMMARY: 1 error

如果明确不关注 API error，可以关闭：

/usr/local/musa/bin/mt-compute-sanitizer \
  --report-api-errors=no \
  --log-file no_api_error.log \
  ./your_test_program --case-args

关闭后，被测程序自己的 API 返回值仍可能出现在应用输出中，但 sanitizer 不再把 API error 计入 ERROR SUMMARY。

命令行参考

基本命令

# 打印帮助信息
/usr/local/musa/bin/mt-compute-sanitizer --help

# 打印版本
/usr/local/musa/bin/mt-compute-sanitizer --version

# 使用默认检查运行（global memory 越界 + API 错误）
/usr/local/musa/bin/mt-compute-sanitizer --log-file san.log ./your_program

# 启用显存泄漏检查
/usr/local/musa/bin/mt-compute-sanitizer --leak-check=full --log-file leak.log ./your_program

# 禁用 API 错误检查
/usr/local/musa/bin/mt-compute-sanitizer --report-api-errors=no --log-file san.log ./your_program

常用选项

选项	默认值	说明
-h / --help	-	打印帮助信息。
-v / --version	-	打印版本信息。
--tool memcheck	memcheck	指定检查工具，当前通常使用默认 memcheck。
--log-file <file>	stdout	指定 sanitizer 日志文件，建议总是使用。
--leak-check=<no/full>	no	开启 MUSA allocation leak 检查。
--report-api-errors=<all/no>	all	控制是否报告 MUSA API error。
--print-limit <N>	100	控制最多打印多少条错误详情，0 表示不限制。
--num-callers-host <N>	0	控制 host backtrace 打印层数，0 表示不限制。
--show-backtrace=<yes/no>	yes	控制发生错误时是否显示 host backtrace。
--backtrace-short=<yes/no>	yes	控制是否显示简短 backtrace。
--strip-paths=<yes/no>	yes	控制源码/路径显示是否只保留文件名。

Debugging 工作流

系统化方法

结果判定

普通运行	sanitizer 运行	ERROR SUMMARY	建议结论
PASS	PASS	0 error	当前版本、输入、运行下 sanitizer clean。
PASS	PASS / FAIL	非 0	重点关注，可能存在 OOB、leak 或 API error。
FAIL	FAIL	任意	先分析普通测试失败，不优先归因 sanitizer。
PASS	sanitizer 异常退出	无完整 summary	收敛最小复现，保存 stdout/stderr、sanitizer log 和退出码。

运行后至少检查三项：

echo $?
grep "ERROR SUMMARY" san.log
grep -n "Out of bounds memory access\|Leaked\|Program hit musaError\|ERROR SUMMARY" san.log

注意

进程退出码不能替代 sanitizer 结果判定。ERROR SUMMARY: 0 error 只能说明当前版本、当前输入、当前运行没有检出问题；没有报错不等于程序一定没有问题。

常见错误模式

错误类型	典型关键字	常见原因	修复方向
global memory 越界	Out of bounds memory access	kernel 无边界检查、指针计算错误	检查索引、shape、stride、launch 配置。
显存泄漏	Leaked ... bytes	musaMalloc() 后缺少 musaFree()	补齐释放路径，处理异常分支。
API 错误	Program hit musaError...	无效参数、重复释放、显存不足等	查看 API 名称、错误码和 host backtrace。

v1.1.0 输出控制

控制错误详情数量

v1.1.0 默认最多打印 100 条详细错误。错误很多时，可以降低日志量：

/usr/local/musa/bin/mt-compute-sanitizer \
  --print-limit=10 \
  --log-file san.log \
  ./your_program

如果需要完整详情，可以设置为 0：

/usr/local/musa/bin/mt-compute-sanitizer \
  --print-limit=0 \
  --log-file san.log \
  ./your_program

当错误详情被截断时，summary 仍会统计错误数，并提示使用 --print-limit 调整打印数量。

控制 backtrace 和路径

/usr/local/musa/bin/mt-compute-sanitizer \
  --strip-paths=no \
  --backtrace-short=no \
  --num-callers-host=0 \
  --log-file san.log \
  ./your_program

API error 示例在关闭路径裁剪和短回溯后，会显示更完整的库路径和调用链：

========= COMPUTE-SANITIZER
========= Program hit musaErrorInvalidDevice (error 101) due to "invalid device ordinal" on MUSA API call to musaSetDevice.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame:  [0xe1a65] in /lib/x86_64-linux-gnu/libmusa.so.1
=========         Host Frame: musaSetDevice [0x72fe2] in /usr/local/musa/lib/libmusart.so.5
=========         Host Frame: main [0x1149] in /path/to/your_program
========= ERROR SUMMARY: 1 error

建议：

• 本地定位源码路径时使用 --strip-paths=no。
• 需要更多调用栈时使用 --backtrace-short=no。
• 日志过长时可以用 --show-backtrace=no 暂时关闭回溯，但上报问题时建议保留完整回溯。

功能范围与限制

检测能力

功能	说明
global memory 越界访问	检测 kernel 中的 global memory 读/写违规。
MUSA allocation leak	检测 musaMalloc() 分配但未释放的显存，需要 --leak-check=full。
MUSA API 错误	检测 Runtime / Driver API 返回值错误，默认开启。

已知限制

类型	限制或注意事项	建议
版本匹配	当前工具与 MUSA Toolkit / Driver 版本有依赖关系。	使用前确认 MUSA 版本并记录 Driver、Toolkit、sanitizer 版本。
安装顺序	需要先安装 MUSA Toolkit，再安装 mt-compute-sanitizer。	严格按安装指南执行。
指令覆盖	atomic、robust LSU、TME 等场景仍在后续计划中。	相关问题需要结合其他手段确认。
Memory 类型	private/shared/local/host pinned 等 memory 类型检测仍可能不完整。	不要只依赖 sanitizer 结论。
大规模任务	sanitizer 有额外内存占用，运行速度会明显变慢。	先将问题收敛到单 kernel、小输入、单进程场景。
性能评估	插桩会改变时序和内存布局。	不要用 sanitizer 结果做性能评估。

问题上报建议

发现问题时，至少保留：

1. 普通运行命令和输出日志。
2. sanitizer 运行命令和输出日志。
3. san.log 原始文件。
4. mthreads-gmi 输出。
5. mt-compute-sanitizer --version 输出。
6. Driver、Toolkit、compute-sanitizer 版本。
7. 是否可稳定复现，复现概率是多少。
8. 是否能缩小输入规模、batch、loop 或单测范围。

推荐日志命名：

_plain.out
_san.out
_san.log
_env.txt

附录：与 Moore Perf Compute 的关系

MUSA Compute Sanitizer 用于优先定位内存访问、显存泄漏和 API 错误；Moore Perf Compute 用于正确性问题收敛后的性能分析。

注意

不要同时运行 MUSA Compute Sanitizer 和 Moore Perf Compute。MUSA Compute Sanitizer 会引入额外开销，可能扭曲性能指标。