MUSA Compute Sanitizer 用户指南

什么是 MUSA Compute Sanitizer？

MUSA Compute Sanitizer (mt-compute-sanitizer) 是一个面向 MUSA GPU 应用程序的运行时 Debugging 工具。

主要功能

功能	说明
越界检查	检测全局内存访问违规
泄漏检查	检测内存泄漏（需显式启用）
MUSA API 错误检查	收集 Runtime 和 Driver API 返回值

内存错误检测

越界检查

可检测以下问题：

GPU Kernel 中的全局内存读/写违规
超出分配内存边界的访问
无效指针解引用

示例：越界错误

有 Bug 的代码
修复后的代码

#include <stdio.h>
#include <stdlib.h>
#include <musa_runtime_api.h>

#define MUSA_CHECK(call) \
do { \
    musaError_t err = call; \
    if (err != musaSuccess) { \
        fprintf(stderr, "MUSA error in file '%s' line %d: %s\n", \
                __FILE__, __LINE__, musaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
} while (0)

__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    // BUG：无边界检查！
    c[i] = a[i] + b[i];
}

int main(int argc, char **argv) {
    int n = 256;
    size_t size = n * sizeof(float);

    float *h_a = (float *)malloc(size);
    float *h_b = (float *)malloc(size);
    float *h_c = (float *)malloc(size);

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = rand() / (float)RAND_MAX;
        h_b[i] = rand() / (float)RAND_MAX;
    }

    // 设备内存分配
    float *d_a, *d_b, *d_c;
    MUSA_CHECK(musaMalloc((void**)&d_a, size));
    MUSA_CHECK(musaMalloc((void**)&d_b, size));
    MUSA_CHECK(musaMalloc((void**)&d_c, size));

    // 数据拷贝
    MUSA_CHECK(musaMemcpy(d_a, h_a, size, musaMemcpyHostToDevice));
    MUSA_CHECK(musaMemcpy(d_b, h_b, size, musaMemcpyHostToDevice));

    // BUG：blockSize=257 > n=256，导致越界
    int blockSize = 257;
    int gridSize = 1;
    vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

    MUSA_CHECK(musaDeviceSynchronize());

    // 释放内存
    free(h_a); free(h_b); free(h_c);
    MUSA_CHECK(musaFree(d_a));
    MUSA_CHECK(musaFree(d_b));
    MUSA_CHECK(musaFree(d_c));
    MUSA_CHECK(musaDeviceReset());

    return 0;
}

__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    // 已修复：添加边界检查
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

MUSA Compute Sanitizer 输出：

========= COMPUTE-SANITIZER
==PROF== '/home/mt/vectorAdd' started, pid:571353
==PROF== '/home/mt/vectorAdd' exit, exit code(0)
========= Device memory out of bounds. Address: 0x1006da04c00,
           vectorAdd(float const*, float const*, float*, int)+0xfc
========= ERROR SUMMARY: 1 error

泄漏检查

泄漏检查默认关闭，需通过 --leak-check=full 显式启用。启用后可检测以下问题：

使用 musaMalloc() 分配但未释放的内存
资源泄漏（流、事件等）

可通过以下命令启用：

# 使用 --leak-check=full 启用
mt-compute-sanitizer --leak-check=full ./your_musa_program

示例：内存泄漏

有 Bug 的代码
修复后的代码

#include <stdio.h>
#include <stdlib.h>
#include <musa_runtime_api.h>

// 简化的 MUSA 错误检查宏
#define MUSA_CHECK(call)                                                        \
do {                                                                           \
    musaError_t err = call;                                                    \
    if (err != musaSuccess) {                                                  \
        fprintf(stderr, "MUSA error: %s\n", musaGetErrorString(err));          \
        exit(EXIT_FAILURE);                                                    \
    }                                                                          \
} while (0)

