使用 MUSA Mapping 在编译期实现 CUDA 兼容

什么是 MUSA Mapping？

MUSA Mapping（安装路径 /usr/local/musa/tools/musaMapping）是 MUSA Toolkits 提供的 编译期 CUDA 兼容层。它在 mcc 编译过程中加载 Clang 前端插件 libMusaMapping.so，对源码进行预处理与 AST 重写，将 CUDA 风格的：

#include 头文件路径（如 cuda_runtime.h → musa_runtime.h）
API / 宏 / 类型标识符（如 cudaMalloc → musaMalloc，CUBLAS → MUBLAS）
部分生态专用头（PyTorch、cuDNN、CuTe/CUTLASS 等）

自动替换为 MUSA 对应符号，源文件仍保持 CUDA 写法。

与 musify 的区别

维度	musify	MUSA Mapping
时机	编译前，离线改源文件	编译时，不改磁盘上的源文件
实现	Python 文本匹配（`musify-text`）	Clang 插件（语义感知重写）
典型场景	一次性迁移、可审查 diff	直接编译未改的 `.cu`、在现有 nvcc 构建里接入 `mcc`、大型第三方工程
输出	生成 MUSA 风格源码	生成目标文件 / 可执行文件

两者可配合使用：先用 musify 做粗迁移，再用 MUSA Mapping 编译仍含 CUDA 名的依赖或宏。

