mcc 编译器

1. 概述

mcc 是基于 clang 的 MUSA 编译驱动入口，负责处理同时包含 Host 代码和 Device 代码的源文件，并在需要时完成设备代码嵌入、链接和最终可执行文件生成。

MUSA 程序通常由两部分组成：

运行在 CPU 上的 Host 代码
运行在 MTGPU 上的 Device 代码

mcc 的核心目标是屏蔽这两条编译路径背后的细节，让用户尽量通过一条统一命令完成：

预处理
Host 编译
Device 编译
Device 链接
嵌入设备镜像
Host 链接

1.1 前提条件

mcc 当前运行于 Linux 环境
mcc 默认从 /usr/local/musa 查找 MUSA Toolkit
可以通过 --musa-path 指定非默认安装路径
若需要链接运行时，通常还需要 -lmusart 和 -L/usr/local/musa/lib
建议用户具备一定的 C/C++ 编程基础，了解 clang/clang++ 的基本用法

1.2 MCC 的作用

mcc 的价值不在于替代底层编译器，而在于：

为 MUSA 源文件提供统一驱动入口
管理 Host/Device 双路径编译
管理 MTGPU 目标相关参数
在常规单文件模式与分离式编译模式之间提供一致接口

2. Host 编译器

mcc 保留 clang/clang++ 的大部分通用行为，并推荐统一使用 mcc 作为入口驱动。

在 mcc 的编译模型下：

Host 侧编译器统一为 clang
用户通常不需要显式单独调用 host 编译器
普通 C/C++ 编译选项、头文件路径、库路径、链接行为，大体与 clang/clang++ 保持一致

推荐做法：

统一使用 mcc 作为 Host + Device 的入口
即便只编译 Host 部分，也优先使用 mcc

3. 编译阶段

mcc 的编译阶段采用通用的异构编译驱动模式，但后端目标和部分开关与 MTGPU/MUSA 适配。

3.1 MCC 识别宏

下表列出 mcc 路径下最基础、最稳定可见的识别宏。

宏名	说明
`__MUSACC__`	使用 `mcc` 编译源文件时定义。用于识别当前编译由 `mcc` 路径处理。
`__MUSA_ARCH__`	编译设备端源文件时定义。其值表示当前目标设备架构编号。
`__MUSART_API_VERSION`	MUSA Runtime API 版本信息可用时定义。用于表示运行时 API 版本。
`MUSART_VERSION`	在定义 Runtime API 版本信息时可用，通常作为 `__MUSART_API_VERSION` 的别名。
`MUSA_VERSION`	MUSA SDK 版本信息可用时定义。用于表示当前 SDK 版本。

补充说明：

__MUSACC__ 在 Host 和 Device 编译路径下都可见
__MUSA_ARCH__ 只在设备端编译路径下可见，其值等于当前目标架构编号
MUSART_VERSION 基本可视为 __MUSART_API_VERSION 的别名

3.2 MCC 编译阶段

mcc 支持的阶段可以类比为：

预处理
Host 侧前端编译
Device 侧前端编译
LLVM IR / 对象文件生成
Device 代码链接
Device 镜像嵌入 Host 目标文件
Host 链接

其中最重要的几种工作模式是：

同时编译 Host + Device
只编译 Device
只编译 Host
分离式编译
多架构 fatbinary 编译

3.3 支持的输入文件后缀

输入后缀	说明
`.mu`	MUSA 源文件，包含 Host 与 Device 代码
`.c`	C 源文件
`.cc`, `.cxx`, `.cpp`	C++ 源文件
`.ll`	LLVM IR 文件
`.o`, `.obj`	目标文件
`.a`	静态库
`.so`	动态库

3.4 支持的阶段控制

mcc 当前常用的 phase control 选项如下：

选项	说明
`-c`	只运行编译/汇编，不执行最终链接
`-emit-llvm`	输出 LLVM 表示
`--musa-compile-host-device`	同时编译 Host 与 Device，默认行为
`--musa-device-only`	只编译 Device 部分
`--musa-host-only`	只编译 Host 部分
`-fgpu-rdc`	打开可重定位设备代码，供分离式编译/设备链接使用

说明：

提示

--musa-device-only / --musa-host-only 用于显式选择 Device 或 Host 编译路径
分离式编译时，通常与 -fgpu-rdc 一起使用

4. 编译轨迹

mcc 的总体编译轨迹如下：

编译轨迹

编译流程说明：

