MUSIFY迁移工具实战 | 摩尔线程 GPU 学习中心

一、什么是MUSIFY？

MUSIFY是摩尔线程提供的代码迁移工具，能够将CUDA代码自动转换为MUSA代码。它通过文本替换和语义理解，将CUDA特有的API调用、头文件引用和编程模式转换为MUSA的对应实现。

对于已有大量CUDA代码的开发者来说，MUSIFY极大降低了迁移到摩尔线程平台的成本。你不需要手动逐行修改代码，MUSIFY可以批量处理整个项目目录。

1.1 MUSIFY能做什么？

将 #include <cuda_runtime.h> 替换为 #include <musa_runtime.h>
将 cudaMalloc 替换为 musaMalloc
将 cudaMemcpy 替换为 musaMemcpy
将 .cu 文件重命名为 .mu 文件
替换所有CUDA命名空间和类型为MUSA对应项
处理cuDNN → muDNN、NCCL → MCCL的API转换
支持自定义排除规则，跳过不需要转换的文件

💡

MUSIFY的工作原理

MUSIFY使用Aho-Corasick多模式匹配算法进行高效的文本替换。它维护了一份完整的CUDA→MUSA API映射表，包含数千个替换规则，覆盖了CUDA Runtime API、Driver API、cuDNN、NCCL、cuBLAS等各个库的API映射。

二、安装MUSIFY

安装系统依赖

bash

# MUSIFY需要python和ahocorapy库
sudo apt install python-is-python3 -y

# 安装Python依赖
pip install ahocorapy

安装MUSIFY工具

bash

# 方法1：通过MUSA SDK安装（如果SDK已安装，musify通常已包含）
which musify-text
# 如果输出路径说明已安装

# 方法2：从MUSA SDK包中手动安装
# musify-text 位于 /usr/local/musa/bin/musify-text

# 方法3：从源码获取
# MUSIFY包含在MUSA SDK的工具集中
# 确保 /usr/local/musa/bin 在PATH中
export PATH=/usr/local/musa/bin:$PATH

# 验证安装
musify-text --help

三、基本使用方法

3.1 单文件转换

bash

# 语法：musify-text [选项] -- 文件列表

# 转换单个.cu文件（就地修改）
musify-text --inplace -- vector_add.cu

# 转换后，文件中的CUDA API会被替换为MUSA API
# 但文件扩展名不会自动更改，需要手动重命名
mv vector_add.cu vector_add.mu

3.2 多文件批量转换

bash

# 转换多个文件
musify-text --inplace -- a.cu b.cpp c.cuh

# 批量转换目录下所有CUDA相关文件
musify-text --inplace -- $(find ./src -name '*.cu' -o -name '*.cuh' -o -name '*.cpp' -o -name '*.h')

# 更安全的方式（处理文件名中的空格）
find ./src -name '*.cu' -o -name '*.cuh' -o -name '*.cpp' -o -name '*.h' | \
  xargs musify-text --inplace --

3.3 常用选项

bash

# --inplace    : 就地修改文件（原文件会被覆盖）
# --print      : 仅打印修改后的内容，不写入文件（预览模式）
# --exclude    : 排除指定模式的替换

# 预览模式：查看转换结果但不修改文件
musify-text --print -- vector_add.cu

# 将转换结果输出到新文件
musify-text --print -- vector_add.cu > vector_add.mu

四、实际迁移案例

下面以一个完整的CUDA向量加法项目为例，演示MUSIFY的迁移过程。

4.1 原始CUDA代码

假设我们有一个CUDA项目，结构如下：

bash

cuda_project/
├── CMakeLists.txt
├── include/
│   └── utils.cuh
├── src/
│   ├── main.cpp
│   └── vector_add.cu
└── README.md

原始CUDA代码（vector_add.cu）：

cpp

// vector_add.cu — 原始CUDA代码
#include <cuda_runtime.h>
#include <stdio.h>
#include "utils.cuh"

__global__ void vectorAdd(float *a, 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() {
    int n = 1024;
    size_t size = n * sizeof(float);

    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;

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

    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);

    cudaDeviceSynchronize();
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 验证
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("CUDA error: %s\n", cudaGetErrorString(err));
    }

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    return 0;
}

4.2 执行MUSIFY转换

bash

# 第1步：备份原始项目
cp -r cuda_project/ cuda_project_backup/

# 第2步：运行MUSIFY转换所有源文件
cd cuda_project/
musify-text --inplace -- $(find . -name '*.cu' -o -name '*.cuh' -o -name '*.cpp' -o -name '*.h')

