nvJitLink

nvJitLink 库的用户指南。

1. 简介

JIT Link API 是一组 API，可以在运行时用于链接 GPU 设备代码。

这些 API 接受多种格式的输入，包括主机对象、主机库、fatbin（包括可重定位的 ptx）、设备 cubin、PTX、索引文件或 LTO-IR。 输出是一个链接的 cubin，可以通过 CUDA 驱动程序 API 的 cuModuleLoadData 和 cuModuleLoadDataEx 加载。

当给定 LTO-IR 或包含 LTO-IR 的更高级别格式时，也可以执行链接时优化。

如果输入不包含 GPU 汇编代码，则首先编译，然后再链接。

此库中的功能类似于 CUDA 驱动程序中的 cuLink* API，但具有以下优势

• cuLink* API 已被弃用，不建议与 LTO-IR 一起使用

• 支持链接时优化

• 允许用户将运行时链接与作为 CUDA 工具包发布一部分支持的最新工具包版本一起使用。 如果应用程序在系统上安装了较旧的驱动程序，则 CUDA 驱动程序 API 中可能不提供此支持。 有关更多详细信息，请参阅 CUDA 兼容性。

• 客户端可以获得细粒度的控制，并且可以在链接期间指定低级编译器选项。

2. 入门指南

2.1. 系统要求

JIT Link 库需要以下系统配置

• 非 Windows 平台的 POSIX 线程支持。

• GPU：任何具有 CUDA 计算能力 3.5 或更高的 GPU。

• CUDA 工具包和驱动程序。

2.2. 安装

JIT Link 库是 CUDA 工具包版本的一部分，组件按如下方式组织在 CUDA 工具包安装目录中

• 在 Windows 上

  • include\nvJitLink.h

  • lib\x64\nvJitLink.dll

  • lib\x64\nvJitLink_static.lib

  • doc\pdf\nvJitLink_User_Guide.pdf

• 在 Linux 上

  • include/nvJitLink.h

  • lib64/libnvJitLink.so

  • lib64/libnvJitLink_static.a

  • doc/pdf/nvJitLink_User_Guide.pdf

3. 用户界面

本章介绍 JIT Link API。 API 的基本用法在 基本用法 中进行了解释。

• 错误代码

• 链接

• 支持的链接选项

3.1. 错误代码

枚举

nvJitLinkResult

枚举类型 nvJitLinkResult 定义 API 调用结果代码。

3.1.1. 枚举

enum nvJitLinkResult

枚举类型 nvJitLinkResult 定义 API 调用结果代码。

nvJitLink API 返回 nvJitLinkResult 代码以指示结果。

值

enumerator NVJITLINK_SUCCESS

enumerator NVJITLINK_ERROR_UNRECOGNIZED_OPTION

enumerator NVJITLINK_ERROR_MISSING_ARCH

enumerator NVJITLINK_ERROR_INVALID_INPUT

enumerator NVJITLINK_ERROR_PTX_COMPILE

enumerator NVJITLINK_ERROR_NVVM_COMPILE

enumerator NVJITLINK_ERROR_INTERNAL

enumerator NVJITLINK_ERROR_THREADPOOL

enumerator NVJITLINK_ERROR_UNRECOGNIZED_INPUT

enumerator NVJITLINK_ERROR_FINALIZE

3.2. 链接

枚举

nvJitLinkInputType

枚举类型 nvJitLinkInputType 定义可以传递给 nvJitLinkAdd* API 的输入类型。

函数

nvJitLinkResult nvJitLinkAddData(nvJitLinkHandle handle, nvJitLinkInputType inputType, const void *data, size_t size, const char *name)

nvJitLinkAddData 将数据映像添加到链接中。

nvJitLinkResult nvJitLinkAddFile(nvJitLinkHandle handle, nvJitLinkInputType inputType, const char *fileName)

nvJitLinkAddFile 从文件读取数据并将其链接到链接中。

nvJitLinkResult nvJitLinkComplete(nvJitLinkHandle handle)

nvJitLinkComplete 执行实际链接。

nvJitLinkResult nvJitLinkCreate(nvJitLinkHandle *handle, uint32_t numOptions, const char **options)