输入源文件（.mu / .cu / .cpp）→ 预处理
预处理后 → 拆分为 Host 和 Device 两条编译路径
Device 路径：
- Device 前端编译 → LLVM IR
- MTGPU 代码生成 → Device 目标文件
- （可选）Device 链接 → GPU 二进制 / Fatbinary
- 将 GPU 二进制嵌入 Host 目标文件
Host 路径：
- Host 前端编译 → LLVM IR → Host 目标文件
Host 目标文件 + 嵌入的 GPU 二进制 → Host 链接
最终输出：可执行文件或共享库

4.1 单文件默认模式

默认情况下，mcc 会：

从一个源文件中拆出 Host 路径和 Device 路径
先生成 Device 目标代码
把 Device 镜像嵌入到 Host 目标文件
再完成最终链接

4.2 多架构编译

可以通过重复传递 --offload-arch 生成多架构 fatbinary，例如：

mcc axpy.mu \
  --offload-arch=<arch1> \
  --offload-arch=<arch2> \
  -lmusart -L/usr/local/musa/lib -o axpy

这类命令会把多个设备架构版本一并打入最终产物。

5. MCC 命令选项

本章按通用异构编译驱动的方式整理 mcc 常用选项。未列出的通用编译选项，通常可参考 clang/LLVM 行为。

5.1 命令选项类型与记号说明

mcc 的命令行选项风格整体继承 clang/LLVM，可以分成以下几类：

文件与路径类选项
编译阶段控制类选项
编译器/链接器行为控制类选项
Host/Linker/Backend 透传类选项
MUSA/MTGPU 专用扩展类选项

常见记号约定：

file：文件名
dir：目录名
path：路径
arch：MTGPU 目标架构，例如 mp_21
N：数值型参数

5.2 命令选项说明

5.2.1 文件和路径选项

选项	说明
`-o <file>`	指定输出文件名
`-x <language>`	显式指定输入语言
`-D<macro>=<value>`	定义预处理宏
`-U<macro>`	取消预定义宏
`-include <file>`	在主文件前强制包含头文件
`--musa-path=<path>`	指定 MUSA Toolkit 根目录
`--musa-path-ignore-env`	忽略环境变量中的 Toolkit 路径信息
`--musa-inc-path=<path>`	指定 MUSA 头文件目录
`--musa-libdevice-path=<path>`	指定 libdevice 路径
`--musa-rt-lib-path=<path>`	指定 MUSA runtime 库路径
`-I <dir>`	增加头文件搜索路径
`-isystem <dir>`	增加 system 头文件搜索路径
`-L <dir>`	增加库搜索路径
`-MF <file>`	指定依赖文件输出名
`-MP`	为依赖中的目标生成 phony target

5.2.2 阶段控制选项

选项	说明
`-c`	编译生成目标文件，不做最终链接
`-emit-llvm`	输出 LLVM IR/bitcode 形式
`-E`	仅预处理
`-M`	生成 make 依赖
`-MM`	生成不含 system headers 的依赖
`-MD`	编译同时生成依赖
`-MMD`	编译同时生成不含 system headers 的依赖
`--musa-compile-host-device`	同时编译 Host 和 Device
`--musa-device-only`	只编译 Device
`--musa-host-only`	只编译 Host
`-fgpu-rdc`	生成可重定位设备代码

说明：

提示

mcc 支持 -c、-E、-M、-MM、-MD、-MMD 这组 clang 风格阶段控制选项
-fatbin / -mubin 类中间产物在实际使用中可通过 --save-temps 获取

5.2.3 编译器和链接器行为选项

选项	说明
`--help`, `-help`	显示帮助信息
`--help-hidden`	显示隐藏选项
`-v`	打印子命令和详细执行过程
`-###`	只打印将执行的子命令，不真正执行
`--version`	打印版本信息
`--save-temps`	保留中间文件
`-l<name>`	链接指定库
`-nogpuinc`	不自动添加 GPU 相关默认 include 路径
`-nogpulib`	不自动链接设备端默认库

5.2.4 MUSA/MTGPU 目标与 Offload 控制选项

选项	说明
`-mtgpu`	将源码按 MTGPU/MUSA 路径处理；对 `.cu` 文件尤为重要
`--offload-arch=<arch>`	指定设备架构，可重复指定实现多架构编译
`-cuda_wrapper`	启用兼容封装模式，兼容既有 API 调用方式

5.2.5 优化选项

选项	说明
`-O1` / `-O2` / `-O3`	优化等级
`-O`	等价于 `-O2`

注意：

mcc 的默认优化等级是 -O2

5.2.6 前端选项