# 第3步：重命名文件扩展名
find . -name '*.cu' -exec bash -c 'mv "$0" "${0%.cu}.mu"' {} \;
find . -name '*.cuh' -exec bash -c 'mv "$0" "${0%.cuh}.muh"' {} \;

# 第4步：更新CMakeLists.txt（手动修改）
# 将 .cu → .mu, nvcc → mcc, CUDA相关的CMake变量替换

4.3 转换后的MUSA代码

MUSIFY转换后的代码（vector_add.mu）：

cpp

// vector_add.mu — MUSIFY自动转换后的代码
#include <musa_runtime.h>     // cuda_runtime.h → musa_runtime.h
#include <stdio.h>
#include "utils.muh"           // utils.cuh → utils.muh

__global__ void vectorAdd(float *a, 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() {
    int n = 1024;
    size_t size = n * sizeof(float);

    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;

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

    musaMalloc(&d_a, size);         // cudaMalloc → musaMalloc
    musaMalloc(&d_b, size);
    musaMalloc(&d_c, size);

    musaMemcpy(d_a, h_a, size, musaMemcpyHostToDevice);  // cuda → musa
    musaMemcpy(d_b, h_b, size, musaMemcpyHostToDevice);

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);

    musaDeviceSynchronize();        // cudaDeviceSynchronize → musaDeviceSynchronize
    musaMemcpy(h_c, d_c, size, musaMemcpyDeviceToHost);

    musaError_t err = musaGetLastError();    // cudaError_t → musaError_t
    if (err != musaSuccess) {                // cudaSuccess → musaSuccess
        printf("MUSA error: %s\n", musaGetErrorString(err));
    }

    musaFree(d_a);                  // cudaFree → musaFree
    musaFree(d_b);
    musaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    return 0;
}

✅

关键替换规则一览

cuda_runtime.h → musa_runtime.h
cudaMalloc → musaMalloc
cudaMemcpy → musaMemcpy
cudaFree → musaFree
cudaError_t → musaError_t
cudaSuccess → musaSuccess
cudaDeviceSynchronize → musaDeviceSynchronize
cudaMemcpyHostToDevice → musaMemcpyHostToDevice
cudnn → mudnn
nccl → mccl

五、排除规则

有些文件或代码片段不应该被转换。MUSIFY支持排除规则来处理这些情况：

5.1 排除特定文件

bash

# 排除第三方库目录
musify-text --inplace -- $(find ./src -name '*.cu' | grep -v third_party)

# 使用排除模式
musify-text --inplace --exclude "vendor/*" -- $(find . -name '*.cu')

5.2 排除特定替换模式

bash

# 某些CUDA API在MUSA中还不支持时，可以选择跳过
# 例如排除CUDA Graph相关API的替换
musify-text --inplace --exclude-pattern "cudaGraph" -- source.cu

5.3 需要手动修改的场景

MUSIFY不能处理所有情况，以下场景需要手动调整：

CMakeLists.txt / Makefile：构建系统配置文件需要手动修改编译器路径和编译选项
第三方依赖：如果项目依赖只提供CUDA版本的第三方库，需要找到MUSA版本或自行迁移
内联PTX汇编：CUDA的PTX汇编代码不能直接转换为MUSA汇编
特定CUDA特性：某些CUDA独有特性（如Cooperative Groups的高级用法）可能需要手动适配

⚠️

迁移建议

1. 务必先备份：使用 --inplace 会直接修改原文件，请先备份。
2. 先预览再执行：用 --print 选项预览结果，确认无误后再用 --inplace。
3. 逐步验证：转换后先编译确认没有语法错误，再逐步运行单元测试。
4. 关注编译警告：mcc编译时的警告通常提示了需要手动修改的地方。

六、编译与验证迁移结果

bash

# 编译迁移后的MUSA代码
mcc -O2 vector_add.mu -o vector_add_musa

# 运行验证
./vector_add_musa

# 如果是CMake项目，修改CMakeLists.txt：
# 将 enable_language(CUDA) 改为 enable_language(MUSA)（如支持）
# 或直接用mcc替代nvcc进行编译

# 简单Makefile示例
cat > Makefile << 'EOF'
CC = mcc
CFLAGS = -O2

all: vector_add

vector_add: vector_add.mu
	$(CC) $(CFLAGS) $< -o $@

clean:
	rm -f vector_add
EOF

make

✅

本课小结

MUSIFY是摩尔线程提供的CUDA代码自动迁移工具，通过 musify-text --inplace 命令可以批量将CUDA API转换为MUSA API。迁移流程为：备份 → MUSIFY转换 → 重命名文件 → 修改构建系统 → 编译验证。对于大部分CUDA项目，MUSIFY可以自动完成95%以上的代码转换工作，剩余部分需要手动微调。