跳到主要内容

MUSA Compute Sanitizer 用户指南

什么是 MUSA Compute Sanitizer？​

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

主要功能​

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

内存错误检测​

越界检查​

可检测以下问题：

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

示例：越界错误​

#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;
}

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

示例：内存泄漏​

#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;
}

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 调用错误。

int main() {
    float *d_data;
    musaError_t err = musaMalloc(&d_data, -1);  // 无效大小
    if (err != musaSuccess) {
        printf("musaMalloc failed: %d\n", err);
    }
    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 memoryGPU 内存耗尽分配过多
invalid device无效 GPU 设备设备未初始化

附录 B：集成工作流​

MUSA Compute Sanitizer + Moore Perf Compute 工作流​

注意

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

版权所有 2024 摩尔线程，本文档受国际版权法保护。

摩尔线程和摩尔线程徽标是摩尔线程智能科技 (北京)有限责任公司的注册商标。

免责声明

本文档提供有关摩尔线程产品的信息。本文档并未授权任何知识产权的许可，并未以明示或暗示，或以禁止反言或其他方式授予任何知识产权许可。