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 支持的阶段可以类比为：

1. 预处理
2. Host 侧前端编译
3. Device 侧前端编译
4. LLVM IR / 对象文件生成
5. Device 代码链接
6. Device 镜像嵌入 Host 目标文件
7. 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	打开可重定位设备代码，供分离式编译/设备链接使用

说明：

tip

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

---

4. 编译轨迹

mcc 的总体编译轨迹如下：

编译流程说明：

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

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	生成可重定位设备代码

说明：

tip

• 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	假定不会使用超大地址偏移，争取更多优化机会

说明：

tip

• -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

说明：

tip

• 其他编译驱动中的 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 参数的链接顺序是否正确。建议采用以下顺序：

1. 源代码或目标文件放在所有 -l 链接参数之前。
2. 越基础的库，通过 -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

说明：

tip

• 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

这个示例的核心思路是：

1. 同一组 .cu 源文件分别加入 test_device 和 test_host
2. test_device 用 --musa-device-only 和 -fgpu-rdc 编译成 device object
3. 用 lld -shared 把所有 device object 链成一个 device.o
4. 准备一个空源文件，把 device.o 嵌入成 fatbin.o
5. test_host 用 --musa-host-only 编译成 host object
6. 最后把 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

---

相关文档

• Driver API: Driver API Reference - MUSA Driver API 参考（底层硬件访问接口）
• Runtime API: Runtime API Reference - MUSA Runtime API 参考（高级编程接口）
• Math API: Math API Reference - MUSA Math API 参考）
• 编程指南: C++ 语言支持 | 第一个内核 | 编程模型
• 其他工具: MUSA Runtime - MUSA 运行时用户手册
• Musify：CUDA 到 MUSA 代码迁移工具