选项	说明
`-Wno-shared-var-init`	忽略 shared memory 初始化引发的报错，需用户自行确保安全性

5.2.7 后端选项

Backend 选项统一通过 -mllvm 透传给 LLVM/MTGPU 后端。

选项	说明
`-mllvm -mtgpu-maxregcnt=<N>`	限制 MTGPU TEMP 寄存器数
`-mllvm -mtgpu-if-convert`	小块 `if` 转 predicate 执行
`-mllvm -mtgpu-internalize-symbols=<bool>`	内部化符号
`-mllvm -mtgpu-enable-const-calc`	打开常量计算程序生成优化
`-mllvm -mtgpu-load-store-opt`	打开 shared memory load/store 合并优化
`-mllvm -enable-ldma-index`	允许 ldma 使用非零 imm-offset index
`-mllvm -mtgpu-prera-internal-opt`	RA 前 internal register 优化
`-mllvm -mtgpu-postra-internal-opt`	RA 后 internal register 优化
`-mllvm -internal-opt-avoid-use-nearest`	调整 internal register 优化区间
`-mllvm -mtgpu-load-cluster-mutation`	RA 后 load cluster mutation
`-mllvm -mtgpu-memory-sched-mutation`	RA 后 memory-specific schedule mutation
`-mllvm -mtgpu-tiny-offset-hint`	假定不会使用超大地址偏移，争取更多优化机会

说明：

提示

-mtgpu-prera-internal-opt 和 -mtgpu-postra-internal-opt 不能同时启用

5.2.8 特殊控制选项

选项	说明
`-mathx-disable-fp64`	在不支持硬件 FP64 的目标上禁用 FP64，减小代码体积、加快编译

5.2.9 透传选项

选项	说明
`-Xclang <arg>`	透传给 clang 前端；当前分离式编译 embed 设备镜像时会直接使用
`-Xlinker <arg>`	透传给 host linker

说明：

提示

其他编译驱动中的 host/linker/backend 透传选项，在 mcc 下并非全部一一对应
当前最常见且已验证的用法是 -Xclang -fcuda-include-gpubinary -Xclang all.d.o

5.2.10 示例

编译 `.mu` 文件

mcc axpy.mu -lmusart -L/usr/local/musa/lib -O2 -o axpy

编译 `.cu` 文件

mcc axpy.cu -mtgpu -lmusart -L/usr/local/musa/lib -O2 -o axpy

显式指定语言

mcc -x musa axpy.cu -lmusart -L/usr/local/musa/lib -O2 -o axpy

兼容封装模式

mcc axpy.cu -mtgpu -cuda_wrapper -lcuda2musa -L/usr/local/musa/lib -O2 -o axpy

编译提示：

mcc 编译器工具链默认启用了 --as-needed 链接器选项。这意味着链接器只会链接实际被引用的库，而不会链接命令行中指定的全部库。

运行可执行文件时，如果系统默认库搜索路径中没有 /usr/local/musa/lib，需要设置：

export LD_LIBRARY_PATH=/usr/local/musa/lib:$LD_LIBRARY_PATH

如果遇到链接错误，特别是“未定义引用”错误，建议首先检查 -l 参数的链接顺序是否正确。建议采用以下顺序：

源代码或目标文件放在所有 -l 链接参数之前。
越基础的库，通过 -l 指定链接时越靠后放置。

在某些复杂情况下，也可以显式使用 -Wl,--no-as-needed 覆盖默认行为。

6. 扩展属性

6.1 `mtgpu_num_usreg`

__attribute__((mtgpu_num_usreg(256))) void KernelName(...) { ... }

作用：

为 MTGPU 后端提供 TEMP/ATTR 使用上限提示
属于优化 hint，不保证严格按给定值分配

6.2 `mtgpu_unroll_threshold`

__attribute__((mtgpu_unroll_threshold(400))) void KernelName(...) { ... }

作用：

为循环展开阈值提供 hint

6.3 `mtgpu_tiny_offset`

__global__ __attribute__((mtgpu_tiny_offset)) void kernel_name(...) { ... }

作用：

告知编译器该 kernel 的访存 offset 不会超过 INT_MAX
帮助后端选择更激进的优化

7. 分离式编译和多架构编译

7.1 概述

分离式编译用于以下场景：

多个源文件分别生成设备目标文件
单独执行设备链接
将链接后的设备镜像嵌入 Host 目标文件
最后统一进行 Host 链接

典型流程如下：

7.2 关键命令

设备端编译

