nvFatbin

nvFatbin库用户指南。

1. 简介

Fatbin Creator API 是一组可在运行时使用的 API，用于将多个 CUDA 对象合并为一个 CUDA 胖二进制文件（fatbin）。

这些API接受多种格式的输入，包括设备cubin、PTX或LTO-IR。输出是一个fatbin，可以通过CUDA Driver API的cuModuleLoadData加载。

该库的功能类似于CUDA工具包中的fatbinary离线工具，具有以下优势：

• 支持运行时创建fatbin。

• 客户端可以对输入过程进行细粒度控制。

• 支持直接从内存输入，而无需将输入写入文件。

2. 快速入门

2.1. 系统要求

Fatbin Creator库不需要特殊的系统配置，也不需要GPU。

2.2. 安装

Fatbin Creator库是CUDA工具包发布的一部分，其组件在CUDA工具包安装目录中的组织结构如下：

• 在Windows上：

  • include\nvFatbin.h

  • lib\x64\nvFatbin.dll

  • lib\x64\nvFatbin_static.lib

  • doc\pdf\nvFatbin_User_Guide.pdf

• 在Linux系统上：

  • include/nvFatbin.h

  • lib64/libnvfatbin.so

  • lib64/libnvfatbin_static.a

  • doc/pdf/nvFatbin_User_Guide.pdf

3. 用户界面

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

• 错误代码

• 创建

• 支持的选项

3.1. 错误代码

枚举

nvFatbinResult

枚举类型nvFatbinResult定义了API调用的结果代码。

Functions

const char * nvFatbinGetErrorString(nvFatbinResult result)

nvFatbinGetErrorString 为每个错误代码返回对应的错误描述字符串。

3.1.1. 枚举

enum nvFatbinResult

枚举类型nvFatbinResult定义了API调用的结果代码。

nvFatbin API返回nvFatbinResult代码以指示结果。

取值：

enumerator NVFATBIN_SUCCESS

enumerator NVFATBIN_ERROR_INTERNAL

enumerator NVFATBIN_ERROR_ELF_ARCH_MISMATCH

enumerator NVFATBIN_ERROR_ELF_SIZE_MISMATCH

enumerator NVFATBIN_ERROR_MISSING_PTX_VERSION

enumerator NVFATBIN_ERROR_NULL_POINTER

enumerator NVFATBIN_ERROR_COMPRESSION_FAILED

enumerator NVFATBIN_ERROR_COMPRESSED_SIZE_EXCEEDED

enumerator NVFATBIN_ERROR_UNRECOGNIZED_OPTION

enumerator NVFATBIN_ERROR_INVALID_ARCH

enumerator NVFATBIN_ERROR_INVALID_NVVM

enumerator NVFATBIN_ERROR_EMPTY_INPUT

enumerator NVFATBIN_ERROR_MISSING_PTX_ARCH

enumerator NVFATBIN_ERROR_PTX_ARCH_MISMATCH

enumerator NVFATBIN_ERROR_MISSING_FATBIN

enumerator NVFATBIN_ERROR_INVALID_INDEX

enumerator NVFATBIN_ERROR_IDENTIFIER_REUSE

enumerator NVFATBIN_ERROR_INTERNAL_PTX_OPTION

3.1.2. 功能

const char *nvFatbinGetErrorString(nvFatbinResult result)

nvFatbinGetErrorString 为每个错误代码返回对应的错误描述字符串。

Parameters

result – [输入] 错误代码

Returns

• 如果结果为 NVFATBIN_SUCCESS，则返回 nullptr

• 一个字符串，如果结果不是NVFATBIN_SUCCESS

3.2. Fatbinary创建

Functions

nvFatbinResult nvFatbinAddCubin(nvFatbinHandle handle, const void *code, size_t size, const char *arch, const char *identifier)

nvFatbinAddCubin 将 CUDA 二进制文件添加到 fatbinary 中。

nvFatbinResult nvFatbinAddIndex(nvFatbinHandle handle, const void *code, size_t size, const char *identifier)

nvFatbinAddIndex 向 fatbinary 添加一个索引文件。