工作原理

  CUDA 风格源码 (.cu / .c / .h)
           │
           ▼
  mcc + libMusaMapping.so (-fplugin, -x musa)
           │
           ├── 读取 mapping/*.json、custom_defines.h
           ├── 重写 #include、标识符、宏
           └── 编译为 MUSA 设备 / Host 目标代码
           │
           ▼
  链接 libmusart、libmusa 等 → 可执行文件

插件在编译时会：

根据 MUSAMAPPING_PATH 加载映射表；
在预处理阶段处理 #include 等（MusaMappingPPCallback）；
在 AST 阶段做标识符替换（MusaMappingASTConsumer）；
注入 custom_defines.h 中的宏定义，保证宏展开前已完成 CUDA→MUSA 映射。

安装与目录结构

安装 MUSA SDK 后，工具位于 /usr/local/musa/tools/musaMapping/。插件完整路径为：

/usr/local/musa/tools/musaMapping/libMusaMapping.so

目录结构如下：

/usr/local/musa/tools/musaMapping/
├── libMusaMapping.so      # Clang 插件（核心）
├── mcc_wrapper            # 兼容 nvcc 调用方式的包装脚本
├── custom_defines.h       # 宏映射表（与 JSON 对应，约 3 万行）
├── mapping/
│   ├── general.json       # 标识符 / 宏（约 3 万条）
│   ├── include.json       # CUDA SDK 头文件映射
│   ├── torch-include.json # PyTorch CUDA 头 → torch_musa
│   ├── dnn-header.json    # cuDNN 等
│   └── other-include.json # CuTe/CUTLASS（cute/ → mute/）等
└── cmake/Modules/         # CMake 兼容模块（可将 CUDA 库目标重定向到 MUSA，见下文限制）

前提条件：

已安装 MUSA SDK（mcc、libmusart.so、MUSA 头文件）；
编译时需能加载 libMusaMapping.so（与 mcc 同套 LLVM/Clang 构建）。

快速开始

以下示例在 不修改 CUDA 源码 的前提下，用 MUSA Mapping 编译并运行一个向量 +1 内核。

示例源码 `add_one.cu`

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void add_one(float* data, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        data[i] += 1.0f;
    }
}

int main(void) {
    const int n = 4;
    float h[4] = {0.f, 1.f, 2.f, 3.f};
    float* d_data = NULL;

    if (cudaMalloc((void**)&d_data, n * sizeof(float)) != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed\n");
        return 1;
    }
    if (cudaMemcpy(d_data, h, n * sizeof(float), cudaMemcpyHostToDevice) != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy H2D failed\n");
        return 1;
    }

    const int threads = 256;
    const int blocks = (n + threads - 1) / threads;
    add_one<<<blocks, threads>>>(d_data, n);
    if (cudaGetLastError() != cudaSuccess) {
        fprintf(stderr, "kernel launch failed\n");
        return 1;
    }

    if (cudaMemcpy(h, d_data, n * sizeof(float), cudaMemcpyDeviceToHost) != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy D2H failed\n");
        return 1;
    }
    cudaFree(d_data);
    cudaDeviceSynchronize();

    printf("result:");
    for (int i = 0; i < n; ++i) {
        printf(" %.1f", h[i]);
    }
    printf("\n");
    return 0;
}

编译与运行

export MUSAMAPPING_PATH=/usr/local/musa/tools/musaMapping
export MUSA_INCLUDE_PATH=/usr/local/musa/include

mcc -x musa -I. \
  -fplugin=$MUSAMAPPING_PATH/libMusaMapping.so \
  --offload-arch=mp_22 \
  -DCCCL_DISABLE_NVFP8_SUPPORT \
  -I/usr/local/musa/include \
  add_one.cu \
  -L/usr/local/musa/lib -Wl,--no-as-needed -lmusart -lmusa \
  -o add_one

./add_one

预期输出：

result: 1.0 2.0 3.0 4.0

--offload-arch 须与目标 GPU 一致（可用 musaInfo 查看 compute capability，例如 2.2 对应 mp_22，3.1 对应 mp_31）。下文示例沿用 mp_22，请按实际硬件替换。

不要把插件当成运行时库

libMusaMapping.so 只在 编译期 通过 -fplugin 加载。链接阶段请使用 -lmusart（建议同时加 -lmusa），不要把 libMusaMapping.so 写入 target_link_libraries。用到 cuSPARSE 时再加 -lmusparse。

mcc 默认开启链接器 --as-needed，且 libmusart.so 不会把 libmusa 写进 DT_NEEDED。因此只写 -lmusart 时，可执行文件往往不会拉上 libmusa / libpthread；在部分发行版上，首次调用 musaMalloc / cudaMalloc 会在驱动初始化阶段段错误。这不是 Mapping 插件的编译参数，也不是 mcc 的必选开关。按 mcc 编译器的说明，用 -Wl,--no-as-needed -lmusart -lmusa 即可（libmusa 会间接带上 pthread）。显式加 -pthread 也能规避，但不是 mcc 文档要求的必备参数。

Makefile 示例

MUSAMAPPING_PATH ?= /usr/local/musa/tools/musaMapping
MCC ?= /usr/local/musa/bin/mcc

PLUGIN := -fplugin=$(MUSAMAPPING_PATH)/libMusaMapping.so
MUSA_FLAGS := -x musa -I. $(PLUGIN) --offload-arch=mp_22 -DCCCL_DISABLE_NVFP8_SUPPORT
MAPPING_ENV := MUSAMAPPING_PATH=$(MUSAMAPPING_PATH) \
               MUSA_INCLUDE_PATH=/usr/local/musa/include

add_one: add_one.cu
	$(MAPPING_ENV) $(MCC) $(MUSA_FLAGS) -I/usr/local/musa/include add_one.cu \
	  -L/usr/local/musa/lib -Wl,--no-as-needed -lmusart -lmusa -o $@

关于编译日志中的提示

编译时可能出现类似 add_one.cu:2:10 的简短提示，表示插件正在处理 #include <cuda_runtime.h> 的映射，一般不影响编译结果。

编译方式

1. 直接使用 mcc（推荐用于手写工程）

必备参数：

参数 / 环境变量	说明
`-fplugin=.../libMusaMapping.so`	加载 MUSA Mapping 插件
`-x musa`	按 MUSA 语言模式编译
`-I.`	当前目录，便于插件解析 `custom_defines.h`
`MUSAMAPPING_PATH`	映射表根目录（`mapping/`、`custom_defines.h`）
`MUSA_INCLUDE_PATH`	MUSA 头文件根目录
`--offload-arch=mp_XX`	目标 GPU 架构（按实际硬件选择，如 `mp_22` / `mp_31`）
链接 `-lmusart`（建议加 `-lmusa`）	运行时库；不要链接 `libMusaMapping.so`。`mcc` 默认 `--as-needed`，建议 `-Wl,--no-as-needed -lmusart -lmusa`

export MUSAMAPPING_PATH=/usr/local/musa/tools/musaMapping
export MUSA_INCLUDE_PATH=/usr/local/musa/include

mcc -x musa -I. -fplugin=$MUSAMAPPING_PATH/libMusaMapping.so \
    --offload-arch=mp_22 \
    -I/usr/local/musa/include \
    -c kernel.cu -o kernel.o

警告

普通 -c 编译路径下，需显式传入 -fplugin 与 -x musa。仅调用 mcc_wrapper 而不带 -fatbin/-ptx 时，包装脚本不会自动注入插件。

2. 使用 mcc_wrapper（兼容 nvcc 调用约定）

mcc_wrapper 用于替代构建系统中的 nvcc：过滤部分 CUDA 专有参数，并在 -fatbin / -ptx 场景下自动追加插件与 -x musa。

/usr/local/musa/tools/musaMapping/mcc_wrapper -fatbin -c kernel.cu -o kernel.fatbin

包装脚本还会：

过滤 --extended-lambda、--expt-relaxed-constexpr 等 MUSA 暂不支持的 nvcc 参数；
在编译设备代码时按设备信息追加 --offload-arch（如 mp_22、mp_31，以实际探测为准）；
为 torch_musa、mutlass 等生态追加常用 include 路径（可按部署环境调整脚本）。

警告

把 CMAKE_CXX_COMPILER 设为 mcc_wrapper 时，请保持 CMAKE_C_COMPILER 为 gcc/cc 等真正的 C 编译器，不要设成 mcc。mcc 会按 C++ 模式处理 .c，导致 CMake 的 C 编译器检测失败。含 MPI 的工程用 mcc_wrapper 做 C++ 编译器时，FindMPI 的 try_compile 也可能失败，需要为 MPI 单独指定头文件与库路径。

3. CMake 集成概要

对存量 CUDA 工程，优先让 .cu 走与上文相同的 mcc -x musa -fplugin=... 命令，主机 C/C++ 继续使用 gcc/g++。不要启用系统 CMake 的 CUDA 语言，以免注入 nvcc 风格的 -x cu。

cmake_minimum_required(VERSION 3.13)
project(my_app LANGUAGES CXX)

set(MUSAMAPPING_PATH /usr/local/musa/tools/musaMapping)
set(MUSA_INCLUDE_PATH /usr/local/musa/include)

add_custom_command(
  OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/main.cu.o
  COMMAND ${CMAKE_COMMAND} -E env
    MUSAMAPPING_PATH=${MUSAMAPPING_PATH}
    MUSA_INCLUDE_PATH=${MUSA_INCLUDE_PATH}
    /usr/local/musa/bin/mcc -x musa -I.
      -fplugin=${MUSAMAPPING_PATH}/libMusaMapping.so
      --offload-arch=mp_22
      -DCCCL_DISABLE_NVFP8_SUPPORT
      -I${MUSA_INCLUDE_PATH}
      -c ${CMAKE_CURRENT_SOURCE_DIR}/main.cu
      -o ${CMAKE_CURRENT_BINARY_DIR}/main.cu.o
  DEPENDS main.cu
  WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
  VERBATIM)

add_executable(my_app ${CMAKE_CURRENT_BINARY_DIR}/main.cu.o)
set_target_properties(my_app PROPERTIES LINKER_LANGUAGE CXX)
target_link_directories(my_app PRIVATE /usr/local/musa/lib)
target_link_options(my_app PRIVATE -Wl,--no-as-needed)
target_link_libraries(my_app musart musa)

若构建系统把编译器写成 nvcc，可将 nvcc 替换为 mcc_wrapper，但普通 -c 仍须在编译命令中看到 -fplugin 与 -x musa。

SDK 还提供 musaMapping/cmake/Modules，配置时可传入：

cmake -S . -B build \
  -DCMAKE_MODULE_PATH=/usr/local/musa/tools/musaMapping/cmake/Modules \
  -DCMAKE_CXX_COMPILER=/usr/local/musa/tools/musaMapping/mcc_wrapper \
  -DCMAKE_C_COMPILER=gcc

这些模块用于把 CUDA::cudart 等目标重定向到 libmusart.so 等 MUSA 库，并可能改写 CUDA 语言探测。请以 实际编译命令 为准：系统 CMake 的 project(... LANGUAGES CUDA) 常会传入 -x cu；此时 --offload-arch=mp_XX 会被当成 CUDA 架构而失败。set(CMAKE_CUDA_STANDARD 11) 暂不生效，需要指定语言标准时请用 CMAKE_CXX_STANDARD 或显式编译参数。

含 #include <cuda_runtime.h> 的 主机 .cpp 若只用 g++ 编译、未加载插件，也会找不到 CUDA 头。这类文件同样需要走 Mapping（或改为包含 MUSA 头）。

映射规则

命名规律（general.json / custom_defines.h）

CUDA 前缀 / 模式	MUSA 对应示例
`cuda` / `CUDA`	`musa` / `MUSA`
`cu*`（Driver API）	`mu*`
`cublas*`	`mublas*`
`cudnn*`	`mudnn*`
`nccl*`	`mccl*`
`nv` / `NV`	`mt` / `MT`
`cute/`（CuTe）	`mute/`

分文件职责

文件	条数（约）	用途
`mapping/general.json`	30,000+	宏、函数名、类型名等标识符
`mapping/include.json`	100+	`cuda_runtime.h` → `musa_runtime.h` 等 SDK 头
`mapping/torch-include.json`	16	`ATen/cuda/` → `ATen/musa/`、`torch_musa` 路径
`mapping/dnn-header.json`	4	`cudnn.h` → `mudnn.h` 等
`mapping/other-include.json`	190+	CUTLASS/CuTe 头路径
`custom_defines.h`	与 general 同步	以 `#define` 形式在编译期注入

扩展映射时，应同时维护 JSON 与 custom_defines.h（若涉及宏），并重新部署到 tools/musaMapping。

环境变量

变量	说明
`MUSAMAPPING_PATH`	映射工具根目录，默认 `/usr/local/musa/tools/musaMapping`
`MUSA_HOME`	SDK 根目录，默认 `/usr/local/musa`；`mcc_wrapper` 用其定位插件路径

典型使用场景

第三方 CUDA 工程：对 .cu 使用 mcc -x musa -fplugin=...；若构建系统调用 nvcc，可替换为 mcc_wrapper，但普通 -c 仍须显式传入插件与 -x musa；源码保持 CUDA 命名；
PyTorch 扩展 / torch_musa：torch-include.json 将 c10/cuda、ATen/cuda 等映射到 MUSA 实现；
推理框架（如 vLLM、SGLang）：通过包装脚本过滤 Python 扩展相关编译参数；
CUTLASS / mutlass：other-include.json 将 cute/ 映射为 mute/；
与 musify 组合：musify 处理业务代码，Mapping 编译仍含 CUDA 符号的依赖或生成代码。

局限性与注意事项

项目	说明
非运行时模拟	仅做源码级映射，不保证所有 CUDA API 在 MUSA 上语义完全一致
映射表覆盖	未收录的 API / 头文件需自行扩展 JSON 或改源码
架构参数	`--offload-arch` 须与目标 GPU 匹配；可用 `musaInfo` 查看 compute capability
CMake CUDA 语言	系统 CMake 的 `project(... LANGUAGES CUDA)` 常注入 `-x cu`，与 Mapping 所需的 `-x musa` 冲突；更稳妥的做法是主机用 `gcc`/`g++`，`.cu` 显式调用 `mcc`
CMake CUDA 标准	`CMAKE_CUDA_STANDARD` 暂不生效，请改用显式编译参数或项目 C++ 标准配置
插件不是运行时库	`libMusaMapping.so` 仅通过 `-fplugin` 加载；链接请使用 `-lmusart`，建议同时 `-lmusa`，并按 mcc 文档处理 `--as-needed`
主机文件中的 CUDA 头	未加载插件的 `g++` 编译若 `#include <cuda_runtime.h>`，会找不到该头文件；主机侧同样需要 Mapping，或改为包含 MUSA 头
FP8 等特性	当前常配合 `-DCCCL_DISABLE_NVFP8_SUPPORT` 使用
插件与 mcc 版本	`libMusaMapping.so` 须与 SDK 自带 `mcc`/Clang 版本匹配
无插件源码	发行包仅含二进制插件，逻辑修改需在 musa-mapping 源码工程中完成

故障排除

Q1: 找不到 cuda_runtime.h / 未发生头文件映射

确认已加载插件并设置路径：

# 必须同时存在
-fplugin=$MUSAMAPPING_PATH/libMusaMapping.so
-x musa
export MUSAMAPPING_PATH=/usr/local/musa/tools/musaMapping
export MUSA_INCLUDE_PATH=/usr/local/musa/include

若仍失败：检查 mapping/include.json 是否包含对应头、musa_runtime.h 是否存在于 /usr/local/musa/include。若报错来自主机 .cpp（未加 -fplugin 的 g++），该编译单元也会找不到 cuda_runtime.h。

Q2: 使用了 mcc_wrapper 但仍报 CUDA 符号未定义

mcc_wrapper 仅在 -fatbin / -ptx 时自动注入插件。普通 -c 请改用完整 mcc 命令（见上文「直接使用 mcc」），或核对实际编译命令中是否已包含 -fplugin 与 -x musa。

Q3: unsupported CUDA gpu architecture: mp_XX

未加 -x musa 时（包括系统 CMake 注入了 -x cu），mcc 会按 CUDA 架构解析 --offload-arch。请使用本文推荐的 mcc -x musa ... 组合，并指定与硬件匹配的 MUSA 架构（如 mp_22、mp_31）。

Q4: 链接阶段 undefined reference to cuda* / 混用 C/C++

确认链接 -lmusart（建议加 -lmusa）及 -L/usr/local/musa/lib，不要链接 libMusaMapping.so；
若程序在 musaMalloc / cudaMalloc 处段错误，优先按 mcc 文档加上 -Wl,--no-as-needed -lmusa（必要时再加 -pthread），而不是把它当成 Mapping 的编译开关；
用到 cuSPARSE 时增加 -lmusparse；
多文件工程注意 extern "C" 与 .cu/.c 混编时的符号修饰；
单文件 .cu 示例可避免 Host/Device 分文件链接问题。

Q5: 与 musify 如何选择？

希望仓库里长期保存 MUSA 源码、做 Code Review → 优先 musify；
希望零改动编译上游 CUDA 工程、或在现有 nvcc 构建中接入 → 优先 MUSA Mapping；
大型迁移 → musify 改业务代码 + Mapping 编译未改依赖。

Q6: CMake 检测 C 编译器失败，或 FindMPI 失败

请勿将 CMAKE_C_COMPILER 设为 mcc。C 使用 gcc，.cu 使用带插件的 mcc。mcc_wrapper 作为 C++ 编译器时，可为 MPI 显式指定 include 与 library，避免依赖 FindMPI 的 try_compile。

使用 MUSA Mapping 在编译期实现 CUDA 兼容

什么是 MUSA Mapping？

工作原理

安装与目录结构

快速开始

示例源码 `add_one.cu`

编译与运行

Makefile 示例

编译方式

1. 直接使用 mcc（推荐用于手写工程）

2. 使用 mcc_wrapper（兼容 nvcc 调用约定）

3. CMake 集成概要

映射规则

命名规律（general.json / custom_defines.h）

分文件职责

环境变量

典型使用场景

局限性与注意事项

故障排除

相关文档

这篇文章对您有帮助吗？

什么是 MUSA Mapping？​

工作原理​

安装与目录结构​

快速开始​

示例源码 add_one.cu​

编译与运行​

Makefile 示例​

编译方式​

1. 直接使用 mcc（推荐用于手写工程）​

2. 使用 mcc_wrapper（兼容 nvcc 调用约定）​

3. CMake 集成概要​

映射规则​

命名规律（general.json / custom_defines.h）​

分文件职责​

环境变量​

典型使用场景​

局限性与注意事项​

故障排除​

相关文档​

这篇文章对您有帮助吗？

什么是 MUSA Mapping？

工作原理

安装与目录结构

快速开始

示例源码 `add_one.cu`

编译与运行

Makefile 示例

编译方式

1. 直接使用 mcc（推荐用于手写工程）

2. 使用 mcc_wrapper（兼容 nvcc 调用约定）

3. CMake 集成概要

映射规则

命名规律（general.json / custom_defines.h）

分文件职责

环境变量

典型使用场景

局限性与注意事项

故障排除

相关文档