// MUSA 核函数：向量加法
__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    // 1. 简化的向量长度（小数据量，方便演示）
    int n = 1024;
    size_t size = n * sizeof(float);

    // 2. 主机内存分配与初始化
    float *h_a = (float *)malloc(size);
    float *h_b = (float *)malloc(size);
    float *h_c = (float *)malloc(size);
    for (int i = 0; i < n; i++) {
        h_a[i] = i;
        h_b[i] = 2*i;
    }

    // 3. 设备内存分配（关键：后续未释放）
    float *d_a, *d_b, *d_c;
    MUSA_CHECK(musaMalloc((void**)&d_a, size));  // 未释放
    MUSA_CHECK(musaMalloc((void**)&d_b, size));  // 未释放
    MUSA_CHECK(musaMalloc((void**)&d_c, size));  // 未释放

    // 4. 数据拷贝 + 核函数执行
    MUSA_CHECK(musaMemcpy(d_a, h_a, size, musaMemcpyHostToDevice));
    MUSA_CHECK(musaMemcpy(d_b, h_b, size, musaMemcpyHostToDevice));
    vectorAdd<<<4, 256>>>(d_a, d_b, d_c, n);
    MUSA_CHECK(musaDeviceSynchronize());

    // 5. 结果拷贝回主机
    MUSA_CHECK(musaMemcpy(h_c, d_c, size, musaMemcpyDeviceToHost));

    // 6. 简单验证
    printf("h_c[0] = %f (expected: 0), h_c[100] = %f (expected: 300)\n", h_c[0], h_c[100]);

    // 7. 仅释放主机内存，**遗漏设备内存释放**
    free(h_a);
    free(h_b);
    free(h_c);

    // 关键遗漏：没有调用 musaFree 释放 d_a/d_b/d_c
    // MUSA_CHECK(musaFree(d_a));
    // MUSA_CHECK(musaFree(d_b));
    // MUSA_CHECK(musaFree(d_c));

    MUSA_CHECK(musaDeviceReset());

    return 0;
}

#include <stdio.h>
#include <stdlib.h>
#include <musa_runtime_api.h>

// 简化的 MUSA 错误检查宏
#define MUSA_CHECK(call)                                                        \
do {                                                                           \
    musaError_t err = call;                                                    \
    if (err != musaSuccess) {                                                  \
        fprintf(stderr, "MUSA error: %s\n", musaGetErrorString(err));          \
        exit(EXIT_FAILURE);                                                    \
    }                                                                          \
} while (0)

// MUSA 核函数：向量加法
__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    // 1. 简化的向量长度（小数据量，方便演示）
    int n = 1024;
    size_t size = n * sizeof(float);

    // 2. 主机内存分配与初始化
    float *h_a = (float *)malloc(size);
    float *h_b = (float *)malloc(size);
    float *h_c = (float *)malloc(size);
    for (int i = 0; i < n; i++) {
        h_a[i] = i;
        h_b[i] = 2*i;
    }

    // 3. 设备内存分配（关键：后续未释放）
    float *d_a, *d_b, *d_c;
    MUSA_CHECK(musaMalloc((void**)&d_a, size));
    MUSA_CHECK(musaMalloc((void**)&d_b, size));
    MUSA_CHECK(musaMalloc((void**)&d_c, size));

    // 4. 数据拷贝 + 核函数执行
    MUSA_CHECK(musaMemcpy(d_a, h_a, size, musaMemcpyHostToDevice));
    MUSA_CHECK(musaMemcpy(d_b, h_b, size, musaMemcpyHostToDevice));
    vectorAdd<<<4, 256>>>(d_a, d_b, d_c, n);
    MUSA_CHECK(musaDeviceSynchronize());

    // 5. 结果拷贝回主机
    MUSA_CHECK(musaMemcpy(h_c, d_c, size, musaMemcpyDeviceToHost));

    // 6. 简单验证
    printf("h_c[0] = %f (expected: 0), h_c[100] = %f (expected: 300)\n", h_c[0], h_c[100]);

    // 7. 释放主机内存
    free(h_a);
    free(h_b);
    free(h_c);

    // 8. 释放设备内存
    MUSA_CHECK(musaFree(d_a));
    MUSA_CHECK(musaFree(d_b));
    MUSA_CHECK(musaFree(d_c));

    MUSA_CHECK(musaDeviceReset());

    return 0;
}