nvJitLinkCreate 使用给定的输入选项创建 nvJitLinkHandle 的实例，并设置输出参数 handle。

nvJitLinkResult nvJitLinkDestroy(nvJitLinkHandle *handle)

nvJitLinkDestroy 释放与给定句柄关联的内存，并将其设置为 NULL。

nvJitLinkResult nvJitLinkGetErrorLog(nvJitLinkHandle handle, char *log)

nvJitLinkGetErrorLog 将任何错误消息放入日志中。

nvJitLinkResult nvJitLinkGetErrorLogSize(nvJitLinkHandle handle, size_t *size)

nvJitLinkGetErrorLogSize 获取错误日志的大小。

nvJitLinkResult nvJitLinkGetInfoLog(nvJitLinkHandle handle, char *log)

nvJitLinkGetInfoLog 将任何信息消息放入日志中。

nvJitLinkResult nvJitLinkGetInfoLogSize(nvJitLinkHandle handle, size_t *size)

nvJitLinkGetInfoLogSize 获取信息日志的大小。

nvJitLinkResult nvJitLinkGetLinkedCubin(nvJitLinkHandle handle, void *cubin)

nvJitLinkGetLinkedCubin 获取链接的 cubin。

nvJitLinkResult nvJitLinkGetLinkedCubinSize(nvJitLinkHandle handle, size_t *size)

nvJitLinkGetLinkedCubinSize 获取链接的 cubin 的大小。

nvJitLinkResult nvJitLinkGetLinkedPtx(nvJitLinkHandle handle, char *ptx)

nvJitLinkGetLinkedPtx 获取链接的 ptx。

nvJitLinkResult nvJitLinkGetLinkedPtxSize(nvJitLinkHandle handle, size_t *size)

nvJitLinkGetLinkedPtxSize 获取链接的 ptx 的大小。

nvJitLinkResult nvJitLinkVersion(unsigned int *major, unsigned int *minor)

nvJitLinkVersion 返回 nvJitLink 的当前版本。

类型定义

nvJitLinkHandle

nvJitLinkHandle 是链接的单元，也是程序的 opaque 句柄。

3.2.1. 枚举

enum nvJitLinkInputType

枚举类型 nvJitLinkInputType 定义可以传递给 nvJitLinkAdd* API 的输入类型。

值

enumerator NVJITLINK_INPUT_NONE

enumerator NVJITLINK_INPUT_CUBIN

enumerator NVJITLINK_INPUT_PTX

enumerator NVJITLINK_INPUT_LTOIR

enumerator NVJITLINK_INPUT_FATBIN

enumerator NVJITLINK_INPUT_OBJECT

enumerator NVJITLINK_INPUT_LIBRARY

enumerator NVJITLINK_INPUT_INDEX

enumerator NVJITLINK_INPUT_ANY

3.2.2. 函数

static inline nvJitLinkResult nvJitLinkAddData(nvJitLinkHandle handle, nvJitLinkInputType inputType, const void *data, size_t size, const char *name)

nvJitLinkAddData 将数据映像添加到链接中。

参数

• handle – [in] nvJitLink 句柄。

• inputType – [in] 输入类型。

• data – [in] 指向内存中数据映像的指针。

• size – [in] 数据的大小。

• name – [in] 输入对象的名称。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkAddFile(nvJitLinkHandle handle, nvJitLinkInputType inputType, const char *fileName)

nvJitLinkAddFile 从文件读取数据并将其链接到链接中。

参数

• handle – [in] nvJitLink 句柄。

• inputType – [in] 输入类型。

• fileName – [in] 文件名。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkComplete(nvJitLinkHandle handle)

nvJitLinkComplete 执行实际链接。

参数

handle – [in] nvJitLink 句柄。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkCreate(nvJitLinkHandle *handle, uint32_t numOptions, const char **options)

nvJitLinkCreate 使用给定的输入选项创建 nvJitLinkHandle 的实例，并设置输出参数 handle。

它支持 支持的链接选项 中列出的选项。

另请参阅

nvJitLinkDestroy

参数

• handle – [out] nvJitLink 句柄的地址。

