Runtime API 与 Driver API

大多数应用用 Runtime API 已能获得完整性能；Driver API 的价值是更显式地控制 context、module、JIT、VMM 和较新的高级特性。

Runtime 本身构建在 Driver 之上，两者可以在同一进程混用，但必须理解当前 context 和资源归属。

Driver 对象关系

cuInit
  └─ CUdevice：物理/逻辑设备描述
      └─ CUcontext：资源命名空间与当前执行环境
          ├─ CUmodule：已加载 GPU 代码容器
          │   ├─ CUfunction：kernel handle
          │   └─ global/constant symbols
          ├─ CUstream / CUevent / CUgraph
          └─ device allocations / VMM mappings

多数 Runtime 应用使用 device 的 primary context。若一个插件另建 context，框架创建的 pointer 与 stream 不能自然跨过去。库代码通常应保留调用方 current context，或明确 retain/release primary context。

最小 Driver 启动流程

#include <cuda.h>

CU_CHECK(cuInit(0));

CUdevice device{};
CU_CHECK(cuDeviceGet(&device, 0));

CUcontext context{};
CU_CHECK(cuDevicePrimaryCtxRetain(&context, device));
CU_CHECK(cuCtxSetCurrent(context));

CUmodule module{};
CU_CHECK(cuModuleLoad(&module, "kernels.fatbin"));

CUfunction kernel{};
CU_CHECK(cuModuleGetFunction(&kernel, module, "saxpy"));

void* args[] = {&a, &x, &y, &n};
CU_CHECK(cuLaunchKernel(kernel,
                       grid, 1, 1,
                       block, 1, 1,
                       0, stream, args, nullptr));

Driver launch 参数数组存放的是“每个参数值的地址”，而 device pointer 本身也是一个值。参数类型、大小、对齐和 kernel ABI 必须完全匹配；名字还要处理 C++ name mangling，常见做法是 kernel 入口使用 extern "C" 或通过 NVRTC lowered name 获取真实符号。

Context 栈与线程

Driver API 的 current context 与 host thread 有关。复杂插件应采用 RAII 恢复调用前状态：

class ScopedContext {
public:
    explicit ScopedContext(CUcontext next) {
        CU_CHECK(cuCtxGetCurrent(&previous_));
        CU_CHECK(cuCtxSetCurrent(next));
    }
    ~ScopedContext() { cuCtxSetCurrent(previous_); }
private:
    CUcontext previous_{};
};

真实析构函数不应抛异常；上例仅表达关系。还需规定多线程是否共享 context、谁负责最终 release，以及 shutdown 时还有没有异步工作。

Module 与代码格式

发行包常采用“已知架构 cubin + 一个合适 PTX fallback”。只放 PTX 会增加首次 JIT 延迟，并让最终代码质量依赖用户 Driver；只放 cubin 则无法覆盖未打包的新架构。

NVRTC 运行时编译流程

CUDA C++ source string + headers/options
              │ nvrtcCreateProgram
              ▼
        nvrtcCompileProgram
         ├─ log on failure
         ├─ PTX
         ├─ cubin（受目标/选项支持约束）
         ├─ LTO IR
         └─ Tile IR（Tile 路线）
              │
       Driver API / nvJitLink
              ▼
       CUmodule → CUfunction → launch

精简错误处理骨架：

nvrtcProgram prog{};
NVRTC_CHECK(nvrtcCreateProgram(&prog, source, "kernel.cu",
                               0, nullptr, nullptr));

const char* options[] = {"--std=c++17", "--gpu-architecture=compute_90"};
nvrtcResult rc = nvrtcCompileProgram(prog, 2, options);

size_t log_size = 0;
NVRTC_CHECK(nvrtcGetProgramLogSize(prog, &log_size));
std::string log(log_size, '\0');
if (log_size) NVRTC_CHECK(nvrtcGetProgramLog(prog, log.data()));
if (rc != NVRTC_SUCCESS) throw std::runtime_error(log);

size_t ptx_size = 0;
NVRTC_CHECK(nvrtcGetPTXSize(prog, &ptx_size));
std::vector<char> ptx(ptx_size);
NVRTC_CHECK(nvrtcGetPTX(prog, ptx.data()));

无论成功失败都应保存编译 log；warning 也可能暴露寄存器、弃用接口或架构选择问题。

模板 Kernel 与 Lowered Name

NVRTC 编译模板实例时，源码名经过 C++ mangling。流程是：

1. nvrtcAddNameExpression(prog, "&kernel<float>")；
2. 调用 nvrtcCompileProgram；
3. 在销毁 program 前调用 nvrtcGetLoweredName；
4. 用返回名字执行 cuModuleGetFunction。

name expression 必须在 compile 前注册，因为它还会触发所需模板实例化。

运行时链接

多个 JIT 单元、静态 device library 或 LTO IR 可以用 nvJitLink 组合：

user specialization PTX ─┐
prebuilt device library ──┼─ nvJitLink ─→ optimized cubin ─→ load
generated LTO IR ─────────┘

这适合数据库表达式、渲染 shader-like kernel、ML 算子特化和插件系统。链接选项、日志、输入版本与最终 cubin 都应进入诊断信息。

JIT Cache 设计