nvFatbinResult nvFatbinAddLTOIR(nvFatbinHandle handle, const void *code, size_t size, const char *arch, const char *identifier, const char *optionsCmdLine)

nvFatbinAddLTOIR 将LTOIR添加到fatbinary中。

nvFatbinResult nvFatbinAddPTX(nvFatbinHandle handle, const char *code, size_t size, const char *arch, const char *identifier, const char *optionsCmdLine)

nvFatbinAddPTX 将PTX添加到fatbinary中。

nvFatbinResult nvFatbinAddReloc(nvFatbinHandle handle, const void *code, size_t size)

nvFatbinAddReloc 将可重定位的PTX条目从主机对象添加到fatbinary中。

nvFatbinResult nvFatbinCreate(nvFatbinHandle *handle_indirect, const char **options, size_t optionsCount)

nvFatbinCreate 创建一个新的句柄

nvFatbinResult nvFatbinDestroy(nvFatbinHandle *handle_indirect)

nvFatbinDestroy 销毁该句柄。

nvFatbinResult nvFatbinGet(nvFatbinHandle handle, void *buffer)

nvFatbinGet 返回已完成的 fatbinary。

nvFatbinResult nvFatbinSize(nvFatbinHandle handle, size_t *size)

nvFatbinSize 返回 fatbinary 的大小。

nvFatbinResult nvFatbinVersion(unsigned int *major, unsigned int *minor)

nvFatbinVersion 返回当前 nvFatbin 的版本

类型定义

nvFatbinHandle

nvFatbinHandle是fatbin创建的基本单元，也是程序的不透明句柄。

3.2.1. 功能

nvFatbinResult nvFatbinAddCubin(nvFatbinHandle handle, const void *code, size_t size, const char *arch, const char *identifier)

nvFatbinAddCubin 向 fatbinary 添加一个 CUDA 二进制文件。

用户需负责确保所有字符串格式正确。

Parameters

• handle – [in] nvFatbin句柄。

• code – [in] 该cubin文件。

• size – [in] cubin文件的大小。

• arch – [输入] 该cubin文件所对应的架构。

• identifier – [输入] cubin的名称，在使用cuobjdump等工具提取fatbin时很有用。

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_INVALID_ARCH

• NVFATBIN_ERROR_ELF_ARCH_MISMATCH

• NVFATBIN_ERROR_ELF_SIZE_MISMATCH

• NVFATBIN_ERROR_COMPRESSION_FAILED,

• NVFATBIN_ERROR_UNRECOGNIZED_OPTION

• NVFATBIN_ERROR_COMPRESSED_SIZE_EXCEEDED

• NVFATBIN_ERROR_EMPTY_INPUT

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinAddIndex(nvFatbinHandle handle, const void *code, size_t size, const char *identifier)

nvFatbinAddIndex 向 fatbinary 添加一个索引文件。

用户需负责确保所有字符串格式正确。

Parameters

• handle – [in] nvFatbin句柄。

• code – [输入] 索引值。

• size – [in] 索引的大小。

• identifier – [输入] 索引名称，在使用cuobjdump等工具提取fatbin时很有用。

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_INVALID_INDEX

• NVFATBIN_ERROR_COMPRESSION_FAILED,

• NVFATBIN_ERROR_UNRECOGNIZED_OPTION

• NVFATBIN_ERROR_COMPRESSED_SIZE_EXCEEDED

• NVFATBIN_ERROR_EMPTY_INPUT

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinAddLTOIR(nvFatbinHandle handle, const void *code, size_t size, const char *arch, const char *identifier, const char *optionsCmdLine)

nvFatbinAddLTOIR 将LTOIR添加到fatbinary中。

用户需负责确保所有字符串格式正确。

Parameters

• handle – [in] nvFatbin句柄。

• code – [in] LTOIR代码。

• size – [in] LTOIR代码的大小。

• arch – [输入] 此LTOIR所对应的架构。

• identifier – [in] LTOIR的名称，在使用cuobjdump等工具提取fatbin时很有用。

• optionsCmdLine – [输入] JIT编译期间使用的选项。

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_NULL_POINTER

• NVFATBIN_ERROR_INVALID_ARCH

• NVFATBIN_ERROR_COMPRESSION_FAILED,