• numOptions – [in] 传递的选项数量。

• options – [in] 大小为 numOptions 的选项字符串数组。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_UNRECOGNIZED_OPTION

• NVJITLINK_ERROR_MISSING_ARCH

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkDestroy(nvJitLinkHandle *handle)

nvJitLinkDestroy 释放与给定句柄关联的内存，并将其设置为 NULL。

另请参阅

nvJitLinkCreate

参数

handle – [in] nvJitLink 句柄的地址。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkGetErrorLog(nvJitLinkHandle handle, char *log)

nvJitLinkGetErrorLog 将任何错误消息放入日志中。

用户负责分配足够的空间来容纳 log。

另请参阅

nvJitLinkGetErrorLogSize

参数

• handle – [in] nvJitLink 句柄。

• log – [out] 错误日志。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkGetErrorLogSize(nvJitLinkHandle handle, size_t *size)

nvJitLinkGetErrorLogSize 获取错误日志的大小。

另请参阅

nvJitLinkGetErrorLog

参数

• handle – [in] nvJitLink 句柄。

• size – [out] 错误日志的大小。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkGetInfoLog(nvJitLinkHandle handle, char *log)

nvJitLinkGetInfoLog 将任何信息消息放入日志中。

用户负责分配足够的空间来容纳 log。

另请参阅

nvJitLinkGetInfoLogSize

参数

• handle – [in] nvJitLink 句柄。

• log – [out] 信息日志。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkGetInfoLogSize(nvJitLinkHandle handle, size_t *size)

nvJitLinkGetInfoLogSize 获取信息日志的大小。

另请参阅

nvJitLinkGetInfoLog

参数

• handle – [in] nvJitLink 句柄。

• size – [out] 信息日志的大小。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkGetLinkedCubin(nvJitLinkHandle handle, void *cubin)

nvJitLinkGetLinkedCubin 获取链接的 cubin。

用户负责分配足够的空间来容纳 cubin。

另请参阅

nvJitLinkGetLinkedCubinSize

参数

• handle – [in] nvJitLink 句柄。

• cubin – [out] 链接的 cubin。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkGetLinkedCubinSize(nvJitLinkHandle handle, size_t *size)

nvJitLinkGetLinkedCubinSize 获取链接的 cubin 的大小。

另请参阅

nvJitLinkGetLinkedCubin

参数

• handle – [in] nvJitLink 句柄。

• size – [out] 链接的 cubin 的大小。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkGetLinkedPtx(nvJitLinkHandle handle, char *ptx)

nvJitLinkGetLinkedPtx 获取链接的 ptx。

仅在使用 -lto 选项时，链接的 PTX 才可用。 用户负责分配足够的空间来容纳 ptx。

另请参阅

nvJitLinkGetLinkedPtxSize

参数

• handle – [in] nvJitLink 句柄。

• ptx – [out] 链接的 PTX。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkGetLinkedPtxSize(nvJitLinkHandle handle, size_t *size)

nvJitLinkGetLinkedPtxSize 获取链接的 ptx 的大小。

仅在使用 -lto 选项时，链接的 PTX 才可用。

另请参阅

nvJitLinkGetLinkedPtx

参数

• handle – [in] nvJitLink 句柄。

• size – [out] 链接的 PTX 的大小。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

nvJitLinkResult nvJitLinkVersion(unsigned int *major, unsigned int *minor)

nvJitLinkVersion 返回 nvJitLink 的当前版本。

参数

• major – [out] 主版本号。

• minor – [out] 次版本号。

返回值

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

3.2.3. 类型定义

typedef struct nvJitLink *nvJitLinkHandle

nvJitLinkHandle 是链接的单元，也是程序的 opaque 句柄。

要链接输入，必须首先使用 nvJitLinkCreate() 创建 nvJitLinkHandle 的实例。

3.3. 支持的链接选项

nvJitLink 支持以下链接选项。

选项名称以单破折号 (-) 为前缀。 采用值的选项具有赋值运算符 (=)，后跟选项值，没有空格，例如 "-arch=sm_90"。

支持的选项包括

