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_SUCCESS
muModuleLoadData 能成功加载编译产物
内核能够正常启动并完成执行
输出结果为 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。
具体可用选项以当前版本支持情况为准。

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

函数	原型	作用	主要返回值
`mtrtcAddNameExpression`	`mtrtcResult mtrtcAddNameExpression(mtrtcProgram prog, const char *const name_expression);`	注册源码级表达式，供编译后查询 lowered name	`MTRTC_SUCCESS`、`MTRTC_ERROR_INVALID_INPUT`、`MTRTC_ERROR_INVALID_PROGRAM`、`MTRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION`
`mtrtcGetLoweredName`	`mtrtcResult mtrtcGetLoweredName(mtrtcProgram prog, const char const name_expression, const char *lowered_name);`	查询已注册表达式对应的 lowered name	`MTRTC_SUCCESS`、`MTRTC_ERROR_INVALID_INPUT`、`MTRTC_ERROR_INVALID_PROGRAM`、`MTRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION`、`MTRTC_ERROR_NAME_EXPRESSION_NOT_VALID`

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

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

3.3.4 编译产物与构建日志

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

使用这些接口时请注意：

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

4. 附录

4.1 NVRTC 接口对应关系

#	NVRTC API / 能力	MTRTC API / 能力	状态
1	`nvrtcVersion`	`mtrtcVersion`	支持
2	`nvrtcCreateProgram`	`mtrtcCreateProgram`	支持
3	`nvrtcCompileProgram`	`mtrtcCompileProgram`	支持，选项为子集
4	`nvrtcGetCUBINSize` / `nvrtcGetCUBIN`	`mtrtcGetFatBinSize` / `mtrtcGetFatBin`	接口对应，用于获取可加载二进制产物
5	`nvrtcAddNameExpression` / `nvrtcGetLoweredName`	`mtrtcAddNameExpression` / `mtrtcGetLoweredName`	支持，`static __global__` 除外
6	`--device-c` + `jitLink`	无	不支持

4.2 相关文档

4.3 名词解释

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

4.4 已知约束

MTRTC 5.2.0 当前有以下约束：

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

1. 概述​

1.1 什么是 MTRTC？​

1.2 为什么使用 MTRTC？​

1.3 关键特性​

1.4 MTRTC 如何工作？​

1.4.1 典型使用路径​

1.4.2 系统架构​

1.4.3 核心机制​

额外头文件输入​

名称表达式与底层符号名（lowered name）​

2. 快速开始​

2.1 准备工作​

硬件要求​

软件要求​

2.2 安装​

2.3 使用​

2.3.1 基本用法​

2.3.2 SAXPY 最小示例​

2.3.3 进阶示例：额外头文件与 lowered name​

3. API​

3.1 错误处理​

3.1.1 mtrtcResult​

3.1.2 mtrtcGetErrorString​

3.2 通用信息查询​

3.3 编译​

3.3.1 mtrtcProgram​

3.3.2 program 创建、销毁与编译​

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

3.3.4 编译产物与构建日志​

4. 附录​

4.1 NVRTC 接口对应关系​

4.2 相关文档​

4.3 名词解释​

4.4 已知约束​

这篇文章对您有帮助吗？