MUSA Compute Sanitizer 输出：

========= COMPUTE-SANITIZER
==PROF== '/home/mt/vectorAdd' started, pid:574766
h_c[0] = 0.000000 (expected: 0), h_c[100] = 300.000000 (expected: 300)
==PROF== '/home/mt/vectorAdd' exit, exit code(0)

========= Leaked 4096 bytes at 0x1006da04000
========= /lib/x86_64-linux-gnu/libmusa.so.1(+0x4e03b3) [0x7f97b602a3b3]
========= /lib/x86_64-linux-gnu/libmusa.so.1(muMemAlloc_v2+0x7f3) [0x7f97b5f86c23]
========= /usr/local/musa/lib/libmusart.so.5(+0x38d99) [0x7f97b6e53d99]
========= /usr/local/musa/lib/libmusart.so.5(musaMalloc+0x308) [0x7f97b6ea5bf8]
========= /home/mt/vectorAdd.mu:39[0x401388]
========= /lib/x86_64-linux-gnu/libc.so.6(+0x29d90) [0x7f97b6c0cd90]

========= Leaked 4096 bytes at 0x1006da05000
========= /lib/x86_64-linux-gnu/libmusa.so.1(+0x4e03b3) [0x7f97b602a3b3]
========= /lib/x86_64-linux-gnu/libmusa.so.1(muMemAlloc_v2+0x7f3) [0x7f97b5f86c23]
========= /usr/local/musa/lib/libmusart.so.5(+0x38d99) [0x7f97b6e53d99]
========= /usr/local/musa/lib/libmusart.so.5(musaMalloc+0x308) [0x7f97b6ea5bf8]
========= /home/mt/vectorAdd.mu:40[0x4013e8]
========= /lib/x86_64-linux-gnu/libc.so.6(+0x29d90) [0x7f97b6c0cd90]

========= Leaked 4096 bytes at 0x1006da06000
========= /lib/x86_64-linux-gnu/libmusa.so.1(+0x4e03b3) [0x7f97b602a3b3]
========= /lib/x86_64-linux-gnu/libmusa.so.1(muMemAlloc_v2+0x7f3) [0x7f97b5f86c23]
========= /usr/local/musa/lib/libmusart.so.5(+0x38d99) [0x7f97b6e53d99]
========= /usr/local/musa/lib/libmusart.so.5(musaMalloc+0x308) [0x7f97b6ea5bf8]
========= /home/mt/vectorAdd.mu:41[0x401448]
========= /lib/x86_64-linux-gnu/libc.so.6(+0x29d90) [0x7f97b6c0cd90]

========= ERROR SUMMARY: 3 errors

MUSA API 错误检查

此功能默认启用，可检测以下问题：

MUSA Runtime API 返回值错误
MUSA Driver API 返回值错误
无效 API 使用

可通过以下命令关闭：

# 使用 --report-api-errors=no 禁用
mt-compute-sanitizer --report-api-errors=no ./your_musa_program

示例：API 错误

这个示例通过向 musaMalloc 传入异常的分配大小，演示如何定位 API 调用错误。

有 Bug 的代码
修复后的代码

int main() {
    float *d_data;
    musaError_t err = musaMalloc(&d_data, -1);  // 无效大小
    if (err != musaSuccess) {
        printf("musaMalloc failed: %d\n", err);
    }
    return 0;
}

int main() {
    float *d_data = nullptr;
    size_t size = 1024 * sizeof(float);

    musaError_t err = musaMalloc((void **)&d_data, size);
    if (err != musaSuccess) {
        printf("musaMalloc failed: %s\n", musaGetErrorString(err));
        return 1;
    }

    musaFree(d_data);
    return 0;
}