• -arch=sm_<N> 传递 SM 架构值。有关 <N> 的有效值，请参阅 nvcc。如果仅生成 PTX，则可以使用 compute_<N> 值代替。这是一个必需选项。

• -maxrregcount=<N> 最大寄存器计数。

• -time 将计时信息打印到 InfoLog。

• -verbose 将详细消息打印到 InfoLog。

• -lto 执行链接时优化。

• -ptx 在链接后而不是 cubin 后发出 ptx；仅在使用 -lto 时支持

• -O<N> 优化级别。仅接受 0 和 3。

• -g 生成调试信息。

• -lineinfo 生成行信息。

• -ftz=<n> 刷新为零。

• -prec-div=<n> 精确除法。

• -prec-sqrt=<n> 精确平方根。

• -fma=<n> 快速乘加。

• -kernels-used=<name> 传递已使用的内核列表；列表中未包含的任何内核都可以删除。此选项可以多次指定。

• -variables-used=<name> 传递已使用的变量列表；列表中未包含的任何变量都可以删除。此选项可以多次指定。

• -optimize-unused-variables 通常，设备代码优化受限于不知道主机代码引用的内容。使用此选项，它可以假定如果变量在设备代码中未被引用，则可以将其删除。

• -Xptxas=<opt> 将 <opt> 传递给 ptxas。此选项可以多次调用。

• -split-compile=<N> 拆分编译最大线程计数。使用 0 以使用所有可用的处理器。值 1 禁用拆分编译（默认）。

• -split-compile-extended=<N> 一种更激进的拆分编译形式，仅在 LTO 模式下可用。接受最大线程计数值。使用 0 以使用所有可用的处理器。值 1 禁用扩展拆分编译（默认）。注意：此选项可能会影响已编译二进制文件的性能。

• -jump-table-density=<N> 当执行 LTO 时，指定 switch 语句中的 case 密度百分比，并将其用作确定是否将使用跳转表（brx.idx 指令）来实现 switch 语句的最小阈值。默认值为 101。百分比范围为 0 到 101（含）。

• -no-cache 不缓存 nvJitLink 的中间步骤。

• -device-stack-protector 在设备代码中启用堆栈金丝雀。堆栈金丝雀使利用某些类型的涉及堆栈局部变量的内存安全漏洞变得更加困难。编译器使用启发式方法来评估每个函数中此类漏洞的风险。只有那些被认为是高风险的函数才会使用堆栈金丝雀。

4. 基本用法

本文档的此部分使用一个简单的示例来说明如何使用 JIT Link API 链接程序。为了简洁和可读性，未显示 API 返回值的错误检查。

此示例假设我们要为 sm_80 链接，但应使用系统上安装的任何架构。我们可以创建链接器并获取其句柄，如图 1所示。

图 1. 链接器创建和程序初始化

nvJitLink_t linker;
const char* link_options[] = { "-arch=sm_80" };
nvJitLinkCreate(&linker, 1, link_options);

假设我们已经有两个可重定位的输入文件（a.o 和 b.o），可以使用 nvcc -dc 命令创建。我们可以添加输入文件，如图 2所示。

图 2. 链接器的输入

nvJitLinkAddFile(linker, NVJITLINK_INPUT_OBJECT, "a.o");
nvJitLinkAddFile(linker, NVJITLINK_INPUT_OBJECT, "b.o");

现在可以完成实际链接，如图 3所示。

图 3. PTX 程序的链接

nvJitLinkComplete(linker);

现在可以获得链接的 GPU 汇编代码。为了获得此代码，我们首先为其分配内存。为了分配内存，我们需要查询链接的 GPU 汇编代码映像的大小，如图 4所示。

图 4. 查询链接的汇编映像的大小

nvJitLinkGetLinkedCubinSize(linker, &cubinSize);

现在可以查询链接的 GPU 汇编代码映像，如图 5所示。然后可以通过将此映像传递给 CUDA 驱动程序 API 在 GPU 上执行此映像。

图 5. 查询链接的汇编映像

elf = (char*) malloc(cubinSize);
nvJitLinkGetLinkedCubin(linker, (void*)elf);

当不再需要链接器时，可以销毁它，如图 6所示。

图 6. 销毁链接器

