nvJitLink

nvJitLink库的用户指南。

1. 简介

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

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

当提供包含LTO-IR的LTO中间表示或更高级格式时，也可以执行链接时优化。

如果输入内容不包含GPU汇编代码，则会先进行编译再链接。

该库的功能与CUDA驱动中的cuLink* API类似，具有以下优势：

• cuLink* API 已弃用与 LTO-IR 配合使用

• 支持链接时间优化

• 允许用户使用运行时链接功能，该功能支持作为CUDA Toolkit发布的一部分的最新工具包版本。如果应用程序在系统中安装了较旧的驱动程序上运行，则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_用户指南.pdf

3. 用户界面

本章介绍JIT Link API。API的基本用法在Basic Usage中有详细说明。

• 错误代码

• 链接

• 支持的链接选项

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的输入类型。

Functions

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是链接的基本单位，也是程序的不透明句柄。

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 将数据映像添加到链接中。

Parameters

• handle – [输入] nvJitLink句柄。

• inputType – [in] 输入类型。

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

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

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

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkAddFile 从文件中读取数据并进行链接。

Parameters

• handle – [输入] nvJitLink句柄。

• inputType – [in] 输入类型。

• fileName – [in] 文件名。

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

static inline nvJitLinkResult nvJitLinkComplete(nvJitLinkHandle handle)

nvJitLinkComplete 执行实际的链接操作。

Parameters

handle – [输入] nvJitLink句柄。

Returns

• 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

Parameters

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

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

• options – [输入] 由numOptions个选项字符串组成的数组。

Returns

• 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

Parameters

handle – [输入] nvJitLink句柄的地址。

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkGetErrorLog 将任何错误信息记录到日志中。

用户需负责分配足够的空间来存储log。

另请参阅

nvJitLinkGetErrorLogSize

Parameters

• handle – [输入] nvJitLink句柄。

• log – [out] 错误日志。

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkGetErrorLogSize 获取错误日志的大小。

另请参阅

nvJitLinkGetErrorLog

Parameters

• handle – [输入] nvJitLink句柄。

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

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

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

用户需负责分配足够的空间来存储log。

另请参阅

nvJitLinkGetInfoLogSize

Parameters

• handle – [输入] nvJitLink句柄。

• log – [out] 信息日志。

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkGetInfoLogSize 获取信息日志的大小。

另请参阅

nvJitLinkGetInfoLog

Parameters

• handle – [输入] nvJitLink句柄。

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

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkGetLinkedCubin 获取已链接的cubin文件。

用户需负责分配足够的空间来存储cubin。

另请参阅

nvJitLinkGetLinkedCubinSize

Parameters

• handle – [输入] nvJitLink句柄。

• cubin – [out] 链接后的cubin文件。

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkGetLinkedCubinSize 获取已链接cubin的大小。

另请参阅

nvJitLinkGetLinkedCubin

Parameters

• handle – [输入] nvJitLink句柄。

• size – [out] 关联cubin文件的大小。

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkGetLinkedPtx 获取已链接的ptx。

链接PTX仅在启用-lto选项时可用。用户需自行分配足够空间来存储ptx。

另请参阅

nvJitLinkGetLinkedPtxSize

Parameters

• handle – [输入] nvJitLink句柄。

• ptx – [输出] 链接的PTX。

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkGetLinkedPtxSize 获取已链接ptx的大小。

链接PTX仅在启用-lto选项时可用。

另请参阅

nvJitLinkGetLinkedPtx

Parameters

• handle – [输入] nvJitLink句柄。

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

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

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

nvJitLinkVersion 返回当前 nvJitLink 的版本。

Parameters

• major – [out] 主版本号。

• minor – [out] 次要版本号。

Returns

• NVJITLINK_SUCCESS

• NVJITLINK_ERROR_INVALID_INPUT

• NVJITLINK_ERROR_INTERNAL

3.2.3. 类型定义

typedef struct nvJitLink *nvJitLinkHandle

nvJitLinkHandle是链接的单位，也是程序的不透明句柄。

要链接输入，首先需要使用nvJitLinkCreate()创建一个nvJitLinkHandle实例。

3.3. 支持的链接选项

nvJitLink支持以下链接选项。

选项名称以单个短横线(-)为前缀。需要赋值的选项会带有一个等号(=)后接选项值，中间没有空格，例如"-arch=sm_90"。

支持的选项包括：

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

• -maxrregcount= 最大寄存器数量。

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

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

• -lto 执行链接时优化。

• -ptx 在链接后生成ptx而非cubin；仅在使用-lto时支持

• -O 优化级别。仅支持0和3。

• -g 生成调试信息。

• -lineinfo 生成行信息。

• -ftz= 刷新到零。

• -prec-div= 精确除法。

• -prec-sqrt= 精确平方根。

• -fma= 快速乘法加法。

• -kernels-used= 传入使用的内核列表；不在列表中的内核将被移除。该选项可以多次指定。

• -variables-used= 传递已使用的变量列表；不在列表中的变量将被移除。该选项可以多次指定。

• -optimize-unused-variables 通常设备代码优化会受到不知道主机代码引用哪些变量的限制。使用此选项后，可以假设如果某个变量在设备代码中未被引用，则可以将其移除。

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

• -split-compile= 拆分编译的最大线程数。使用0表示使用所有可用处理器。值为1时禁用拆分编译（默认值）。

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

• -jump-table-density= 在执行LTO时，指定switch语句中的分支密度百分比，并将其作为决定是否使用跳转表(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命令创建。我们可以按照Figure 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版本（x >= 1）替换并使用nvJitLink共享库。反之，您不能使用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源文件（例如lto_89和lto_90）可以正常工作，只要最终链接是所有被链接架构中最新的。也就是说，对于任何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公司（“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对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。

6.4.2. OpenCL

OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。

6.4.3. 商标

NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。

© 2022-2022 NVIDIA公司及附属机构。保留所有权利。