mcc -c -fgpu-rdc -mtgpu --musa-device-only \
  --offload-arch=<arch> file.cu -o file.d.o

设备端链接

lld -flavor gnu --no-undefined -shared file1.d.o file2.d.o -o all.d.o

准备空源文件并嵌入设备镜像

mcc empty.cu -mtgpu -c -fgpu-rdc --musa-host-only \
  -Xclang -fcuda-include-gpubinary \
  -Xclang all.d.o \
  -o empty.embed.o

说明：

提示

empty.cu 是额外准备的空文件
这一步的目标是把 all.d.o 中的设备镜像嵌入到 Host 目标文件中

Host 端编译

mcc -c -fgpu-rdc -mtgpu --musa-host-only file.cu -o file.h.o

最终 Host 链接

最终把以下文件一起链接：

file1.h.o
file2.h.o
empty.embed.o
Host 侧依赖库
MUSA runtime 库

例如：

mcc file1.h.o file2.h.o empty.embed.o -lmusart -L/usr/local/musa/lib -o app

7.3 示例

以下示例保持分离式编译设计中的基本结构：

第 1 步：设备端编译

mcc -c -fgpu-rdc -mtgpu --musa-device-only \
  --offload-arch=<arch> a.cu -o a.d.o
mcc -c -fgpu-rdc -mtgpu --musa-device-only \
  --offload-arch=<arch> b.cu -o b.d.o

第 2 步：设备端链接

lld -flavor gnu --no-undefined -shared a.d.o b.d.o -o all.d.o

第 3 步：嵌入设备镜像

mcc empty.cu -mtgpu -c -fgpu-rdc --musa-host-only \
  -Xclang -fcuda-include-gpubinary \
  -Xclang all.d.o \
  -o empty.embed.o

第 4 步：Host 端编译

mcc -c -fgpu-rdc -mtgpu --musa-host-only a.cu -o a.h.o
mcc -c -fgpu-rdc -mtgpu --musa-host-only b.cu -o b.h.o

第 5 步：最终链接

mcc a.h.o b.h.o empty.embed.o -lmusart -L/usr/local/musa/lib -o app

7.4 CMake 工程适配示例

下面给出一个最小 CMake 工程示例，演示如何把分离式编译流程接入 CMake。

示例目录结构：

example/
├── CMakeLists.txt
├── main.cu
├── test.cu
└── test.cuh

示例源码

test.cuh

#pragma once

class Test {
public:
    __host__ __device__ void Foo();
};

test.cu

#include "test.cuh"

__host__ __device__ void Test::Foo() {
}

main.cu

#include "test.cuh"
#include <musa_runtime.h>

__global__ void test() {
    Test test;
    test.Foo();
}

int main() {
    test <<< 1, 1 >>>();
    musaError_t err = musaGetLastError();
    if (err != musaSuccess) {
        return 1;
    }
    err = musaDeviceSynchronize();
    return err == musaSuccess ? 0 : 1;
}

CMakeLists.txt

这个示例的核心思路是：

同一组 .cu 源文件分别加入 test_device 和 test_host
test_device 用 --musa-device-only 和 -fgpu-rdc 编译成 device object
用 lld -shared 把所有 device object 链成一个 device.o
准备一个空源文件，把 device.o 嵌入成 fatbin.o
test_host 用 --musa-host-only 编译成 host object
最后把 host object 和 fatbin.o 一起链接成最终程序

示例配置如下：

cmake_minimum_required(VERSION 3.18)

project(Test LANGUAGES CXX)

set(CMAKE_CXX_COMPILER /usr/local/musa/bin/mcc)

set(TEMP_DIR "${CMAKE_BINARY_DIR}")
set(MUSA_ARCH mp_31 CACHE STRING "MUSA offload architecture")

# 源码文件列表，修改这里以适配工程
# 这些源文件会用到两次，一次编译device代码，一次编译host代码
set(SOURCES
    main.cu
    test.cu
)

# device侧代码的编译
# mcc xxx.mu -x musa -c -fgpu-rdc --musa-device-only --offload-arch=<arch>
add_library(test_device OBJECT)
target_compile_options(test_device PRIVATE
    -x musa -c -fgpu-rdc --musa-device-only --offload-arch=${MUSA_ARCH})

# host侧代码的编译
# mcc xxx.mu -mtgpu -c -fgpu-rdc --musa-host-only
add_library(test_host OBJECT)
target_compile_options(test_host PRIVATE
    -mtgpu -c -fgpu-rdc --musa-host-only)