nvJitLinkDestroy(&linker);

5. 兼容性

nvJitLink 库在同一发行版的小版本之间兼容，但在主版本之间可能不兼容。库版本本身必须 >= 输入的最大版本，并且共享库版本必须 >= 与其链接的版本。

例如，如果您的 nvJitLink 库版本为 12.x，其中 x >= 1，则可以链接使用 12.0 创建的对象和使用 12.1 创建的对象。如果它是与 12.1 链接的，那么您可以替换并使用任何版本为 12.x 的 nvJitLink 共享库，其中 x >= 1。另一方面，您不能使用 12.0 链接 12.1 对象，也不能使用 12.0 nvJitLink 库来运行 12.1 代码。

跨主版本（例如 11.x 与 12.x）的链接适用于 ELF 和 PTX 输入，但不适用于 LTOIR 输入。如果使用 LTO，则仅在主版本内保证兼容性。

将扩展 ISA 源（如 sm_90a）与任何其他 sm 版本链接将始终失败。

只要最终链接是所有被链接的架构中最新的，链接来自不同架构（例如 compute_89 和 compute_90）的 PTX 源就可以工作。也就是说，对于任何 compute_X 和 compute_Y，如果目标是 sm_N，其中 N >= max(X,Y)，则链接有效。

只要最终链接是所有被链接的架构中最新的，链接来自不同架构（例如 lto_89 和 lto_90）的 LTO 源就可以工作。也就是说，对于任何 lto_X 和 lto_Y，如果目标是 sm_N，其中 N >= max(X,Y)，则链接有效。

与非 PTX、非 LTO 源的链接仅限于链接兼容的架构，例如 sm_70 和 sm_75 可以相互链接，但不能与 sm_80 链接。

6. 示例：设备 LTO（链接时优化）

本节演示设备链接时优化 (LTO)。有两个 LTO IR 单元。第一个单元是使用 nvcc 离线生成的，方法是将架构指定为 '-arch lto_XX'（请参阅 offline.cu）。生成的 LTO IR 打包在 fatbinary 中。

第二个单元是使用 NVRTC 在线生成的，方法是指定标志 '-dlto'（请参阅 online.cpp）。

然后将这两个单元传递给 libnvJitLink* API 函数，这些函数将 LTO IR 链接在一起，在链接的 IR 上运行优化器，并生成 cubin（请参阅 online.cpp）。然后将 cubin 加载到 GPU 上并执行。

6.1. 代码 (offline.cu)

__device__  float compute(float a, float x, float y) {
  return a * x + y;
}

6.2. 代码 (online.cpp)

#include <nvrtc.h>
#include <cuda.h>
#include <nvJitLink.h>
#include <nvrtc.h>
#include <iostream>

#define NUM_THREADS 128
#define NUM_BLOCKS 32

#define NVRTC_SAFE_CALL(x)                                        \
  do {                                                            \
    nvrtcResult result = x;                                       \
    if (result != NVRTC_SUCCESS) {                                \
      std::cerr << "\nerror: " #x " failed with error "           \
                << nvrtcGetErrorString(result) << '\n';           \
      exit(1);                                                    \
    }                                                             \
  } while(0)
#define CUDA_SAFE_CALL(x)                                         \
  do {                                                            \
    CUresult result = x;                                          \
    if (result != CUDA_SUCCESS) {                                 \
      const char *msg;                                            \
      cuGetErrorName(result, &msg);                               \
      std::cerr << "\nerror: " #x " failed with error "           \
                << msg << '\n';                                   \
      exit(1);                                                    \
    }                                                             \
  } while(0)
#define NVJITLINK_SAFE_CALL(h,x)                                  \
  do {                                                            \
    nvJitLinkResult result = x;                                   \
    if (result != NVJITLINK_SUCCESS) {                            \
      std::cerr << "\nerror: " #x " failed with error "           \
                << result << '\n';                                \
      size_t lsize;                                               \
      result = nvJitLinkGetErrorLogSize(h, &lsize);               \
      if (result == NVJITLINK_SUCCESS && lsize > 0) {             \
        char *log = (char*)malloc(lsize);                         \
    result = nvJitLinkGetErrorLog(h, log);                        \
    if (result == NVJITLINK_SUCCESS) {                            \
      std::cerr << "error: " << log << '\n';                      \
      free(log);                                                  \
    }                                                             \
      }                                                           \
      exit(1);                                                    \
    }                                                             \
  } while(0)