正确的 cache key 至少包含：

source bytes
+ included header bytes/version
+ compile options/macros
+ target architecture
+ NVRTC/Toolkit version
+ Driver version or codegen-relevant version
+ linked libraries/LTO inputs
+ application ABI/schema version

缓存写入应原子化，多个进程并发编译同一个 key 时要避免产物损坏。失败结果可以短期负缓存，但升级 Driver/配置后必须失效。

Driver 与 Toolkit 兼容模型

最重要的方向是：

较新 Driver 通常可运行较旧 Toolkit 构建的应用   ← backward compatibility

较新 Toolkit 产物跑在较旧 Driver：
├─ 可能落在同一 major 的 minor compatibility 范围
├─ 新特性可能仍要求对应新 Driver
└─ 新 PTX 版本可能无法被旧 Driver JIT 接受

CUDA 13.x 的精确最低 Driver 和新功能要求应查该 Toolkit 的 Release Notes/Compatibility Guide。不要把“程序能启动”理解为所有新 API、PTX ISA 和架构功能都可用。

PTX Forward Compatibility 的边界

PTX 提供面向未来 GPU 的 JIT 路径，但有两个方向的版本约束：

• 新 GPU 的新 Driver 通常能把较旧 PTX 编译到新架构；
• 新 NVRTC/NVCC 生成的新 PTX 版本，旧 Driver 可能不认识。

部署前可强制测试 PTX 路径，例如临时禁用已有 cubin 的选择，确认包中确实含可 JIT 的 PTX；同时测试冷 cache 首次启动延迟。

架构目标策略

假设部署到已知的 A、B 两代 GPU，并希望覆盖未来设备：

fatbin
├─ sm_A cubin：A 直接运行
├─ sm_B cubin：B 直接运行
└─ compute_B PTX：未来兼容 fallback

不要机械包含所有历史架构，会显著增大二进制。根据真实用户设备、Toolkit 已支持范围和需要的特性制定矩阵。

延迟加载与冷启动

进程首次 CUDA 调用可能包含 Driver 初始化、context 创建、module load、PTX JIT、library handle 初始化和内存池建立。应把启动路径拆开测量：

process start
  → cuInit/context
  → module/library load
  → JIT/cache lookup
  → first allocation
  → first kernel
  → steady state

CUDA Lazy Loading 可延后部分 module/kernel 加载，减少未使用代码的初始化与内存成本，但会把延迟移到首次使用位置。低延迟服务可显式预热关键路径。

Runtime/Driver 混用规则

• 先确认两边操作的是同一个 current/primary context；
• Runtime cudaStream_t 与 Driver CUstream 可互操作的场景仍要遵循类型/API 规定；
• 不在一层销毁另一层仍持有的对象；
• 错误检查分别使用 cudaGetErrorString、cuGetErrorName/cuGetErrorString；
• 库初始化和 context 切换应有清晰线程模型；
• VMM pointer 可被 Runtime launch 使用，但映射/访问权仍由 Driver 管理。

插件与 ABI 边界

跨动态库边界传递时，优先传稳定的 C ABI 数据：

struct GpuBufferView {
    void* data;
    std::uint64_t bytes;
    int device;
    std::uint32_t dtype;
    std::uint32_t abi_version;
};

还应传入或约定 caller stream，而不是插件私建默认 stream。不要跨编译器/标准库边界传 std::vector、异常对象或拥有型 C++ CUDA wrapper。

部署检查表

构建阶段

• 明确 host compiler、C++ 标准、Toolkit 和目标 SM；
• 保存 nvcc --version、编译参数和 git revision；
• 检查 fatbin 中的 cubin/PTX 目标；
• 对运行时编译保存 source hash、日志与 cache key；
• 静态/动态链接方式满足 CUDA library 的许可与依赖要求。

发布阶段

• 列出最低 Driver、GPU Compute Capability、OS/架构；
• 在最低 Driver 和最新 Driver 上测试；
• 覆盖有/无 JIT cache、只走 PTX fallback、只走 cubin；
• 验证容器是否挂载 host Driver library，而不是把内核 Driver 塞进镜像；
• 确认 MIG/MPS、多进程与显存限额环境下的行为。

运行阶段

• 日志记录 device UUID/name、CC、Driver、Runtime、加载代码目标；
• 所有 feature 先 query，再选择 fallback；
• 对 JIT、module load、allocation 和首个 kernel 分段计时；
• 异步错误在同步边界附带最近操作上下文；
• 不把用户机器上的 JIT cache 当作唯一产物存储。

故障定位

练习

1. 用 NVRTC 编译模板 SAXPY，通过 lowered name 加载 float 和 double 两个实例。
2. 生成 sm_XX cubin 与 compute_XX PTX，分别测 module load 和首次 launch 时间。
3. 写一个插件接受 caller 的 device pointer 与 stream，验证 primary context 与错误 context 两条路径。
4. 设计一个 JIT cache，并用两个并发进程测试相同 key 的原子写入。

官方参考

• CUDA Programming Guide: The CUDA Driver API
• CUDA Driver API Reference
• NVRTC Documentation
• nvJitLink Documentation
• NVCC Compiler Driver
• CUDA Compatibility
• CUDA Toolkit Release Notes