Skip to main content

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 支持通过字符串提供额外头文件不依赖磁盘文件,便于封装
底层符号名查询通过 mtrtcAddNameExpressionmtrtcGetLoweredName 获取底层符号名(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 个阶段:

  1. 创建 program:用源码字符串和可选额外头文件创建 mtrtcProgram
  2. 注册名称表达式:如需查询模板内核或全局变量符号,请在编译前调用 mtrtcAddNameExpression
  3. 执行编译:通过 mtrtcCompileProgram 传入支持的选项子集,生成二进制产物。
  4. 获取结果:读取构建日志、二进制产物和 lowered name。
  5. 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 SDK5.2.0
头文件mtrtc.hmusa.h
运行库libmtrtc.solibmusa.so

2.2 安装

MTRTC 随 MUSA SDK 提供,无需单独安装。开始前,请确认编译环境可以找到 mtrtc.hmusa.h,运行环境可以加载 libmtrtc.solibmusa.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.024.036.048.0。如果您只是验证编译、加载和启动链路,这个示例已经足够。

您可以通过以下现象判断示例是否运行成功:

  • mtrtcCompileProgram 返回 MTRTC_SUCCESS
  • muModuleLoadData 能成功加载编译产物
  • 内核能够正常启动并完成执行
  • 输出结果为 12.024.036.048.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
通用信息查询mtrtcVersionmtrtcGetNumSupportedArchsmtrtcGetSupportedArchs
编译mtrtcCreateProgrammtrtcDestroyProgrammtrtcCompileProgrammtrtcAddNameExpressionmtrtcGetLoweredNamemtrtcGetFatBinSizemtrtcGetFatBinmtrtcGetProgramLogSizemtrtcGetProgramLog

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_FAILUREprogram 创建失败
MTRTC_ERROR_INVALID_INPUT输入参数非法
MTRTC_ERROR_INVALID_PROGRAMprogram 句柄非法
MTRTC_ERROR_INVALID_OPTION编译选项非法或不被识别
MTRTC_ERROR_COMPILATION编译失败
MTRTC_ERROR_BUILTIN_OPERATION_FAILUREbuiltin 或 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 通用信息查询

本节介绍版本号和支持架构查询接口。

函数原型作用成功返回失败返回
mtrtcVersionmtrtcResult mtrtcVersion(int *major, int *minor);查询 MTRTC 主版本号和次版本号MTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUT
mtrtcGetNumSupportedArchsmtrtcResult mtrtcGetNumSupportedArchs(int *numArchs);查询当前支持的架构数量MTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUT
mtrtcGetSupportedArchsmtrtcResult mtrtcGetSupportedArchs(int *supportedArchs);按升序写出当前支持的架构列表MTRTC_SUCCESSMTRTC_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 创建、销毁与编译

函数原型作用主要返回值
mtrtcCreateProgrammtrtcResult mtrtcCreateProgram(mtrtcProgram *prog, const char *src, const char *name, int numHeaders, const char *const *headers, const char *const *includeNames);用源码字符串和可选额外头文件创建 programMTRTC_SUCCESSMTRTC_ERROR_OUT_OF_MEMORYMTRTC_ERROR_PROGRAM_CREATION_FAILUREMTRTC_ERROR_INVALID_INPUTMTRTC_ERROR_INVALID_PROGRAM
mtrtcDestroyProgrammtrtcResult mtrtcDestroyProgram(mtrtcProgram *prog);销毁 programMTRTC_SUCCESSMTRTC_ERROR_INVALID_PROGRAM
mtrtcCompileProgrammtrtcResult mtrtcCompileProgram(mtrtcProgram prog, int numOptions, const char *const *options);编译 program,并把生成的二进制保存在 program 内部MTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUTMTRTC_ERROR_INVALID_PROGRAMMTRTC_ERROR_INVALID_OPTIONMTRTC_ERROR_COMPILATION

mtrtcCreateProgram 的关键约束如下:

  • prog 不能为空;当 progNULL 时返回 MTRTC_ERROR_INVALID_PROGRAM
  • name 可以为 NULL 或空字符串;这两种情况下会使用 default_program
  • numHeaders 必须大于等于 0
  • numHeaders0 时,headersincludeNames 可以为 NULL
  • numHeaders 大于 0 时,headersincludeNames 都必须非空。
  • 每个 header 源码和 include name 都必须非空。
  • include name 在同一个 program 内必须唯一。
  • 这些额外头文件在后续 mtrtcCompileProgram 中可按对应的 include name 被源码引用。

mtrtcCompileProgram 的关键约束如下:

  • options 可以在 numOptions0 时为 NULL
  • 如果没有显式传入架构选项,当前实现默认使用 mp_31
  • 具体可用选项以当前版本支持情况为准。

3.3.3 名称表达式与底层符号名

函数原型作用主要返回值
mtrtcAddNameExpressionmtrtcResult mtrtcAddNameExpression(mtrtcProgram prog, const char *const name_expression);注册源码级表达式,供编译后查询 lowered nameMTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUTMTRTC_ERROR_INVALID_PROGRAMMTRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION
mtrtcGetLoweredNamemtrtcResult mtrtcGetLoweredName(mtrtcProgram prog, const char *const name_expression, const char **lowered_name);查询已注册表达式对应的 lowered nameMTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUTMTRTC_ERROR_INVALID_PROGRAMMTRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATIONMTRTC_ERROR_NAME_EXPRESSION_NOT_VALID

使用这两个接口时请注意:

  • mtrtcAddNameExpression 必须在 mtrtcCompileProgram 之前调用。
  • 重复注册同一个表达式是允许的。
  • mtrtcGetLoweredName 必须使用与注册时完全一致的原始字符串。
  • mtrtcGetLoweredName 返回的字符串由 program 持有。
  • 返回的 lowered name 在下一次 mtrtcCompileProgram 之后,或 mtrtcDestroyProgram 之后失效。
  • 如果已注册的表达式本身无效,mtrtcCompileProgram 可能直接因这些表达式而返回 MTRTC_ERROR_COMPILATION

3.3.4 编译产物与构建日志

函数原型作用主要返回值
mtrtcGetFatBinSizemtrtcResult mtrtcGetFatBinSize(mtrtcProgram prog, size_t *fatbinSizeRet);查询当前二进制产物大小MTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUTMTRTC_ERROR_INVALID_PROGRAM
mtrtcGetFatBinmtrtcResult mtrtcGetFatBin(mtrtcProgram prog, char *fatbin);读取当前二进制产物MTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUTMTRTC_ERROR_INVALID_PROGRAM
mtrtcGetProgramLogSizemtrtcResult mtrtcGetProgramLogSize(mtrtcProgram prog, size_t *logSizeRet);查询构建日志大小MTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUTMTRTC_ERROR_INVALID_PROGRAM
mtrtcGetProgramLogmtrtcResult mtrtcGetProgramLog(mtrtcProgram prog, char *log);读取构建日志MTRTC_SUCCESSMTRTC_ERROR_INVALID_INPUTMTRTC_ERROR_INVALID_PROGRAM

使用这些接口时请注意:

  • mtrtcGetFatBin 的目标缓冲区长度至少要达到 mtrtcGetFatBinSize 的返回值。
  • mtrtcGetFatBin 返回的数据按原始字节处理,不保证以 NUL 结尾。
  • mtrtcGetProgramLogSize 返回值包含末尾 NUL。
  • mtrtcGetProgramLog 的目标缓冲区长度至少要达到 mtrtcGetProgramLogSize 的返回值。
  • 当前实现返回的构建日志字符串以 NUL 结尾。

4. 附录

4.1 NVRTC 接口对应关系

#NVRTC API / 能力MTRTC API / 能力状态
1nvrtcVersionmtrtcVersion支持
2nvrtcCreateProgrammtrtcCreateProgram支持
3nvrtcCompileProgrammtrtcCompileProgram支持,选项为子集
4nvrtcGetCUBINSize / nvrtcGetCUBINmtrtcGetFatBinSize / mtrtcGetFatBin接口对应,用于获取可加载二进制产物
5nvrtcAddNameExpression / nvrtcGetLoweredNamemtrtcAddNameExpression / mtrtcGetLoweredName支持,static __global__ 除外
6--device-c + jitLink不支持

4.2 相关文档

4.3 名词解释

术语全称说明
MTRTCMUSA Runtime CompilationMUSA 的运行时编译库
programmtrtcProgram持有源码、编译产物、构建日志和 lowered name 结果的句柄
lowered name-编译后解析得到的底层符号名,用于查找 C++ 内核或全局变量
name expression-编译前注册的源码层表达式,例如 fill_const<float>
Driver APIMUSA Driver API负责加载模块、查找符号和启动内核的接口
MPMUSA Processormp_NN 架构命名中的前缀,用于 -arch / --gpu-architecture

4.4 已知约束

MTRTC 5.2.0 当前有以下约束:

  • 异常路径行为不保证与 NVRTC 完全兼容。
  • 参数校验和部分错误码行为与 NVRTC 不完全一致。
  • 当前不支持 jitLink,因此不支持 --device-c
  • static __global__ 不支持 name expression 注册。
  • __MUSACC_RTC__ 当前未预定义。