const char *lto_saxpy = "                                       \n\
extern __device__ float compute(float a, float x, float y);     \n\
                                                                \n\
extern \"C\" __global__                                         \n\
void saxpy(float a, float *x, float *y, float *out, size_t n)   \n\
{                                                               \n\
  size_t tid = blockIdx.x * blockDim.x + threadIdx.x;           \n\
  if (tid < n) {                                                \n\
    out[tid] = compute(a, x[tid], y[tid]);                      \n\
  }                                                             \n\
}                                                               \n";

int main(int argc, char *argv[])
{
  size_t numBlocks = 32;
  size_t numThreads = 128;
  // Create an instance of nvrtcProgram with the code string.
  nvrtcProgram prog;
  NVRTC_SAFE_CALL(
    nvrtcCreateProgram(&prog,                       // prog
                       lto_saxpy,                   // buffer
                       "lto_saxpy.cu",              // name
                       0,                           // numHeaders
                       NULL,                        // headers
                       NULL));                      // includeNames

  // specify that LTO IR should be generated for LTO operation
  const char *opts[] = {"-dlto",
                        "--relocatable-device-code=true"};
  nvrtcResult compileResult = nvrtcCompileProgram(prog,  // prog
                                                  2,     // numOptions
                                                  opts); // options
  // Obtain compilation log from the program.
  size_t logSize;
  NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
  char *log = new char[logSize];
  NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
  std::cout << log << '\n';
  delete[] log;
  if (compileResult != NVRTC_SUCCESS) {
    exit(1);
  }
  // Obtain generated LTO IR from the program.
  size_t LTOIRSize;
  NVRTC_SAFE_CALL(nvrtcGetLTOIRSize(prog, &LTOIRSize));
  char *LTOIR = new char[LTOIRSize];
  NVRTC_SAFE_CALL(nvrtcGetLTOIR(prog, LTOIR));
  // Destroy the program.
  NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));

  CUdevice cuDevice;
  CUcontext context;
  CUmodule module;
  CUfunction kernel;
  CUDA_SAFE_CALL(cuInit(0));
  CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
  CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));

  // Load the generated LTO IR and the LTO IR generated offline
  // and link them together.
  nvJitLinkHandle handle;
  // Dynamically determine the arch to link for
  int major = 0;
  int minor = 0;
  CUDA_SAFE_CALL(cuDeviceGetAttribute(&major,
                   CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevice));
  CUDA_SAFE_CALL(cuDeviceGetAttribute(&minor,
                   CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevice));
  int arch = major*10 + minor;
  char smbuf[16];
  sprintf(smbuf, "-arch=sm_%d", arch);
  const char *lopts[] = {"-lto", smbuf};
  NVJITLINK_SAFE_CALL(handle, nvJitLinkCreate(&handle, 2, lopts));

  // NOTE: assumes "offline.fatbin" is in the current directory
  // The fatbinary contains LTO IR generated offline using nvcc
  NVJITLINK_SAFE_CALL(handle, nvJitLinkAddFile(handle, NVJITLINK_INPUT_FATBIN,
                                "offline.fatbin"));
  NVJITLINK_SAFE_CALL(handle, nvJitLinkAddData(handle, NVJITLINK_INPUT_LTOIR,
                            (void *)LTOIR, LTOIRSize, "lto_online"));

  // The call to nvJitLinkComplete causes linker to link together the two
  // LTO IR modules (offline and online), do optimization on the linked LTO IR,
  // and generate cubin from it.
  NVJITLINK_SAFE_CALL(handle, nvJitLinkComplete(handle));
  size_t cubinSize;
  NVJITLINK_SAFE_CALL(handle, nvJitLinkGetLinkedCubinSize(handle, &cubinSize));
  void *cubin = malloc(cubinSize);
  NVJITLINK_SAFE_CALL(handle, nvJitLinkGetLinkedCubin(handle, cubin));
  NVJITLINK_SAFE_CALL(handle, nvJitLinkDestroy(&handle));
  CUDA_SAFE_CALL(cuModuleLoadData(&module, cubin));
  CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "saxpy"));

  // Generate input for execution, and create output buffers.
  size_t n = NUM_THREADS * NUM_BLOCKS;
  size_t bufferSize = n * sizeof(float);
  float a = 5.1f;
  float *hX = new float[n], *hY = new float[n], *hOut = new float[n];
  for (size_t i = 0; i < n; ++i) {
    hX[i] = static_cast<float>(i);
    hY[i] = static_cast<float>(i * 2);
  }
  CUdeviceptr dX, dY, dOut;
  CUDA_SAFE_CALL(cuMemAlloc(&dX, bufferSize));
  CUDA_SAFE_CALL(cuMemAlloc(&dY, bufferSize));
  CUDA_SAFE_CALL(cuMemAlloc(&dOut, bufferSize));
  CUDA_SAFE_CALL(cuMemcpyHtoD(dX, hX, bufferSize));
  CUDA_SAFE_CALL(cuMemcpyHtoD(dY, hY, bufferSize));
  // Execute SAXPY.
  void *args[] = { &a, &dX, &dY, &dOut, &n };
  CUDA_SAFE_CALL(
    cuLaunchKernel(kernel,
                   NUM_BLOCKS, 1, 1,    // grid dim
                   NUM_THREADS, 1, 1,   // block dim
                   0, NULL,             // shared mem and stream
                   args, 0));           // arguments
  CUDA_SAFE_CALL(cuCtxSynchronize());
  // Retrieve and print output.
  CUDA_SAFE_CALL(cuMemcpyDtoH(hOut, dOut, bufferSize));

  for (size_t i = 0; i < n; ++i) {
    std::cout << a << " * " << hX[i] << " + " << hY[i]
              << " = " << hOut[i] << '\n';
  }
  // Release resources.
  CUDA_SAFE_CALL(cuMemFree(dX));
  CUDA_SAFE_CALL(cuMemFree(dY));
  CUDA_SAFE_CALL(cuMemFree(dOut));
  CUDA_SAFE_CALL(cuModuleUnload(module));
  CUDA_SAFE_CALL(cuCtxDestroy(context));
  free(cubin);
  delete[] hX;
  delete[] hY;
  delete[] hOut;
  delete[] LTOIR;
  return 0;
}