foreach(file ${SOURCES})
    set_source_files_properties(${file} PROPERTIES LANGUAGE CXX)
    target_sources(test_device PRIVATE ${file})
    target_sources(test_host PRIVATE ${file})
endforeach()

# device侧代码链接为一个device.o
# lld -flavor gnu --no-undefined -shared a.o b.o c.o -o device.o
add_custom_target(all_device
    COMMAND_EXPAND_LISTS
    COMMAND lld -flavor gnu --no-undefined -shared $<TARGET_OBJECTS:test_device> -o ${TEMP_DIR}/device.o
    DEPENDS test_device
    COMMENT "Linking device objects")

# 创建一个空源码，将device.o打包为fatbin
# mcc empty.mu -o fatbin.o -x musa -c -fgpu-rdc --musa-host-only -Xclang -fcuda-include-gpubinary -Xclang device.o
add_custom_target(empty
    COMMAND ${CMAKE_COMMAND} -E touch ${TEMP_DIR}/empty.mu
)
add_custom_target(fatbin
    COMMAND mcc ${TEMP_DIR}/empty.mu -o ${TEMP_DIR}/fatbin.o
            -x musa -c -fgpu-rdc --musa-host-only
            -Xclang -fcuda-include-gpubinary
            -Xclang ${TEMP_DIR}/device.o
)
add_dependencies(fatbin all_device empty)

# host及device fatbin的链接，创建最终可执行程序
# mcc -o a.out a.o b.o c.o fatbin.o -lmusart -L/usr/local/musa/lib
add_custom_target(test ALL
    DEPENDS test_host fatbin
    COMMAND_EXPAND_LISTS
    COMMAND mcc -o test $<TARGET_OBJECTS:test_host> ${TEMP_DIR}/fatbin.o -lmusart -L/usr/local/musa/lib
)

构建方式

mkdir -p build
cd build
cmake .. -DMUSA_ARCH=<arch>
cmake --build .

适配要点

SOURCES 里的每个 .cu 文件会被编译两次：
- 一次进入 test_device
- 一次进入 test_host
test_device 路径需要：
- -x musa
- -c
- -fgpu-rdc
- --musa-device-only
test_host 路径需要：
- -mtgpu
- -c
- -fgpu-rdc
- --musa-host-only
设备链接不通过 mcc 直接完成，而是显式调用：
- lld -flavor gnu --no-undefined -shared
embed 阶段需要一个额外空文件，例如 empty.mu 或 empty.cu
embed 的关键参数是：
- -Xclang -fcuda-include-gpubinary
- -Xclang ${TEMP_DIR}/device.o

可按工程实际修改的部分

SOURCES
目标架构参数，例如：
- -DMUSA_ARCH=mp_31
最终链接库：
- -lmusart
- -L/usr/local/musa/lib
输出文件名：
- device.o
- fatbin.o
- test

1. 概述​

1.1 前提条件​

1.2 MCC 的作用​

2. Host 编译器​

3. 编译阶段​

3.1 MCC 识别宏​

3.2 MCC 编译阶段​

3.3 支持的输入文件后缀​

3.4 支持的阶段控制​

4. 编译轨迹​

4.1 单文件默认模式​

4.2 多架构编译​

5. MCC 命令选项​

5.1 命令选项类型与记号说明​

5.2 命令选项说明​

5.2.1 文件和路径选项​

5.2.2 阶段控制选项​

5.2.3 编译器和链接器行为选项​

5.2.4 MUSA/MTGPU 目标与 Offload 控制选项​

5.2.5 优化选项​

5.2.6 前端选项​

5.2.7 后端选项​

5.2.8 特殊控制选项​

5.2.9 透传选项​

5.2.10 示例​

编译 .mu 文件​

编译 .cu 文件​

显式指定语言​

兼容封装模式​

6. 扩展属性​

6.1 mtgpu_num_usreg​

6.2 mtgpu_unroll_threshold​

6.3 mtgpu_tiny_offset​

7. 分离式编译和多架构编译​

7.1 概述​

7.2 关键命令​

设备端编译​

设备端链接​

准备空源文件并嵌入设备镜像​

Host 端编译​

最终 Host 链接​

7.3 示例​

第 1 步：设备端编译​

第 2 步：设备端链接​

第 3 步：嵌入设备镜像​

第 4 步：Host 端编译​

第 5 步：最终链接​

7.4 CMake 工程适配示例​

示例源码​

CMakeLists.txt​

构建方式​

适配要点​

可按工程实际修改的部分​

相关文档​

这篇文章对您有帮助吗？