• NVFATBIN_ERROR_UNRECOGNIZED_OPTION

• NVFATBIN_ERROR_COMPRESSED_SIZE_EXCEED

• NVFATBIN_ERROR_EMPTY_INPUT

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinAddPTX(nvFatbinHandle handle, const char *code, size_t size, const char *arch, const char *identifier, const char *optionsCmdLine)

nvFatbinAddPTX 将PTX添加到fatbinary中。

用户需确保所有字符串格式正确。大小应包含终止空字符('\0')。如果最后一个字符不是'\0'，系统会自动添加一个，但在此过程中，若代码尚未复制，则会被复制。

Parameters

• handle – [in] nvFatbin句柄。

• code – [in] PTX代码。

• size – [in] PTX代码的大小。

• arch – [输入] 该PTX所对应的架构。

• identifier – [输入] PTX的名称，在使用cuobjdump等工具提取fatbin时很有用。

• optionsCmdLine – [输入] JIT编译期间使用的选项。

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_NULL_POINTER

• NVFATBIN_ERROR_INVALID_ARCH

• NVFATBIN_ERROR_PTX_架构不匹配

• NVFATBIN_ERROR_COMPRESSION_FAILED,

• NVFATBIN_ERROR_UNRECOGNIZED_OPTION

• NVFATBIN_ERROR_COMPRESSED_SIZE_EXCEEDED

• NVFATBIN_ERROR_EMPTY_INPUT

• NVFATBIN_ERROR_MISSING_PTX_VERSION

• NVFATBIN_ERROR_MISSING_PTX_ARCH

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinAddReloc(nvFatbinHandle handle, const void *code, size_t size)

nvFatbinAddReloc 将可重定位的PTX条目从主机对象添加到fatbinary中。

请注意，每个可重定位的ptx源文件必须具有唯一标识符（标识符取自对象的条目）。这是强制要求的，因为每个唯一标识符在每个sm中只能有一个条目。另外还需注意，此操作会忽略句柄选项，而是从每个条目复制主机对象的选项。

Parameters

• handle – [in] nvFatbin句柄。

• code – [in] 主机对象映像。

• size – [in] 主机对象映像代码的大小。

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_NULL_POINTER

• NVFATBIN_ERROR_INVALID_ARCH

• NVFATBIN_ERROR_PTX_架构不匹配

• NVFATBIN_ERROR_COMPRESSION_FAILED,

• NVFATBIN_ERROR_UNRECOGNIZED_OPTION

• NVFATBIN_ERROR_COMPRESSED_SIZE_EXCEEDED

• NVFATBIN_ERROR_EMPTY_INPUT

• NVFATBIN_ERROR_MISSING_PTX_VERSION

• NVFATBIN_ERROR_MISSING_PTX_ARCH

• NVFATBIN_ERROR_IDENTIFIER_REUSE

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinCreate(nvFatbinHandle *handle_indirect, const char **options, size_t optionsCount)

nvFatbinCreate 创建一个新的句柄

Parameters

• handle_indirect – [输出] nvFatbin句柄的地址

• options – [输入] 一个字符串数组，每个元素包含一个选项。

• optionsCount – [in] 选项数量。

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_NULL_POINTER

• NVFATBIN_ERROR_UNRECOGNIZED_OPTION

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinDestroy(nvFatbinHandle *handle_indirect)

nvFatbinDestroy 销毁该句柄。

调用此函数后，再使用任何其他指向该句柄的指针将导致未定义行为。传入的句柄将被设置为nullptr。

Parameters

handle_indirect – [输入] 指向句柄的指针。

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_NULL_POINTER

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinGet(nvFatbinHandle handle, void *buffer)

nvFatbinGet返回已完成的fatbinary。

用户需负责确保缓冲区大小适合fatbinary。在使用此功能前必须调用nvFatbinSize，否则将返回错误。

另请参阅

nvFatbinSize

Parameters

• handle – [in] nvFatbin句柄。

• buffer – [输出] 用于存储fatbinary的内存。

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_NULL_POINTER

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinSize(nvFatbinHandle handle, size_t *size)

nvFatbinSize 返回 fatbinary 的大小。