6.3. 构建说明

假设环境变量 CUDA_PATH 指向 CUDA 工具包安装目录，请按以下方式构建此示例

• 将 offline.cu 编译为包含 LTO IR 的 fatbinary（将 lto_52 更改为不同的 lto_XX 架构，如适用）。

  nvcc -arch lto_52 -rdc=true -fatbin offline.cu
  

• 使用 nvJitLink 共享库（请注意，如果测试未使用 nvrtc，则不需要与 nvrtc 链接）

  • Windows

    cl.exe online.cpp /Feonline ^
          /I "%CUDA_PATH%\include" ^
          "%CUDA_PATH%"\lib\x64\nvrtc.lib ^
          "%CUDA_PATH%"\lib\x64\nvJitLink.lib ^
          "%CUDA_PATH%"\lib\x64\cuda.lib
    

  • Linux

    g++ online.cpp -o online \
          -I $CUDA_PATH/include \
          -L $CUDA_PATH/lib64 \
          -lnvrtc -lnvJitLink -lcuda \
          -Wl,-rpath,$CUDA_PATH/lib64
    

• 使用 nvJitLink 静态库（当与静态库链接时，还需要与 nvptxcompiler_static 链接，但这已隐式包含）

  • Windows

    cl.exe online.cpp /Feonline  ^
          /I "%CUDA_PATH%"\include ^
          "%CUDA_PATH%"\lib\x64\nvrtc_static.lib ^
          "%CUDA_PATH%"\lib\x64\nvrtc-builtins_static.lib ^
          "%CUDA_PATH%"\lib\x64\nvJitLink_static.lib ^
          "%CUDA_PATH%"\lib\x64\nvptxcompiler_static.lib ^
          "%CUDA_PATH%"\lib\x64\cuda.lib user32.lib Ws2_32.lib
    

  • Linux

    g++ online.cpp -o online \
          -I $CUDA_PATH/include \
          -L $CUDA_PATH/lib64 \
          -lnvrtc_static -lnvrtc-builtins_static -lnvJitLink_static -lnvptxcompiler_static -lcuda \
          -lpthread
    

