MTRTC 运行时编译用户指南
1. 概述
1.1 什么是 MTRTC?
MTRTC(MUSA Runtime Compilation) 是 MUSA SDK 提供的运行时编译库,库文件为 libmtrtc.so。MTRTC 接收字符串形式的 MUSA 设备端源码,并在运行时将其编译为可加载的二进制产物,再交给 MUSA Driver API 完成模块加载、符号查找和内核启动。
如果您已经熟悉离线 mcc 编译流程,可以将 MTRTC 理解为运行时场景下的设备端编译入口。MTRTC 的输出形式、支持选项和异常路径,以当前实现为准。
1.2 为什么使用 MTRTC?
MTRTC 主要面向“源码只能在运行时确定”的场景。在这些场景中,纯离线编译往往不够用,应用通常需要在运行时动态生成源码、拼接宏和头文件,或者按模板实例化结果解析 C++ 符号。
如果没有运行时编译库,应用通常只能预编译大量变体,或者在运行时调用外部编译工具处理源码。这会增加编译链路复杂度,也会让部署和调试变得更重。MTRTC 把这条流程收敛为一组库接口,直接提供 program 生命周期管理、构建日志获取、二进制产物获取,以及底层符号名(lowered name)查询。
因此,MTRTC 特别适合运行时生成源码、内存内依赖管理、C++ 符号查找,以及需要把编译产物直接接入 MUSA Driver API 的场景。
1.3 关键特性
| 特性 | 说明 | 好处 |
|---|---|---|
| 运行时编译 | 直接编译字符串形式的设备端源码 | 适合动态生成内核或按需编译 |
| 二进制编译产物 | 编译结果通过 mtrtcGetFatBin(Size) 获取 | 可直接交给 MUSA Driver API 加载 |
| 额外头文件输入 | mtrtcCreateProgram 支持通过字符串提供额外头文件 | 不依赖磁盘文件,便于封装 |
| 底层符号名查询 | 通过 mtrtcAddNameExpression 和 mtrtcGetLoweredName 获取底层符号名(lowered name) | 便于处理模板内核、命名空间内核和全局变量 |
| 可重复编译 | 同一 program 支持再次调用 mtrtcCompileProgram | 适合动态调参或重复构建 |
| 选项支持范围明确 | mtrtcCompileProgram 支持的编译选项以当前版本为准 | 便于建立稳定预期 |
详细 API、参数、返回值和条件可用能力,参见第 3 节。与 NVRTC 的接口对应关系见 4.1 NVRTC 接口对应关系。
1.4 MTRTC 如何工作?
1.4.1 典型使用路径
MTRTC 的工作流程可以概括为:在运行时接收源码,完成编译,再把编译产物直接交给 MUSA Driver API 加载执行。主路径如下:源码字符串 / 额外头文件 -> mtrtcCreateProgram -> mtrtcCompileProgram -> 构建日志 / 二进制产物 / lowered name -> MUSA Driver API 加载执行。
按调用顺序,MTRTC 的典型使用路径可以概括为 5 个阶段:
- 创建 program:用源码字符串和可选额外头文件创建
mtrtcProgram。 - 注册名称表达式:如需查询模板内核或全局变量符号,请在编译前调用
mtrtcAddNameExpression。 - 执行编译:通过
mtrtcCompileProgram传入支持的选项子集,生成二进制产物。 - 获取结果:读取构建日志、二进制产物和 lowered name。
- MUSA Driver API 加载执行:用
muModuleLoadData加载编译产物,再查找符号并启动内核。
1.4.2 系统架构
1.4.1 典型使用路径 说明调用顺序,本节说明组件分工。MTRTC 在系统中的位置可以概括为:Host 程序提供源码输入,libmtrtc.so 负责运行时编译,libmusa.so 负责加载并在 GPU 上执行。
┌──────────────────────────────┐
│ Host 侧程序(C/C++) │
└──────────────┬───────────────┘
│
│ 源码字符串 / 额外头文件
v
┌──────────────────────────────┐
│ libmtrtc.so │
│ 运行时编译 / 生成二进制产物 │
└──────────────┬───────────────┘
│
│ 编译产物
v
┌──────────────────────────────┐
│ libmusa.so │
│ 加载 / 查找符号 / 启动 │
└──────────────┬───────────────┘
│
v
┌──────────────────────────────┐
│ GPU 执行 │
└──────────────────────────────┘
各组件的职责如下:
| 组件 | 主要职责 |
|---|---|
| Host 侧程序 | 准备源码字符串、额外头文件内容和编译选项,并串联整个调用流程。 |
libmtrtc.so | 接收输入、完成运行时编译,并产出构建日志、二进制产物和 lowered name。 |
libmusa.so | 加载编译产物、查找内核或全局变量符号,并发起执行。 |
| GPU | 执行加载后的设备端代码。 |
1.4.3 核心机制
额外头文件输入
MTRTC 支持把源码和额外头文件内容都直接作为字符串传入。头文件内容放在 headers[],头文件名放在 includeNames[],源码再通过 #include "..." 引用对应头文件。
src 字符串 + headers[] + includeNames[]
|
v
mtrtcCreateProgram
|
v
#include "extra_header.h"
在内存中完成解析
这样,运行时编译所需输入可以全部保存在内存中,而不依赖磁盘上的真实头文件。
名称表达式与底层符号名(lowered name)
对于模板内核、命名空间内核、C++ 修饰后的符号,以及 __device__ / __constant__ 变量,您通常需要先获取 lowered name,MUSA Driver API 才能正确查找对应符号。
"fill_const<float>"
|
v
mtrtcAddNameExpression
|
v
mtrtcCompileProgram
|
v
mtrtcGetLoweredName
|
v
muModuleGetFunction / muModuleGetGlobal
使用时请注意:
- 必须在
mtrtcCompileProgram之前调用mtrtcAddNameExpression。 - 只能在编译成功后调用
mtrtcGetLoweredName。 mtrtcGetLoweredName的表达式字符串必须与mtrtcAddNameExpression完全一致。- 返回的 lowered name 指针由
program持有;再次编译或销毁program后,该指针失效。 static __global__当前不支持 name expression 注册。
同一 program 支持重复编译。再次调用 mtrtcCompileProgram 后,旧的二进制产物、构建日志和 lowered name 结果都会失效;已注册的 name expression 会保留,并在下一轮编译中重新解析。
2. 快速开始
2.1 准备工作
硬件要求
- 可用的 MUSA 设备
- 目标设备架构与
-arch/--gpu-architecture选项一致
软件要求
| 要求 | 说明 |
|---|---|
| MUSA SDK | 5.2.0 |
| 头文件 | mtrtc.h、musa.h |
| 运行库 | libmtrtc.so、libmusa.so |
2.2 安装
MTRTC 随 MUSA SDK 提供,无需单独安装。开始前,请确认编译环境可以找到 mtrtc.h 和 musa.h,运行环境可以加载 libmtrtc.so 和 libmusa.so。
2.3 使用
2.3.1 基本用法
推荐按以下顺序调用:
mtrtcCreateProgram
-> mtrtcCompileProgram
-> mtrtcGetProgramLog
-> mtrtcGetFatBin
-> muModuleLoadData
-> muModuleGetFunction / muModuleGetGlobal
-> mtrtcDestroyProgram
使用时请注意:
- 即使编译成功,构建日志中也可能包含警告信息。
mtrtcGetProgramLogSize返回的大小包含末尾 NUL。mtrtcGetFatBinSize返回的是二进制字节数,返回结果不是字符串。- 如果您 需要查找模板内核、命名空间内核或全局变量,请在编译前补充
mtrtcAddNameExpression,并在编译后调用mtrtcGetLoweredName。
2.3.2 SAXPY 最小示例
如果您只想先跑通一个最小内核,建议先从 SAXPY 示例开始。这个示例只覆盖主路径:源码字符串 -> 编译 -> 获取二进制产物 -> MUSA Driver API 加载 -> 启动内核,不包含额外头文件或 lowered name。
#include <mtrtc.h>
#include <musa.h>
#include <cstdio>
#include <vector>
int main() {
const char* src = R"musa(
extern "C" __global__ void saxpy(float a, const float* x, float* y, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = a * x[i] + y[i];
}
)musa";
mtrtcProgram prog = nullptr;
mtrtcCreateProgram(&prog, src, "saxpy.mu", 0, nullptr, nullptr);
const char* options[] = {
"--gpu-architecture=mp_22",
"--std=c++17"
};
mtrtcResult r = mtrtcCompileProgram(prog, 2, options);
size_t logSize = 0;
mtrtcGetProgramLogSize(prog, &logSize);
std::vector<char> log(logSize);
mtrtcGetProgramLog(prog, log.data());
if (r != MTRTC_SUCCESS) {
std::fprintf(stderr, "compile failed: %s\n", mtrtcGetErrorString(r));
std::fprintf(stderr, "log:\n%s\n", log.data());
mtrtcDestroyProgram(&prog);
return 1;
}
size_t fatbinSize = 0;
mtrtcGetFatBinSize(prog, &fatbinSize);
std::vector<char> fatbin(fatbinSize);
mtrtcGetFatBin(prog, fatbin.data());
muInit(0);
MUdevice dev;
muDeviceGet(&dev, 0);
MUcontext ctx;
muCtxCreate(&ctx, 0, dev);
MUmodule mod;
muModuleLoadData(&mod, fatbin.data());
MUfunction saxpyFn;
muModuleGetFunction(&saxpyFn, mod, "saxpy");
std::vector<float> x = {1.0f, 2.0f, 3.0f, 4.0f};
std::vector<float> y = {10.0f, 20.0f, 30.0f, 40.0f};
float a = 2.0f;
int n = static_cast<int>(x.size());
MUdeviceptr dx;
MUdeviceptr dy;
muMemAlloc(&dx, n * sizeof(float));
muMemAlloc(&dy, n * sizeof(float));
muMemcpyHtoD(dx, x.data(), n * sizeof(float));
muMemcpyHtoD(dy, y.data(), n * sizeof(float));
void* args[] = {&a, &dx, &dy, &n};
int blockX = 256;
int gridX = (n + blockX - 1) / blockX;
muLaunchKernel(saxpyFn, gridX, 1, 1, blockX, 1, 1, 0, 0, args, nullptr);
muCtxSynchronize();
muMemcpyDtoH(y.data(), dy, n * sizeof(float));
for (int i = 0; i < n; ++i) {
std::printf("y[%d] = %.1f\n", i, y[i]);
}
muMemFree(dx);
muMemFree(dy);
muModuleUnload(mod);
muCtxDestroy(ctx);
mtrtcDestroyProgram(&prog);
return 0;
}
以上输入下,y 的预期结果为 12.0、24.0、36.0、48.0。如果您只是验证编译、加载和启动链路,这个示例已经足够。
您可以通过以下现象判断示例是否运行成功:
mtrtcCompileProgram返回MTRTC_SUCCESSmuModuleLoadData能成功加载编译产物- 内核能够正常启动并完成执行
- 输出结果为
12.0、24.0、36.0、48.0
2.3.3 进阶示例:额外头文件与 lowered name
如果您需要同时使用额外头文件和 lowered name,可参考以下示例。
#include <mtrtc.h>
#include <musa.h>
#include <cstdio>
#include <vector>
int main() {
const char* src = R"musa(
#include "axpy_constants.h"
extern "C" __global__ void axpy(const float* x, float* y, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = SCALE * x[i] + y[i];
}
template <typename T>
__global__ void fill_const(T* p, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) p[i] = static_cast<T>(SCALE);
}
template __global__ void fill_const<float>(float*, int);
)musa";
const char* header = "#define SCALE 2.5f\n";
const char* headers[] = {header};
const char* includeNames[] = {"axpy_constants.h"};
mtrtcProgram prog = nullptr;
mtrtcCreateProgram(&prog, src, "demo.mu", 1, headers, includeNames);
mtrtcAddNameExpression(prog, "fill_const<float>");
const char* options[] = {
"--gpu-architecture=mp_22",
"--std=c++17"
};
mtrtcResult r = mtrtcCompileProgram(prog, 2, options);
size_t logSize = 0;
mtrtcGetProgramLogSize(prog, &logSize);
std::vector<char> log(logSize);
mtrtcGetProgramLog(prog, log.data());
if (r != MTRTC_SUCCESS) {
std::fprintf(stderr, "compile failed: %s\n", mtrtcGetErrorString(r));
std::fprintf(stderr, "log:\n%s\n", log.data());
mtrtcDestroyProgram(&prog);
return 1;
}
size_t fatbinSize = 0;
mtrtcGetFatBinSize(prog, &fatbinSize);
std::vector<char> fatbin(fatbinSize);
mtrtcGetFatBin(prog, fatbin.data());
const char* lowered = nullptr;
mtrtcGetLoweredName(prog, "fill_const<float>", &lowered);
muInit(0);
MUdevice dev;
muDeviceGet(&dev, 0);
MUcontext ctx;
muCtxCreate(&ctx, 0, dev);
MUmodule mod;
muModuleLoadData(&mod, fatbin.data());
MUfunction fillFn;
muModuleGetFunction(&fillFn, mod, lowered);
muModuleUnload(mod);
muCtxDestroy(ctx);
mtrtcDestroyProgram(&prog);
return 0;
}
3. API
本章介绍 MTRTC 的 API。基本用法参见 2.3 使用。
当前支持的函数如下:
| 类别 | 函数 |
|---|---|
| 错误处理 | mtrtcGetErrorString |
| 通用信息查询 | mtrtcVersion、mtrtcGetNumSupportedArchs、mtrtcGetSupportedArchs |
| 编译 | mtrtcCreateProgram、mtrtcDestroyProgram、mtrtcCompileProgram、mtrtcAddNameExpression、mtrtcGetLoweredName、mtrtcGetFatBinSize、mtrtcGetFatBin、mtrtcGetProgramLogSize、mtrtcGetProgramLog |
3.1 错误处理
本节介绍结果码和错误字符串接口。
3.1.1 mtrtcResult
typedef enum {
MTRTC_SUCCESS = 0,
MTRTC_ERROR_OUT_OF_MEMORY = 1,
MTRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
MTRTC_ERROR_INVALID_INPUT = 3,
MTRTC_ERROR_INVALID_PROGRAM = 4,
MTRTC_ERROR_INVALID_OPTION = 5,
MTRTC_ERROR_COMPILATION = 6,
MTRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
MTRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
MTRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
MTRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
MTRTC_ERROR_INTERNAL_ERROR = 11,
MTRTC_ERROR_INPUT_FILE = 12,
MTRTC_ERROR_OUTPUT_FILE = 13,
MTRTC_ERROR_LINK = 14,
} mtrtcResult;
| 结果码 | 说明 |
|---|---|
MTRTC_SUCCESS | 调用成功 |
MTRTC_ERROR_OUT_OF_MEMORY | 内存分配失败 |
MTRTC_ERROR_PROGRAM_CREATION_FAILURE | program 创建失败 |
MTRTC_ERROR_INVALID_INPUT | 输入参数非法 |
MTRTC_ERROR_INVALID_PROGRAM | program 句柄非法 |
MTRTC_ERROR_INVALID_OPTION | 编译选项非法或不被识别 |
MTRTC_ERROR_COMPILATION | 编译失败 |
MTRTC_ERROR_BUILTIN_OPERATION_FAILURE | builtin 或 libdevice 相关操作失败 |
MTRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION | 编译后再注册 name expression |
MTRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION | 编译前查询 lowered name |
MTRTC_ERROR_NAME_EXPRESSION_NOT_VALID | 表达式未注册或解析失败 |
MTRTC_ERROR_INTERNAL_ERROR | 内部错误 |
MTRTC_ERROR_INPUT_FILE | 输入文件处理失败 |
MTRTC_ERROR_OUTPUT_FILE | 输出文件处理失败 |
MTRTC_ERROR_LINK | 链接或二进制产物生成失败 |
3.1.2 mtrtcGetErrorString
const char *mtrtcGetErrorString(mtrtcResult result);
mtrtcGetErrorString 用于把 mtrtcResult 转成可读字符串。对于未识别的枚举值,该函数返回 MTRTC_ERROR unknown。
3.2 通用信息查询
本节介绍版本号和支持架构查询接口。
| 函数 | 原型 | 作用 | 成功返回 | 失败返回 |
|---|---|---|---|---|
mtrtcVersion | mtrtcResult mtrtcVersion(int *major, int *minor); | 查询 MTRTC 主版本号和次版本号 | MTRTC_SUCCESS | MTRTC_ERROR_INVALID_INPUT |
mtrtcGetNumSupportedArchs | mtrtcResult mtrtcGetNumSupportedArchs(int *numArchs); | 查询当前支持的架构数量 | MTRTC_SUCCESS | MTRTC_ERROR_INVALID_INPUT |
mtrtcGetSupportedArchs | mtrtcResult mtrtcGetSupportedArchs(int *supportedArchs); | 按升序写出当前支持的架构列表 | MTRTC_SUCCESS | MTRTC_ERROR_INVALID_INPUT |
使用 mtrtcGetSupportedArchs 前,请先调用 mtrtcGetNumSupportedArchs 确定数组长度。
3.3 编译
本节介绍 mtrtcProgram 及其相关编译接口。
3.3.1 mtrtcProgram
typedef struct _mtrtcProgram *mtrtcProgram;
mtrtcProgram 是 MTRTC 的编译单元句柄。源码字符串需要先通过 mtrtcCreateProgram 创建为 program,再通过 mtrtcCompileProgram 编译。
3.3.2 program 创建、销毁与编译
| 函数 | 原型 | 作用 | 主要返回值 |
|---|---|---|---|
mtrtcCreateProgram | mtrtcResult mtrtcCreateProgram(mtrtcProgram *prog, const char *src, const char *name, int numHeaders, const char *const *headers, const char *const *includeNames); | 用源码字符串和可选额外头文件创建 program | MTRTC_SUCCESS、MTRTC_ERROR_OUT_OF_MEMORY、MTRTC_ERROR_PROGRAM_CREATION_FAILURE、MTRTC_ERROR_INVALID_INPUT、MTRTC_ERROR_INVALID_PROGRAM |
mtrtcDestroyProgram | mtrtcResult mtrtcDestroyProgram(mtrtcProgram *prog); | 销毁 program | MTRTC_SUCCESS、MTRTC_ERROR_INVALID_PROGRAM |
mtrtcCompileProgram | mtrtcResult mtrtcCompileProgram(mtrtcProgram prog, int numOptions, const char *const *options); | 编译 program,并把生成的二进制保存在 program 内部 | MTRTC_SUCCESS、MTRTC_ERROR_INVALID_INPUT、MTRTC_ERROR_INVALID_PROGRAM、MTRTC_ERROR_INVALID_OPTION、MTRTC_ERROR_COMPILATION |
mtrtcCreateProgram 的关键约束如下:
prog不能为空;当prog为NULL时返回MTRTC_ERROR_INVALID_PROGRAM。name可以为NULL或空字符串;这两种情况下会使用default_program。numHeaders必须大于等于0。- 当
numHeaders为0时,headers和includeNames可以为NULL。 - 当
numHeaders大于0时,headers和includeNames都必须非空。 - 每个 header 源码和 include name 都必须非空。
- include name 在同一个
program内必须唯一。 - 这些额外头文件在后续
mtrtcCompileProgram中可按对应的 include name 被源码引用。
mtrtcCompileProgram 的关键约束如下:
options可以在numOptions为0时为NULL。- 如果没有显式传入架构选项,当前实现默认使用
mp_31。 - 具体可用选项以当前版本支持情况为准。