Parameters

• handle – [in] nvFatbin句柄。

• size – [输出] fatbinary的大小

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_NULL_POINTER

• NVFATBIN_ERROR_INTERNAL

nvFatbinResult nvFatbinVersion(unsigned int *major, unsigned int *minor)

nvFatbinVersion 返回当前 nvFatbin 的版本号

Parameters

• major – [out] 主版本号。

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

Returns

• NVFATBIN_SUCCESS

• NVFATBIN_ERROR_NULL_POINTER

• NVFATBIN_ERROR_INTERNAL

3.2.2. 类型定义

typedef struct _nvFatbinHandle *nvFatbinHandle

nvFatbinHandle是fatbin创建的基本单位，也是程序的不透明句柄。

要创建一个fatbin，首先必须使用nvFatbinCreate()创建一个nvFatbinHandle实例。

3.3. 支持的选项

nvFatbin支持以下选项。

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

支持的选项包括：

• -32 将条目设为32位。

• -64 将条目设置为64位。

• -c 无实际效果。（已弃用，将在下一个主要版本中移除。从一开始就未产生任何作用。）

• -compress= 启用(true)/禁用(false)压缩功能(默认值: true)。

• -compress-all 压缩fatbin中的所有内容，即使内容很小。

• -compress-mode= 选择压缩行为。有效选项包括"default"（默认）、"size"（更小的fatbin体积）、"speed"（更快的解压速度，当前与默认相同）、"balance"（在体积和速度之间取得平衡）以及"none"（不压缩，类似于-compress=false但优先级高于-compress=true）。

• -cuda 指定使用CUDA（而非OpenCL）。

• -g 生成调试信息。

• -host= 指定主机操作系统。有效选项为"linux"、"windows"和"mac"(已弃用)。

• -opencl 指定使用OpenCL（而非CUDA）。

4. 基础用法

本文档的这一部分通过一个简单示例，说明如何使用Fatbin Creator API来链接程序。为简洁和可读性起见，未展示对API返回值的错误检查。

本示例假设我们需要创建一个包含sm_52架构CUBIN、sm_61架构PTX以及sm_70架构LTOIR的fatbin文件。我们可以创建一个fatbin生成器实例，并获取其API句柄，如图1所示。

图1. Fatbin Creator创建并初始化程序

nvFatbinHandle handle;
nvFatbinCreate(&handle, nullptr, 0);

假设我们已经将三个输入存储在std::vector中（CUBIN、PTX和LTOIR），这些输入可能来自使用nvrtc创建的代码并存储到向量中。（它们不必一定在向量中，这里只是为了说明需要数据本身及其大小。）我们可以按照图2所示添加这些输入。

图2. fatbin创建工具的输入

nvFatbinAddCubin(handle, cubin.data(), cubin.size(), "52", nullptr);
nvFatbinAddPTX(handle, ptx.data(), ptx.size(), "61", nullptr, nullptr);
nvFatbinAddLTOIR(handle, ltoir.data(), ltoir.size(), "70", nullptr, nullptr);

现在可以获取fatbin文件。为此，我们首先需要为其分配内存。而要分配内存，我们需要查询fatbin的大小，具体操作如图3所示。

图3. 生成的fatbin文件查询大小

nvFatbinSize(linker, &fatbinSize);

现在可以查询fatbin，如图4所示。然后通过将此fatbin传递给CUDA驱动API，即可在GPU上执行。

图4. 查询创建的fatbin

void* fatbin = malloc(fatbinSize);
nvFatbinGet(handle, fatbin);

当不再需要fatbin创建器时，可以按照图5所示的方式销毁它。

图5. 销毁fatbin创建器

nvFatbinDestroy(&handle);

5. 兼容性

nvFatbin库在不同版本间保持兼容。该库的主版本号必须大于或等于输入文件的最大主版本号。

例如，如果您的nvFatbin库版本至少为12.x，您可以将使用11.8创建的cubin和使用12.4创建的cubin合并生成一个fatbin。

6. 示例：运行时fatbin创建

本节演示运行时fatbin的创建。包含两个cubin文件，这些cubin是通过NVRTC在线生成的。