6.4. 声明

6.4.1. 声明

本文档仅供参考，不得视为对产品的特定功能、条件或质量的保证。NVIDIA Corporation（“NVIDIA”）对本文档中包含的信息的准确性或完整性不作任何明示或暗示的陈述或保证，并且对本文档中包含的任何错误不承担任何责任。NVIDIA 对因使用此类信息或因其使用而可能导致的侵犯第三方专利或其他权利的后果或使用不承担任何责任。本文档不是对开发、发布或交付任何材料（下文定义）、代码或功能的承诺。

NVIDIA 保留随时对此文档进行更正、修改、增强、改进和任何其他更改的权利，恕不另行通知。

客户在下订单前应获取最新的相关信息，并应验证此类信息是否为最新和完整。

NVIDIA 产品的销售受订单确认时提供的 NVIDIA 标准销售条款和条件的约束，除非 NVIDIA 和客户的授权代表签署的单独销售协议（“销售条款”）另有约定。NVIDIA 特此明确反对将任何客户一般条款和条件应用于购买本文档中引用的 NVIDIA 产品。本文件未直接或间接地形成任何合同义务。

NVIDIA 产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于 NVIDIA 产品的故障或故障可能合理预期会导致人身伤害、死亡或财产或环境损害的应用。NVIDIA 对在上述设备或应用中包含和/或使用 NVIDIA 产品不承担任何责任，因此，此类包含和/或使用由客户自行承担风险。

NVIDIA 不保证或声明基于本文档的产品将适用于任何特定用途。NVIDIA 不一定对每个产品的所有参数进行测试。客户全权负责评估和确定本文档中包含的任何信息的适用性，确保产品适合客户计划的应用并满足其要求，并为该应用执行必要的测试，以避免应用或产品的缺陷。客户产品设计中的缺陷可能会影响 NVIDIA 产品的质量和可靠性，并可能导致超出本文档中包含的附加或不同的条件和/或要求。NVIDIA 对可能基于或归因于以下原因的任何缺陷、损坏、成本或问题不承担任何责任：(i) 以任何违反本文档的方式使用 NVIDIA 产品或 (ii) 客户产品设计。

本文件未授予 NVIDIA 专利权、版权或本文件项下的其他 NVIDIA 知识产权的任何明示或暗示的许可。NVIDIA 发布的有关第三方产品或服务的信息不构成 NVIDIA 授予使用此类产品或服务的许可，也不构成对其的保证或认可。使用此类信息可能需要获得第三方的专利或其他知识产权项下的第三方许可，或者获得 NVIDIA 的专利或其他知识产权项下的 NVIDIA 许可。

只有在事先获得 NVIDIA 书面批准的情况下，才允许复制本文档中的信息，并且复制必须不得更改，完整复制，完全符合所有适用的出口法律和法规，并附带所有相关的条件、限制和声明。

本文件和所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称和单独称为“材料”）均按“原样”提供。NVIDIA 对材料不作任何明示、暗示、法定或其他方面的保证，并明确否认所有关于非侵权、适销性和特定用途适用性的暗示保证。在法律未禁止的范围内，在任何情况下，NVIDIA 均不对任何损害（包括但不限于任何直接、间接、特殊、偶然、惩罚性或后果性损害）负责，无论其原因和责任理论如何，即使 NVIDIA 已被告知可能发生此类损害。尽管客户可能因任何原因遭受任何损害，NVIDIA 对本文所述产品的累计和总责任应根据产品的销售条款进行限制。

6.4.2. OpenCL

OpenCL 是 Apple Inc. 的商标，已获得 Khronos Group Inc. 的许可使用。

6.4.3. 商标

NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。

© 2022-2022 NVIDIA Corporation 及关联公司。保留所有权利。