nvFatbin
nvFatbin 库的用户指南。
1. 简介
Fatbin Creator API 是一组 API，可以在运行时用于将多个 CUDA 对象组合成一个 CUDA fat binary (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 Toolkit 发行版的一部分，组件在 CUDA 工具包安装目录中组织如下
-
在 Windows 上
include\nvFatbin.hlib\x64\nvFatbin.dlllib\x64\nvFatbin_static.libdoc\pdf\nvFatbin_User_Guide.pdf
-
在 Linux 上
include/nvFatbin.hlib64/libnvfatbin.solib64/libnvfatbin_static.adoc/pdf/nvFatbin_User_Guide.pdf
3. 用户界面
本章介绍 Fatbin Creator API。API 的基本用法在 基本用法 中进行了解释。
3.1. 错误代码
枚举
- nvFatbinResult
-
枚举类型 nvFatbinResult 定义了 API 调用结果代码。
函数
- 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
-
enumerator NVFATBIN_SUCCESS
3.1.2. 函数
-
const char *nvFatbinGetErrorString(nvFatbinResult result)
-
nvFatbinGetErrorString 为每个错误代码返回错误描述字符串。
- 参数
-
result – [in] 错误代码
- 返回值
-
nullptr，如果 result 是 NVFATBIN_SUCCESS
字符串，如果 result 不是 NVFATBIN_SUCCESS
3.2. Fatbinary 创建
函数
- 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 将 CUDA 二进制文件添加到 fatbinary。
用户负责确保所有字符串格式正确。
- 参数
-
handle – [in] nvFatbin 句柄。
code – [in] cubin。
size – [in] cubin 的大小。
arch – [in] 此 cubin 适用的架构。
identifier – [in] cubin 的名称，在使用 cuobjdump 等工具提取 fatbin 时很有用。
- 返回值
-
nvFatbinResult nvFatbinAddIndex(nvFatbinHandle handle, const void *code, size_t size, const char *identifier)
-
nvFatbinAddIndex 将索引文件添加到 fatbinary。
用户负责确保所有字符串格式正确。
- 参数
-
handle – [in] nvFatbin 句柄。
code – [in] 索引。
size – [in] 索引的大小。
identifier – [in] 索引的名称，在使用 cuobjdump 等工具提取 fatbin 时很有用。
- 返回值
-
nvFatbinResult nvFatbinAddLTOIR(nvFatbinHandle handle, const void *code, size_t size, const char *arch, const char *identifier, const char *optionsCmdLine)
-
nvFatbinAddLTOIR 将 LTOIR 添加到 fatbinary。
用户负责确保所有字符串格式正确。
- 参数
-
handle – [in] nvFatbin 句柄。
code – [in] LTOIR 代码。
size – [in] LTOIR 代码的大小。
arch – [in] 此 LTOIR 适用的架构。
identifier – [in] LTOIR 的名称，在使用 cuobjdump 等工具提取 fatbin 时很有用。
optionsCmdLine – [in] JIT 编译期间使用的选项。
- 返回值
-
nvFatbinResult nvFatbinAddPTX(nvFatbinHandle handle, const char *code, size_t size, const char *arch, const char *identifier, const char *optionsCmdLine)
-
nvFatbinAddPTX 将 PTX 添加到 fatbinary。
用户负责确保所有字符串格式正确。大小应包含终止空字符 ('\0')。如果最后一个字符不是 '\0'，则会自动添加一个，但这样做会复制代码（如果尚未复制）。
- 参数
-
handle – [in] nvFatbin 句柄。
code – [in] PTX 代码。
size – [in] PTX 代码的大小。
arch – [in] 此 PTX 适用的架构。
identifier – [in] PTX 的名称，在使用 cuobjdump 等工具提取 fatbin 时很有用。
optionsCmdLine – [in] JIT 编译期间使用的选项。
- 返回值
-
-
nvFatbinResult nvFatbinAddReloc(nvFatbinHandle handle, const void *code, size_t size)
-
nvFatbinAddReloc 将来自主机对象的可重定位 PTX 条目添加到 fatbinary。
请注意，每个可重定位的 ptx 源都必须具有唯一的标识符（标识符取自对象的条目）。强制执行此操作是因为每个 sm 的每个唯一标识符只有一个条目。另请注意，此操作会忽略句柄选项。相反，主机对象的选项从其每个条目中复制过来。
- 参数
-
handle – [in] nvFatbin 句柄。
code – [in] 主机对象镜像。
size – [in] 主机对象镜像代码的大小。
- 返回值
-
-
nvFatbinResult nvFatbinCreate(nvFatbinHandle *handle_indirect, const char **options, size_t optionsCount)
-
nvFatbinCreate 创建新句柄
- 参数
-
handle_indirect – [out] nvFatbin 句柄的地址
options – [in] 字符串数组，每个字符串包含一个选项。
optionsCount – [in] 选项数量。
- 返回值
-
nvFatbinResult nvFatbinDestroy(nvFatbinHandle *handle_indirect)
-
nvFatbinDestroy 销毁句柄。
在此调用之后使用任何其他指向句柄的指针将导致未定义的行为。传入的句柄将设置为 nullptr。
- 参数
-
handle_indirect – [in] 指向句柄的指针。
- 返回值
-
nvFatbinResult nvFatbinGet(nvFatbinHandle handle, void *buffer)
-
nvFatbinGet 返回已完成的 fatbinary。
用户负责确保缓冲区的大小适合
fatbinary。您必须在使用此函数之前调用 nvFatbinSize，否则它将返回错误。另请参阅
- 参数
-
handle – [in] nvFatbin 句柄。
buffer – [out] 用于存储 fatbinary 的内存。
- 返回值
-
nvFatbinResult nvFatbinSize(nvFatbinHandle handle, size_t *size)
-
nvFatbinSize 返回 fatbinary 的大小。
- 参数
-
handle – [in] nvFatbin 句柄。
size – [out] fatbinary 的大小
- 返回值
-
nvFatbinResult nvFatbinVersion(unsigned int *major, unsigned int *minor)
-
nvFatbinVersion 返回 nvFatbin 的当前版本
- 参数
-
major – [out] 主版本号。
minor – [out] 次版本号。
- 返回值
3.2.2. 类型定义
-
typedef struct _nvFatbinHandle *nvFatbinHandle
-
nvFatbinHandle 是 fatbin 创建的单元，以及程序的不透明句柄。
要创建 fatbin，必须首先使用 nvFatbinCreate() 创建 nvFatbinHandle 的实例。
3.3. 支持的选项
nvFatbin 支持以下选项。
选项名称以单破折号 (-) 为前缀。 接受值的选项具有赋值运算符 (=)，后跟选项值，不带空格，例如 "-host=windows"。
支持的选项包括
-32生成 32 位条目。-64生成 64 位条目。-c无效。（已弃用，将在下一个主要版本中删除。从一开始就没有任何作用。）-compress=<bool>启用 (true) / 禁用 (false) 压缩（默认值：true）。-compress-all压缩 fatbin 中的所有内容，即使它很小。-compress-mode=<mode>选择压缩行为。 有效选项包括 “default”、“size”（较小的 fatbin 大小）、“speed”（更快的解压缩速度，目前与 default 相同）、“balance”（介于 size 和 speed 之间）和 “none”（不压缩，类似于 -compress=false，但优先级高于 -compress=true）。-cuda指定 CUDA（而不是 OpenCL）。-g生成调试信息。-host=<name>指定主机操作系统。 有效选项包括 “linux”、“windows” 和 “mac”（已弃用）。-opencl指定 OpenCL（而不是 CUDA）。
4. 基本用法
本文档的这一部分使用一个简单的示例来解释如何使用 Fatbin Creator API 来链接程序。 为了简洁和可读性，未显示 API 返回值的错误检查。
此示例假设我们要创建一个 fatbin，其中包含用于 sm_52 的 CUBIN、用于 sm_61 的 PTX 和用于 sm_70 的 LTOIR。 我们可以创建一个 fatbin creator 实例并获取其 API 句柄，如图 1所示。
图 1. Fatbin Creator 的创建和程序的初始化
nvFatbinHandle handle;
nvFatbinCreate(&handle, nullptr, 0);
假设我们已经有三个存储在 std::vector 中的输入（CUBIN、PTX 和 LTOIR），它们可能来自使用 nvrtc 创建并存储到向量中的代码。 （它们不必在向量中，这仅仅说明需要数据本身及其大小。）我们可以添加输入，如图 2所示。
图 2. Fatbin creator 的输入
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所示。 然后可以通过将其传递给 CUDA Driver API，在 GPU 上执行此 fatbin。
图 4. 查询创建的 fatbin
void* fatbin = malloc(fatbinSize);
nvFatbinGet(handle, fatbin);
当不再需要 fatbin creator 时，可以销毁它，如图 5所示。
图 5. 销毁 fatbin creator
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 或驱动程序，此示例也需要兼容的 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 Toolkit 安装目录，请按如下方式构建此示例
-
使用 nvFatbin 共享库（请注意，如果测试不使用 nvrtc 或运行代码，则它不需要与 nvrtc 或 CUDA Driver 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 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 已被告知可能发生此类损害。 尽管客户可能因任何原因而遭受任何损害，但 NVIDIA 对本文所述产品的客户的累计总责任应根据产品的销售条款进行限制。
6.3.2. OpenCL
OpenCL 是 Apple Inc. 的商标，经 Khronos Group Inc. 许可使用。
6.3.3. 商标
NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。 其他公司和产品名称可能是与其相关的各自公司的商标。
© 2023 NVIDIA Corporation 和/或其附属公司。 保留所有权利。