这两个cubin文件随后被传递给nvFatbin* API函数，该函数将这些cubin文件打包成一个fatbin。

请注意，此示例需要兼容的GPU、驱动程序及NVRTC才能运行，尽管该库本身并不依赖这些组件。

6.1. 代码 (online.cpp)

#include <nvrtc.h>
#include <cuda.h>
#include <nvFatbin.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 NVFATBIN_SAFE_CALL(x)                            \
do                                                       \
{                                                        \
   nvFatbinResult result = x;                            \
   if (result != NVFATBIN_SUCCESS)                       \
   {                                                     \
      std::cerr << "\nerror: " #x " failed with error "  \
               << nvFatbinGetErrorString(result) << '\n';\
      exit(1);                                           \
   }                                                     \
} while (0)

const char *fatbin_saxpy = "                                  \n\
__device__  float compute(float a, float x, float y) {        \n\
return a * x + y;                                             \n\
}                                                             \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";

size_t process(const void* input, const char* input_name, void** output, const char* arch)
{
// Create an instance of nvrtcProgram with the code string.
nvrtcProgram prog;
NVRTC_SAFE_CALL(
nvrtcCreateProgram(&prog,         // prog
      (const char*) input,                       // buffer
      input_name,                  // name
      0,                           // numHeaders
      NULL,                        // headers
      NULL));                      // includeNames

const char *opts[1];
opts[0] = arch;
nvrtcResult compileResult = nvrtcCompileProgram(prog,  // prog
               1,     // 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 CUBIN from the program.
size_t CUBINSize;
NVRTC_SAFE_CALL(nvrtcGetCUBINSize(prog, &CUBINSize));
char *CUBIN = new char[CUBINSize];
NVRTC_SAFE_CALL(nvrtcGetCUBIN(prog, CUBIN));
// Destroy the program.
NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));
*output = (void*) CUBIN;
return CUBINSize;
}

int main(int argc, char *argv[])
{
void* known = NULL;
size_t known_size = process(fatbin_saxpy, "fatbin_saxpy.cu", &known, "-arch=sm_52");

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));

// Dynamically determine the arch to make one of the entries of the fatbin with
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);

void* dynamic = NULL;
size_t dynamic_size = process(fatbin_saxpy, "fatbin_saxpy.cu", &dynamic, smbuf);
sprintf(smbuf, "%d", arch);

// Load the dynamic CUBIN and the statically known arch CUBIN
// and put them in a fatbin together.
nvFatbinHandle handle;
const char* fatbin_options[] = {"-cuda"};
NVFATBIN_SAFE_CALL(nvFatbinCreate(&handle, fatbin_options, 1));

NVFATBIN_SAFE_CALL(nvFatbinAddCubin(handle,
                           (void *)dynamic, dynamic_size, smbuf, "dynamic"));
NVFATBIN_SAFE_CALL(nvFatbinAddCubin(handle,
                           (void *)known, known_size, "52", "known"));

size_t fatbinSize;
NVFATBIN_SAFE_CALL(nvFatbinSize(handle, &fatbinSize));
void *fatbin = malloc(fatbinSize);
NVFATBIN_SAFE_CALL(nvFatbinGet(handle, fatbin));
NVFATBIN_SAFE_CALL(nvFatbinDestroy(&handle));

CUDA_SAFE_CALL(cuModuleLoadData(&module, fatbin));
CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "saxpy"));

// Generate input for execution, and create output buffers.
#define NUM_THREADS 128
#define NUM_BLOCKS 32
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));
delete[] hX;
delete[] hY;
delete[] hOut;
// Release resources.
free(fatbin);
delete[] ((char*)known);
delete[] ((char*)dynamic);

return 0;
}

6.2. 构建说明

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

• 使用nvFatbin共享库（注意：如果测试未使用nvrtc或运行代码，则无需链接nvrtc或CUDA驱动程序API）：

  • Windows系统:

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

  • Linux:

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

• 使用nvFatbin静态库：

  • 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\nvfatbin_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 -lnvfatbin_static -lnvptxcompiler_static -lcuda \
          -lpthread
    

6.3. 通知

6.3.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.3.2. OpenCL

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

6.3.3. 商标

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

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