MUSA Compute Sanitizer 输出：

========= COMPUTE-SANITIZER
==PROF== '/home/mt/vectorAdd' started, pid:574601
==PROF== '/home/mt/vectorAdd' exit, exit code(0)
========= Program hit musaErrorMemoryAllocation (error 2) due to
           "out of memory" on MUSA Runtime API call to musaMalloc.
========= /lib/x86_64-linux-gnu/libmusa.so.1(+0x505d9b) [0x7f5a823c7d9b]
========= /home/mt/musa_runtime.h:794[0x40117d]
========= ERROR SUMMARY: 1 error

建议重点关注以下信息：

musaMalloc：出错的 API 名称
error 2 和 "out of memory"：错误码及错误原因
ERROR SUMMARY：本次检测到的错误数量

此示例中传入 -1 会被解释为异常大的分配请求，因此输出中表现为 out of memory。

如需先了解基本使用流程，可参考快速入门。

命令行参考

基本命令

# 打印帮助信息
mt-compute-sanitizer --help

# 使用默认检查运行（越界 + API 错误检查）
mt-compute-sanitizer ./your_program

# 启用泄漏检查运行
mt-compute-sanitizer --leak-check=full ./your_program

# 禁用 API 错误检查
mt-compute-sanitizer --report-api-errors=no ./your_program

# 启用所有检查
mt-compute-sanitizer --leak-check=full --report-api-errors=all ./your_program

选项

选项	默认值	描述
`--leak-check=<no/full>`	no	启用内存泄漏检测
`--report-api-errors=<all/no>`	all	启用 MUSA API 错误检查
`--help`	-	打印帮助信息

Debugging 工作流

系统化方法

常见错误模式

错误类型	原因	修复
越界	Kernel 中无边界检查	添加 `if (idx < n)` 检查
内存泄漏	缺少 `musaFree()`	为所有分配添加 `musaFree()`
API 错误	无效参数	检查 API 文档

功能范围

检测能力

功能	说明
全局内存越界访问	检测 Kernel 中的全局内存读/写违规
内存泄漏	检测使用 `musaMalloc()` 分配但未释放的内存（需 `--leak-check=full`）
MUSA API 错误	检测 Runtime/Driver API 返回值错误

已知限制

大型应用程序的性能开销可能很大

附录 A：错误代码参考

内存错误

错误	描述	示例
越界	访问超出分配内存	`array[100]` 当大小为 50
内存泄漏	已分配但未释放	`musaMalloc()` 无 `musaFree()`

API 错误

错误	描述	常见原因
invalid argument	传递无效参数	Size = 0, NULL 指针
out of memory	GPU 内存耗尽	分配过多
invalid device	无效 GPU 设备	设备未初始化

附录 B：集成工作流

MUSA Compute Sanitizer + Moore Perf Compute 工作流

warning

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

什么是 MUSA Compute Sanitizer？​

主要功能​

内存错误检测​

越界检查​

示例：越界错误​

泄漏检查​

示例：内存泄漏​

MUSA API 错误检查​

示例：API 错误​

命令行参考​

基本命令​

选项​

Debugging 工作流​

系统化方法​

常见错误模式​

功能范围​

检测能力​

已知限制​

附录 A：错误代码参考​

内存错误​

API 错误​

附录 B：集成工作流​

MUSA Compute Sanitizer + Moore Perf Compute 工作流​

这篇文章对您有帮助吗？

什么是 MUSA Compute Sanitizer？

主要功能

内存错误检测

越界检查

示例：越界错误

泄漏检查

示例：内存泄漏

MUSA API 错误检查

示例：API 错误

命令行参考

基本命令

选项

Debugging 工作流

系统化方法

常见错误模式

功能范围

检测能力

已知限制

附录 A：错误代码参考

内存错误

API 错误

附录 B：集成工作流

MUSA Compute Sanitizer + Moore Perf Compute 工作流