使用GPUDirect RDMA开发Linux内核模块

启用GPUDirect RDMA连接NVIDIA GPU的API参考指南。

1. 概述

GPUDirect RDMA是一项在Kepler架构GPU和CUDA 5.0中引入的技术，它利用PCI Express的标准特性，为GPU与第三方对等设备之间的数据交换提供直接路径。第三方设备的示例包括：网络接口、视频采集设备、存储适配器。

GPUDirect RDMA 在 Tesla 和 Quadro GPU 上均可使用。

存在一些限制条件，其中最重要的是两个设备必须共享相同的上游PCI Express根复合体。部分限制取决于所使用的平台，可能会在当前或未来的产品中得到解除。

为了在广泛的硬件设备上启用此功能，需要对设备驱动程序进行一些简单的修改。本文档介绍了该技术，并描述了在Linux上启用与NVIDIA GPU的GPUDirect RDMA连接所需的步骤。

GPUDirect RDMA within the Linux Device Driver Model — Linux设备驱动模型中的GPUDirect RDMA

1.1. GPUDirect RDMA工作原理

在两台设备之间建立GPUDirect RDMA通信时，从PCI Express设备的角度来看，所有物理地址都是相同的。在这个物理地址空间中存在称为PCI BAR的线性窗口。每个设备最多有六个BAR寄存器，因此最多可以有六个活动的32位BAR区域。64位BAR会占用两个BAR寄存器。PCI Express设备对等端设备的BAR地址发起读写操作的方式与对系统内存发起操作的方式相同。

传统上，像BAR窗口这样的资源通过CPU的内存管理单元(MMU)映射为用户或内核地址空间，作为内存映射I/O(MMIO)地址。然而，由于当前操作系统缺乏足够的机制在驱动程序之间交换MMIO区域，NVIDIA内核驱动程序导出函数来执行必要的地址转换和映射。

要为设备驱动程序添加GPUDirect RDMA支持，必须修改内核驱动程序中少量的地址映射代码。这部分代码通常位于现有调用get_user_pages()的附近。

涉及GPUDirect RDMA的API和控制流程与标准DMA传输所使用的非常相似。

更多硬件详情请参阅支持的系统和PCI BAR大小。

1.2. 标准DMA传输

首先，我们概述一个从用户空间发起的标准DMA传输。在此场景中，包含以下组件：

用户空间程序
用户空间通信库
对需要进行DMA传输的设备的内核驱动

一般流程如下：

用户空间程序通过用户空间通信库请求数据传输。该操作接收一个指向数据的指针（虚拟地址）和以字节为单位的大小。
通信库必须确保虚拟地址和大小对应的内存区域已准备好进行传输。如果尚未准备好，则必须由内核驱动程序处理（下一步）。
内核驱动程序从用户空间通信库接收虚拟地址和大小。然后它请求内核将虚拟地址范围转换为物理页列表，并确保这些页面已准备好进行传输。我们将此操作称为内存锁定。
内核驱动程序使用页面列表来配置物理设备的DMA引擎。
通信库启动传输。
传输完成后，通信库最终应清理用于固定内存的所有资源。我们将此操作称为解除内存固定。

1.3. GPUDirect RDMA传输

为了支持GPUDirect RDMA传输的通信，必须对上述流程进行一些修改。首先，引入了两个新组件：

用户空间CUDA库
NVIDIA内核驱动程序

如UVA CUDA内存管理基础所述，使用CUDA库的程序将其地址空间划分为GPU和CPU虚拟地址，通信库需要为它们实现两条独立的路径。

用户空间的CUDA库提供了一个函数，使通信库能够区分CPU和GPU地址。此外，对于GPU地址，它还会返回识别该地址所代表的GPU内存所需的额外元数据。详情请参阅用户空间API。

CPU和GPU地址路径的区别在于内存的固定(pinning)和释放(unpinning)方式。对于CPU内存，这是由Linux内核内置函数(get_user_pages()和put_page())处理的。然而在GPU内存情况下，固定和释放操作必须由NVIDIA内核驱动程序提供的函数来处理。详情请参阅Pinning GPU memory和Unpinning GPU memory。

部分硬件注意事项在支持的系统和PCI BAR大小中有说明。

1.4. CUDA 6.0版本变更

本节我们简要列出CUDA 6.0中的可用变更：

CUDA点对点令牌不再是强制性的。对于调用进程拥有的内存缓冲区（这是典型情况），在内核模式函数nvidia_p2p_get_pages()中可以用零(0)替代令牌。这一新特性旨在让现有的第三方软件栈更容易采用GPUDirect的RDMA技术。
由于上述变更，引入了一个新的API cuPointerSetAttribute()。该API必须用于注册任何未使用点对点令牌的缓冲区。为确保CUDA API在操作可能被RDMA读取用于GPUDirect的内存时具有正确的同步行为，必须使用此API。在这些情况下未使用它可能导致数据损坏。请参阅令牌使用中的变更。
cuPointerGetAttribute() 已扩展为返回一个全局唯一的数字标识符，底层库可以利用该标识符来检测用户级代码中发生的缓冲区重新分配（参见Userspace API）。当无法拦截CUDA分配和释放API时，它提供了一种替代方法来检测重新分配。
内核模式内存固定功能已扩展至可与多进程服务(MPS)协同工作。

截至CUDA 6.0的注意事项：

CUDA统一内存与GPUDirect RDMA结合使用时未得到明确支持。虽然nvidia_p2p_get_pages()返回的页表对托管内存缓冲区有效，并提供了GPU内存在任何给定时刻的映射，但该内存的GPU设备副本可能与不在GPU上的可写页副本不一致。在这种情况下使用页表可能导致访问过时数据或数据丢失，因为对设备内存的DMA写入访问随后会被统一内存运行时覆盖。cuPointerGetAttribute()可用于确定地址是否由统一内存运行时管理。
每次固定设备内存区域时，都会无条件分配新的GPU BAR空间，即使固定的是重叠或重复的设备内存范围，也就是说，系统不会尝试复用映射。这种行为自CUDA 7.0起已更改。

1.5. CUDA 7.0的变更

本节我们简要列出CUDA 7.0中的可用变更：

在IBM POWER8平台上，GPUDirect RDMA不受支持，尽管它没有被明确禁用。
不能保证GPUDirect RDMA在任何给定的ARM64平台上都能正常工作。
针对CUDA 6.0改进了GPU BAR映射的管理。现在当设备内存区域被固定时，GPU BAR空间可能会与预先存在的映射共享。例如在固定重叠或重复的设备内存范围时就会出现这种情况。因此，当取消固定某个区域时，即使只有部分BAR空间被共享，其整个BAR空间也不会被释放。
引入了新的cuPointerGetAttributes() API。当需要检索同一缓冲区的多个属性时（例如在MPI中检查新缓冲区时），该API会非常有用。
cudaPointerGetAttributes() 现在速度更快，因为它在内部利用了 cuPointerGetAttributes()。
在CUDA 6.5中新增了一个示例代码samples/7_CUDALibraries/cuHook，可作为实现CUDA内存分配/释放API拦截框架的模板。

1.6. CUDA 8.0的变更

本节我们简要列出CUDA 8.0中的可用变更：

nvidia_p2p_page_table结构体已扩展包含一个新成员，同时保持二进制兼容性。NVIDIA_P2P_PAGE_TABLE_VERSION宏中的次要版本号已相应更新。
引入了nvidia_p2p_dma_mapping结构体、nvidia_p2p_dma_map_pages()和nvidia_p2p_dma_unmap_pages() API接口，以及NVIDIA_P2P_DMA_MAPPING_VERSION宏。这些API可供第三方设备驱动程序使用，将GPU BAR页面映射或解除映射到其设备的I/O地址空间中。主要应用场景是在某些平台上，用于PCIe点对点传输的PCIe资源I/O地址与CPU访问这些资源时使用的物理地址不同。查看此链接获取使用这些新API的代码示例。
引入了NVIDIA_P2P_PAGE_TABLE_VERSION_COMPATIBLE和NVIDIA_P2P_DMA_MAPPING_VERSION_COMPATIBLE宏。这些宏旨在供第三方设备驱动程序调用，以检查运行时二进制兼容性，例如在数据结构布局发生变化的情况下。
在IBM POWER8平台上，当使用上述API时，据报告GPUDirect RDMA仅在GPU与第三方设备通过受支持的PCIe交换机连接的情况下能正常工作。

1.7. CUDA 10.1的变更

GPUDirect RDMA 在 Jetson AGX Xavier 平台上得到支持。详情请参阅移植到Tegra部分。

1.8. CUDA 11.2的变更

GPUDirect RDMA在基于Drive AGX Xavier Linux的平台上得到支持。详情请参阅移植到Tegra部分。

1.9. CUDA 11.4版本变更

新增了一个内核模块nvidia-peermem，该模块支持基于Mellanox InfiniBand的HCA（主机通道适配器）直接对NVIDIA GPU显存进行点对点读写访问。详情请参阅Using nvidia-peermem。

GPUDirect RDMA在Jetson Orin平台上获得支持。详情请参阅移植到Tegra章节。

已知问题：

目前没有自动加载nvidia-peermem的服务，用户需要手动加载该模块。

1.10. CUDA 12.2版本变更

在从R515到R535分支发布的驱动程序中（下文提到的新版R525和R535版本除外），存在一个竞态错误，可能表现为内核空指针解引用。当GPU调用（此处为I/O）内核驱动程序失效回调（该回调是在调用nvidia_p2p_get_pages时注册的）的同时，I/O驱动程序调用nvidia_p2p_put_pages时，就会发生这种情况。该竞态错误不会影响持久映射的情况，因为在持久映射情况下既不支持也不需要失效回调。

该错误修复需要以下API变更：

nvidia_p2p_get_pages 不再接受NULL回调指针。
相反，当请求持久映射时，应使用nvidia_p2p_put_pages_persistent和nvidia_p2p_get_pages_persistent替代。
这些新的持久API的使用可以通过预处理宏NVIDIA_P2P_CAP_GET_PAGES_PERSISTENT_API来保护，例如在编写可移植驱动程序时。
nvidia-peermem 内核模块已相应更新。
虽然在运行R470分支及更新版本的GPU驱动程序时已弃用，但仍在使用树外nv_peer_mem 模块（https://github.com/Mellanox/nv_peer_memory）且需要持久映射功能的客户将需要切换到nvidia-peermem。

请注意，不需要持久映射的I/O驱动程序无需修改源代码。

上述API变更已部署在R535分支中，具体为535.14及后续版本，同时也已向后移植到R525分支，适用于TeslaRD3（525.105.17）及更高版本。

2. 设计考量

在设计利用GPUDirect RDMA的系统时，需要考虑多个因素。

2.1. 延迟解固定优化

将GPU设备内存固定在BAR中是一项昂贵的操作，耗时可达毫秒级。因此，应用程序的设计应尽量减少这种开销。

使用GPUDirect RDMA最直接的实现方式是在每次传输前固定内存，并在传输完成后立即解除固定。然而，这种做法通常性能较差，因为固定和解除固定内存都是开销较大的操作。不过，执行RDMA传输所需的其余步骤可以快速完成而无需进入内核（DMA列表可以通过MMIO寄存器/命令列表进行缓存和重放）。

因此，延迟解除内存固定是实现高性能RDMA的关键。这意味着即使在传输完成后仍保持内存固定状态。这种做法利用了同一内存区域很可能被用于未来DMA传输的特性，从而通过延迟解除固定来节省反复固定/解除固定的操作开销。

惰性解除固定(pinning)的一个示例实现是维护一组固定的内存区域，仅当这些区域的总大小达到某个阈值时，或者由于BAR空间耗尽导致固定新区域失败时（参考PCI BAR大小），才会解除其中部分区域（例如最近最少使用的区域）的固定状态。

2.2. 注册缓存

通信中间件通常采用一种称为注册缓存（registration cache）或固定缓存（pin-down cache）的优化技术，以最小化固定开销。这类缓存通常已针对主机内存实现，包含延迟解固定、LRU注销等功能。对于网络中间件，此类缓存通常实现在用户空间，因为它们需要与支持用户态消息注入的硬件配合使用。CUDA统一虚拟寻址（UVA）内存布局通过考虑少量设计因素，使得GPU内存固定能够与这些缓存协同工作。在CUDA环境中，这一点尤为重要，因为可固定的内存量可能比主机内存受到更严格的限制。

由于GPU BAR空间通常使用64KB页面进行映射，将区域缓存对齐到64KB边界能更高效地利用资源。尤其当两个内存区域位于同一64KB边界时，它们会分配并返回相同的BAR映射。

注册缓存通常依赖于能够拦截用户应用程序中发生的释放事件，以便它们可以取消固定内存并释放重要的硬件资源，例如在网卡上。要为GPU内存实现类似的机制，实现方案有两种选择：

对所有CUDA分配和释放API进行插桩。
使用标签检查功能来跟踪释放和重新分配。参考Buffer ID Tag Check for A Registration Cache。

这里有一个示例应用7_CUDALibraries/cuHook，展示了如何在运行时拦截对CUDA API的调用，该功能可用于检测GPU内存的分配/释放操作。

虽然拦截CUDA API超出了本文档的范围，但从CUDA 6.0开始提供了一种执行标签检查的方法。该方法涉及在cuPointerGetAttribute()（或如果需要更多属性则使用cuPointerGetAttributes()）中使用CU_POINTER_ATTRIBUTE_BUFFER_ID属性来检测内存缓冲区的释放或重新分配。如果缓冲区地址不再有效，API将在重新分配时返回不同的ID值或返回错误。有关API用法，请参阅Userspace API。

注意

使用标签检查会在每次使用内存缓冲区时向CUDA API引入额外调用，因此当额外延迟不是问题时，这种方法最为合适。

2.3. 取消固定回调

当第三方设备驱动程序使用nvidia_p2p_get_pages()固定GPU页面时，它还必须提供一个回调函数，如果NVIDIA驱动程序需要撤销对映射的访问，它将调用该函数。此回调是同步发生的，这为第三方驱动程序提供了清理和删除对相关页面的任何引用（即等待未完成的DMA完成）的机会。用户回调函数可能会阻塞几毫秒，但建议回调尽可能快地完成。必须小心不要引入死锁，因为在回调中等待GPU执行任何操作是不安全的。

回调函数必须调用nvidia_p2p_free_page_table()（而非nvidia_p2p_put_pages()）来释放page_table所指向的内存。对应的映射内存区域只有在从回调函数返回后才会被NVIDIA驱动解除映射。

请注意，回调函数将在以下两种情况下被调用：

如果用户空间程序在第三方内核驱动有机会使用nvidia_p2p_put_pages()解除内存固定之前，就显式释放了对应的GPU内存，例如通过cuMemFree、cuCtxDestroy等操作。
由于进程提前退出导致的结果。

在后一种情况下，第三方内核驱动程序和NVIDIA内核驱动程序的关闭文件描述符之间可能存在拆卸顺序问题。如果NVIDIA内核驱动程序的文件描述符首先关闭，则会调用nvidia_p2p_put_pages()回调函数。

良好的软件设计至关重要，因为NVIDIA内核驱动程序在调用回调之前会通过锁机制保护自身免受重入问题的影响。第三方内核驱动程序几乎肯定会采取类似的措施，因此如果不仔细考虑，可能会出现死锁或活锁的情况。

2.4. 支持的系统

总体说明

尽管第三方设备与NVIDIA GPU之间实现GPUDirect RDMA的唯一理论要求是它们共享相同的根复合体，但存在一些错误（主要出现在芯片组中）会导致其性能不佳，或在某些配置中完全无法工作。

我们可以根据GPU与第三方设备之间路径上的情况，区分以下三种情形：

仅限PCIe交换机
单CPU/IOH
CPU/IOH <-> QPI/HT <-> CPU/IOH

第一种情况，路径上仅有PCIe交换器，是最优配置，能提供最佳性能。第二种情况涉及单个CPU/IOH，虽然可用但性能较差（特别是点对点读取带宽在某些处理器架构上会受到严重限制）。最后第三种情况，当路径需要穿越QPI/HT链路时，性能可能受到极大限制甚至无法稳定运行。

提示

lspci 可用于检查 PCI 拓扑结构：

$ lspci -t

平台支持

对于IBM POWER8平台，不支持GPUDirect RDMA和P2P功能，但并未明确禁用。在运行时这些功能可能无法正常工作。

GPUDirect RDMA功能在Jetson AGX Xavier平台上从CUDA 10.1开始支持，在基于Linux的Drive AGX Xavier平台上从CUDA 11.2开始支持。详情请参阅移植到Tegra。在ARM64架构上，必要的点对点功能取决于特定平台的硬件和软件。因此，虽然GPUDirect RDMA在非Jetson和非Drive平台上没有被明确禁用，但不能保证其完全正常运行。

IOMMU（输入输出内存管理单元）

GPUDirect RDMA目前依赖于从不同PCI设备视角看所有物理地址都相同这一条件。这意味着它与执行非1:1地址转换的IOMMU不兼容，因此必须禁用IOMMU或将其配置为直通转换模式才能使GPUDirect RDMA正常工作。

2.5. PCI BAR 大小

PCI设备可以请求操作系统/BIOS将一段物理地址空间映射给它们。这些区域通常被称为BAR。NVIDIA GPU目前会暴露多个BAR，其中一些可以支持任意设备内存，这使得GPUDirect RDMA成为可能。可用于GPUDirect RDMA的最大BAR大小因GPU而异。例如，目前Kepler级GPU上可用的最小BAR大小为256MB。其中，32MB目前保留供内部使用。这些大小可能会发生变化。

在某些Tesla级GPU上启用了大型BAR功能，例如BAR1大小设置为16GB或更大。大容量BAR可能会对BIOS造成问题，尤其是在较旧的主板上，这与对32位操作系统的兼容性支持有关。在这些主板上，引导程序可能会在早期POST阶段停止，或者GPU可能配置错误而无法使用。如果出现这种情况，可能需要启用某些特殊的BIOS功能来处理大容量BAR问题。有关大容量BAR支持的更多详情，请咨询您的系统供应商。

2.6. 令牌使用情况

警告

从CUDA 6.0开始，令牌(tokens)应被视为已弃用，尽管它们仍受支持。

如用户空间API和内核API所示，固定和释放内存的一种方法除了需要GPU虚拟地址外，还需要两个令牌。

这些令牌p2pToken和vaSpaceToken是唯一标识GPU虚拟地址空间所必需的。仅靠进程标识符无法识别GPU虚拟地址空间。

令牌在单个CUDA上下文中保持一致（即，在同一CUDA上下文中通过cudaMalloc()获取的所有内存将具有相同的p2pToken和vaSpaceToken）。然而，给定的GPU虚拟地址在其整个生命周期内不一定映射到相同的上下文/GPU。具体示例如下：

cudaSetDevice(0)
ptr0 = cudaMalloc();
cuPointerGetAttribute(&return_data, CU_POINTER_ATTRIBUTE_P2P_TOKENS, ptr0);
// Returns [p2pToken = 0xabcd, vaSpaceToken = 0x1]
cudaFree(ptr0);
cudaSetDevice(1);
ptr1 = cudaMalloc();
assert(ptr0 == ptr1);
// The CUDA driver is free (although not guaranteed) to reuse the VA,
// even on a different GPU
cuPointerGetAttribute(&return_data, CU_POINTER_ATTRIBUTE_P2P_TOKENS, ptr0);
// Returns [p2pToken = 0x0123, vaSpaceToken = 0x2]

也就是说，同一个地址在程序执行的不同时刻传递给cuPointerGetAttribute时，可能会返回不同的令牌。因此，第三方通信库必须对其操作的每个指针调用cuPointerGetAttribute()。

安全影响

这两个令牌作为NVIDIA内核驱动程序的认证机制。如果您知道这些令牌，就可以映射它们对应的地址空间，而NVIDIA内核驱动程序不会执行任何额外检查。64位的p2pToken经过随机化处理，以防止被攻击者猜测到。

当未使用令牌时，NVIDIA驱动程序将Kernel API限制为拥有内存分配的进程。

2.7. 同步与内存排序

GPUDirect RDMA引入了一条新的独立GPU数据流路径，向第三方设备开放，理解这些设备如何与GPU的宽松内存模型交互非常重要。

为了确保映射与CUDA API对该内存的操作保持一致，必须正确注册CUDA内存的BAR映射。
只有CUDA同步和工作提交API提供GPUDirect RDMA操作的内存排序功能。

CUDA API一致性注册

为确保CUDA API内存操作在API调用将控制权返回给调用CPU线程之前对BAR映射可见，注册是必要的。这为使用GPUDirect RDMA映射的设备提供了内存的一致性视图，当在线程中的CUDA API之后调用时。这是CUDA API的一种更为保守的操作模式，会禁用优化，因此可能会对性能产生负面影响。

此行为按每次分配的粒度启用，可以通过调用cuPointerSetAttribute()并设置CU_POINTER_ATTRIBUTE_SYNC_MEMOPS属性来实现，或者在使用旧版路径时为缓冲区获取p2p令牌。更多详情请参阅Userspace API。

一个示例情况是cuMemcpyDtoD()与后续对复制目标进行的GPUDirect RDMA读取操作之间存在写后读依赖关系。作为优化手段，设备到设备的内存复制通常在将复制任务加入GPU调度队列后异步返回调用线程。但在这种情况下，会导致通过BAR映射读取的数据不一致，因此该优化会被禁用，确保在CUDA API返回前完成复制操作。

CUDA 内存排序 API

只有由CPU发起的CUDA API才能确保GPU观察到的GPUDirect内存操作顺序。也就是说，尽管第三方设备已发出所有PCIE事务，但在后续CPU发起的CUDA工作提交或同步API之前，正在运行的GPU内核或复制操作可能会观察到过时数据或乱序到达的数据。为确保内存更新对CUDA内核或复制操作可见，实现方案应保证所有对GPU BAR的写入操作在控制权返回给将调用相关CUDA API的CPU线程之前完成。

网络通信场景中的一个示例情况是，当第三方网络设备完成网络RDMA写操作并将数据写入GPU BAR映射时。尽管通过GPU BAR或CUDA内存复制操作读取回写入的数据将返回新写入的数据，但与该网络写入同时运行的GPU内核可能会观察到过时数据、部分写入的数据或乱序写入的数据。

简而言之，GPU内核与支持GPUDirect操作的并发RDMA完全不兼容，在这种情况下访问被第三方设备覆盖的内存将被视为数据竞争。要解决这种不一致性并消除数据竞争，DMA写操作必须相对于将启动依赖GPU内核的CPU线程完成。

3. 如何执行特定任务

3.1. 显示GPU BAR空间

从CUDA 6.0开始，NVIDIA SMI工具提供了转储BAR1内存使用情况的功能。该功能可用于了解应用程序对BAR空间的使用情况，这是GPUDirect RDMA映射消耗的主要资源。

$ nvidia-smi -q
...
      BAR1 Memory Usage
         Total                       : 256 MiB
         Used                        : 2 MiB
         Free                        : 254 MiB
...

GPU内存以固定大小的块进行锁定，因此此处显示的空间量可能出乎意料。此外，驱动程序会保留一定数量的BAR空间供内部使用，因此并非所有可用内存都能通过GPUDirect RDMA访问。请注意，相同的功能也可以通过nvmlDeviceGetBAR1MemoryInfo() NVML API以编程方式实现。

3.2. 固定GPU内存

要实现正确的行为，需要在内存地址上使用cuPointerSetAttribute()来启用CUDA驱动中适当的同步行为。详情请参阅同步与内存排序。

void pin_buffer(void *address, size_t size)
{
    unsigned int flag = 1;
    CUresult status = cuPointerSetAttribute(&flag, CU_POINTER_ATTRIBUTE_SYNC_MEMOPS, address);
    if (CUDA_SUCCESS == status) {
        // GPU路径
        pass_to_kernel_driver(address, size);
    } else {
        // CPU路径
        // ...
    }
}

这样做是为了让CUDA驱动以特殊方式处理GPU内存缓冲区，从而确保CUDA内存传输始终与主机保持同步。有关cuPointerSetAttribute()的详细信息，请参阅用户空间API。

In the kernel driver, invoke nvidia_p2p_get_pages().

// for boundary alignment requirement
#define GPU_BOUND_SHIFT   16
#define GPU_BOUND_SIZE    ((u64)1 << GPU_BOUND_SHIFT)
#define GPU_BOUND_OFFSET  (GPU_BOUND_SIZE-1)
#define GPU_BOUND_MASK    (~GPU_BOUND_OFFSET)

struct kmd_state {
    nvidia_p2p_page_table_t *page_table;
    // ...
};

void kmd_pin_memory(struct kmd_state *my_state, void *address, size_t size)
{
    // do proper alignment, as required by NVIDIA kernel driver
    u64 virt_start = address & GPU_BOUND_MASK;
    size_t pin_size = (address + size - virt_start + GPU_BOUND_SIZE - 1) & GPU_BOUND_MASK;
    if (!size)
        return -EINVAL;
    int ret = nvidia_p2p_get_pages(0, 0, virt_start, pin_size, &my_state->page_table, free_callback, &my_state);
    if (ret == 0) {
        // Succesfully pinned, page_table can be accessed
    } else {
        // Pinning failed
    }
}

Note how the start address is aligned to a 64KB boundary before calling the pinning functions.

If the function succeeds the memory has been pinned and the page_table entries can be used to program the device’s DMA engine. Refer to Kernel API for details on nvidia_p2p_get_pages().

3.3. 释放GPU内存

在内核驱动程序中，调用 nvidia_p2p_put_pages()。

void unpin_memory(void *address, size_t size, nvidia_p2p_page_table_t *page_table)
{
    nvidia_p2p_put_pages(0, 0, address, size, page_table);
}

有关nvidia_p2p_put_pages()的详细信息，请参阅Kernel API。

从CUDA 6.0开始，应使用零作为令牌参数。请注意，nvidia_p2p_put_pages()必须从与发出相应nvidia_p2p_get_pages()相同的进程上下文中调用。

3.4. 处理空闲回调

如果NVIDIA内核驱动需要撤销映射，则会按照nvidia_p2p_get_pages()调用中的指定调用free_callback(data)。详情请参阅Kernel API和Unpin Callback。

回调函数会等待挂起的传输完成，然后清理页表分配。

void free_callback(void *data)
{
   my_state *state = data;
   wait_for_pending_transfers(state);
   nvidia_p2p_free_pages(state->page_table);
}

NVIDIA内核驱动程序会处理取消映射，因此不应调用nvidia_p2p_put_pages()。

3.5. 注册缓存中的缓冲区ID标签检查

请注意，对于延迟敏感的实现，不建议采用基于缓冲区ID标签检查的解决方案。相反，建议通过检测CUDA分配和释放API来为注册缓存提供回调，从而消除关键路径上的标签检查开销。

当首次遇到设备内存缓冲区并识别为尚未固定时，会创建固定映射并检索关联的缓冲区ID，然后将其一起存储在缓存条目中。cuMemGetAddressRange()函数可用于获取整个分配的起始地址和大小，进而用于固定操作。由于nvidia_p2p_get_pages()需要64K对齐的指针，直接对齐缓存地址很有必要。此外，由于BAR空间当前以64KB为单位进行映射，将整个固定操作按64KB对齐能更高效地利用资源。
```
// struct buf表示注册缓存的一个条目
struct buf {
    CUdeviceptr pointer;
    size_t      size;
    CUdeviceptr aligned_pointer;
    size_t      aligned_size;
    int         is_pinned;
    uint64_t    id; // 固定操作后立即获取的缓冲区ID
};
```

Once created, every time a registration cache entry will be used it must be first checked for validity. One way to do this is to use the Buffer ID provided by CUDA as a tag to check for deallocation or reallocation.

int buf_is_gpu_pinning_valid(struct buf* buf) {
    uint64_t buffer_id;
    int retcode;
    assert(buf->is_pinned);
    // get the current buffer id
    retcode = cuPointerGetAttribute(&buffer_id, CU_POINTER_ATTRIBUTE_BUFFER_ID, buf->pointer);
    if (CUDA_ERROR_INVALID_VALUE == retcode) {
        // the device pointer is no longer valid
        // it could have been deallocated
        return ERROR_INVALIDATED;
    } else if (CUDA_SUCCESS != retcode) {
        // handle more serious errors here
        return ERROR_SERIOUS;
    }
    if (buf->id != buffer_id)
        // the original buffer has been deallocated and the cached mapping should be invalidated and the buffer re-pinned
        return ERROR_INVALIDATED;
    return 0;
}

When the buffer identifier changes the corresponding memory buffer has been reallocated so the corresponding kernel-space page table will not be valid anymore. In this case the kernel-space nvidia_p2p_get_pages() callback would have been invoked. Thus the Buffer IDs provide a tag to keep the pin-down cache consistent with the kernel-space page table without requiring the kernel driver to up-call into the user-space.

If CUDA_ERROR_INVALID_VALUE is returned from cuPointerGetAttribute(), the program should assume that the memory buffer has been deallocated or is otherwise not a valid GPU memory buffer.

在这两种情况下，都必须使相应的缓存条目失效。

// 在注册缓存代码中
if (buf->is_pinned && !buf_is_gpu_pinning_valid(buf)) {
  regcache_invalidate_entry(buf);
  pin_buffer(buf);
}

3.6. 将内核模块链接到nvidia.ko

运行提取脚本：
```
./NVIDIA-Linux-x86_64-.run -x
```
这将提取NVIDIA驱动程序和内核包装器。
进入输出目录：
```
cd /kernel/
```
在此目录下，为您的内核构建NVIDIA模块：
```
make module
```
完成后，内核构建目录下的Module.symvers文件将包含nvidia.ko的符号信息。

在内核模块构建过程中添加以下行进行修改：

KBUILD_EXTRA_SYMBOLS := <内核构建目录路径>/Module.symvers

3.7. 使用nvidia-peermem

NVIDIA GPU驱动包提供了一个内核模块nvidia-peermem，该模块支持基于NVIDIA InfiniBand的主机通道适配器(HCAs)直接对NVIDIA GPU显存进行点对点读写访问。它允许基于GPUDirect RDMA的应用程序在无需将数据复制到主机内存的情况下，通过RDMA互连技术利用GPU计算能力。

此功能支持使用NVIDIA ConnectX®-3 VPI或更新版本的适配器。它适用于InfiniBand和RoCE（基于融合以太网的RDMA）技术。

NVIDIA OFED（开放结构企业版分发），或称MLNX_OFED，在InfiniBand核心与对等内存客户端（如NVIDIA GPU）之间引入了一个API。nvidia-peermem模块通过使用NVIDIA GPU驱动程序提供的点对点API，将NVIDIA GPU注册到InfiniBand子系统中。

内核必须通过额外的补丁或通过MLNX_OFED提供对RDMA对等内存的必要支持，这是加载和使用nvidia-peermem的前提条件。

系统可能已安装并加载了来自GitHub项目的nv_peer_mem模块。安装nvidia-peermem不会影响现有nv_peer_mem模块的功能。但是，要加载并使用nvidia-peermem，用户必须禁用nv_peer_mem服务。此外，建议卸载nv_peer_mem软件包以避免与nvidia-peermem产生冲突，因为同一时间只能加载一个模块。

要停止nv_peer_mem服务：

# service nv_peer_mem stop</screen>

检查停止服务后nv_peer_mem.ko是否仍在加载：

# lsmod | grep nv_peer_mem

如果nv_peer_mem.ko仍然加载着，请使用以下命令卸载：

# rmmod nv_peer_mem

卸载 nv_peer_mem 软件包：

适用于基于DEB的操作系统：

# dpkg -P nvidia-peer-memory

# dpkg -P nvidia-peer-memory-dkms

适用于基于RPM的操作系统：

# rpm -e nvidia_peer_memory

在确保内核支持并安装GPU驱动程序后，可以在终端窗口中使用以下命令以root权限加载nvidia-peermem：

# modprobe nvidia-peermem

注意

注意：如果在安装MLNX_OFED之前已安装NVIDIA GPU驱动程序，则必须卸载并重新安装GPU驱动程序，以确保nvidia-peermem能够使用MLNX_OFED提供的RDMA API进行编译。

4. 参考文献

4.1. UVA CUDA内存管理基础

统一虚拟寻址(UVA)是CUDA 4.0及更高版本在运行64位进程的Fermi和Kepler GPU上默认启用的内存地址管理系统。UVA内存管理的设计为GPUDirect RDMA的运行提供了基础。在支持UVA的配置中，当CUDA运行时初始化时，应用程序的虚拟地址(VA)范围被划分为两个区域：CUDA管理的VA范围和操作系统管理的VA范围。所有CUDA管理的指针都在此VA范围内，且该范围始终位于进程VA空间的前40位内。

CUDA VA Space Addressing — CUDA虚拟地址空间寻址

随后，在CUDA VA空间中，地址可细分为三种类型：

GPU: 一个由GPU内存支持的页面。这将无法从主机访问，并且所讨论的虚拟地址(VA)在主机上永远不会具有物理后备。从CPU解引用指向GPU虚拟地址的指针将触发段错误。
CPU: 一个由CPU内存支持的页面。该页面可通过相同的虚拟地址同时从主机和GPU访问。
FREE: 这些虚拟地址空间由CUDA保留，供未来分配使用。

这种分区方式使得CUDA运行时能够通过内存对象在预留的CUDA虚拟地址空间中的指针值来确定其物理位置。

地址按页面粒度细分为以下几类；同一页面内的所有内存属于同一类型。请注意，GPU页面大小可能与CPU页面不同。CPU页面通常为4KB，而Kepler级GPU的页面为64KB。GPUDirect RDMA仅对此CUDA虚拟地址空间内的GPU页面（由cudaMalloc()创建）进行操作。

4.2. 用户空间API

数据结构

typedef struct CUDA_POINTER_ATTRIBUTE_P2P_TOKENS_st {
    unsigned long long p2pToken;
    unsigned int vaSpaceToken;
} CUDA_POINTER_ATTRIBUTE_P2P_TOKENS;

函数 cuPointerSetAttribute()

CUresult cuPointerSetAttribute(void *data, CUpointer_attribute attribute, CUdeviceptr pointer);

在GPUDirect RDMA范围内，一个有趣的用法是当CU_POINTER_ATTRIBUTE_SYNC_MEMOPS作为attribute参数传递时：

unsigned int flag = 1;
cuPointerSetAttribute(&flag, CU_POINTER_ATTRIBUTE_SYNC_MEMOPS, pointer);

参数

data [in]: 指向一个包含布尔值的unsigned int类型变量的指针。
attribute [in]: 在GPUDirect RDMA范围内应始终使用CU_POINTER_ATTRIBUTE_SYNC_MEMOPS。
pointer [in]: 一个指针。

返回

CUDA_SUCCESS: 如果指针指向GPU内存且CUDA驱动程序能够为整个设备内存分配设置新的行为。
anything else: 如果指针指向CPU内存。

该功能用于显式启用对pointer所指向的整个内存区域的严格同步行为，通过这种方式禁用所有可能导致RDMA和CUDA内存复制操作并发问题的数据传输优化。此API具有CUDA同步行为，因此应视为高开销操作，建议每个缓冲区仅调用一次。

函数 cuPointerGetAttribute()

CUresult cuPointerGetAttribute(const void *data, CUpointer_attribute attribute, CUdeviceptr pointer);

该函数有两个与GPUDirect RDMA相关的不同属性：CU_POINTER_ATTRIBUTE_P2P_TOKENS 和 CU_POINTER_ATTRIBUTE_BUFFER_ID。

警告

CU_POINTER_ATTRIBUTE_P2P_TOKENS 已在 CUDA 6.0 中被弃用

当将CU_POINTER_ATTRIBUTE_P2P_TOKENS作为attribute参数传递时，data是指向CUDA_POINTER_ATTRIBUTE_P2P_TOKENS的指针：

CUDA_POINTER_ATTRIBUTE_P2P_TOKENS tokens;
cuPointerGetAttribute(&tokens, CU_POINTER_ATTRIBUTE_P2P_TOKENS, pointer);

在这种情况下，该函数返回两个令牌用于Kernel API。

参数

data [out]: 结构体 CUDA_POINTER_ATTRIBUTE_P2P_TOKENS 包含两个令牌。
attribute [in]: 在GPUDirect RDMA范围内应始终使用CU_POINTER_ATTRIBUTE_P2P_TOKENS。
pointer [in]: 一个指针。

返回

CUDA_SUCCESS: 如果指针指向GPU内存。
anything else: 如果指针指向CPU内存。

该函数可以在任何时候调用，包括在CUDA初始化之前，并且它具有CUDA同步行为，如CU_POINTER_ATTRIBUTE_SYNC_MEMOPS所示，因此应视为开销较大的操作，每个缓冲区仅应调用一次。

请注意，在用户空间程序的生命周期内，对于相同的pointer值，tokens中设置的值可能会有所不同。具体示例请参阅Tokens Usage。

出于安全考虑，p2pToken中设置的值将被随机化，以防止被攻击者猜测到。

在CUDA 6.0中，引入了一个新属性，可用于检测内存重新分配。

当将CU_POINTER_ATTRIBUTE_BUFFER_ID作为attribute参数传入时，data应指向一个64位无符号整型变量，例如uint64_t。

uint64_t buf_id;
cuPointerGetAttribute(&buf_id, CU_POINTER_ATTRIBUTE_BUFFER_ID, pointer);

参数

data [out]: 指向一个64位变量的指针，用于存储缓冲区ID。
attribute [in]: CU_POINTER_ATTRIBUTE_BUFFER_ID 枚举器。
pointer [in]: 指向GPU内存的指针。

返回

CUDA_SUCCESS: 如果指针指向GPU内存。
anything else: 如果指针指向CPU内存。

以下是一些总体说明：

cuPointerGetAttribute() 和 cuPointerSetAttribute() 是仅属于CUDA驱动API的函数。
特别是，cuPointerGetAttribute()并不等同于cudaPointerGetAttributes()，因为所需功能仅在前者函数中提供。这完全不会限制GPUDirect RDMA的使用范围，因为cuPointerGetAttribute()与CUDA Runtime API兼容。
没有提供与cuPointerGetAttribute()等效的运行时API。这是因为与CUDA运行时API到驱动API调用序列相关的额外开销会引入不必要的负担，而cuPointerGetAttribute()可能位于关键路径上，例如通信库的关键路径。
在可能的情况下，我们建议通过使用cuPointerGetAttributes来合并对cuPointerGetAttribute的多次调用。

函数 ``cuPointerGetAttributes()``

CUresult cuPointerGetAttributes(unsigned int numAttributes, CUpointer_attribute *attributes, void **data, CUdeviceptr ptr);

该函数可用于一次性检查多个属性。最可能与GPUDirect RDMA相关的属性包括CU_POINTER_ATTRIBUTE_BUFFER_ID、CU_POINTER_ATTRIBUTE_MEMORY_TYPE和CU_POINTER_ATTRIBUTE_IS_MANAGED。

4.3. 内核API

以下声明可在NVIDIA驱动程序包中分发的nv-p2p.h头文件中找到。有关下述函数参数和返回值的详细说明，请参阅该头文件中的内联文档。

预处理器宏

NVIDIA_P2P_PAGE_TABLE_VERSION_COMPATIBLE() 和 NVIDIA_P2P_DMA_MAPPING_VERSION_COMPATIBLE() 预处理器宏旨在供第三方设备驱动程序调用，以检查运行时二进制兼容性。

结构体 nvidia_p2p_page

typedef
struct nvidia_p2p_page {
    uint64_t physical_address;
    union nvidia_p2p_request_registers {
        struct {
            uint32_t wreqmb_h;
            uint32_t rreqmb_h;
            uint32_t rreqmb_0;
            uint32_t reserved[3];
        } fermi;
    } registers;
} nvidia_p2p_page_t;

在nvidia_p2p_page结构中，只有physical_address字段与GPUDirect RDMA相关。

结构体 nvidia_p2p_page_table

typedef
struct nvidia_p2p_page_table {
    uint32_t version;
    uint32_t page_size;
    struct nvidia_p2p_page **pages;
    uint32_t entries;
    uint8_t *gpu_uuid;
} nvidia_p2p_page_table_t;

在访问页表的其他字段之前，应使用NVIDIA_P2P_PAGE_TABLE_VERSION_COMPATIBLE()检查页表的version字段。

page_size字段根据nvidia_p2p_page_size_type枚举类型进行编码。

结构体 nvidia_p2p_dma_mapping

typedef
struct nvidia_p2p_dma_mapping {
    uint32_t  version;
    enum nvidia_p2p_page_size_type page_size_type;
    uint32_t  entries;
    uint64_t *dma_addresses;
} nvidia_p2p_dma_mapping_t;

在访问其他字段之前，应将dma映射的版本字段传递给NVIDIA_P2P_DMA_MAPPING_VERSION_COMPATIBLE()。

函数 nvidia_p2p_get_pages()

int nvidia_p2p_get_pages(uint64_t p2p_token, uint32_t va_space_token,
                uint64_t virtual_address,
                uint64_t length,
                struct nvidia_p2p_page_table **page_table,
                void (*free_callback)(void *data),
                void *data);

此函数使第三方设备能够访问GPU虚拟内存范围内底层页面。

警告

这是一项开销较大的操作，应尽可能减少执行频率 - 请参阅Lazy Unpinning Optimization。

函数 nvidia_p2p_put_pages()

int nvidia_p2p_put_pages(uint64_t p2p_token, uint32_t va_space_token,
        uint64_t virtual_address,
        struct nvidia_p2p_page_table *page_table);

此函数释放一组先前对第三方设备可访问的页面。警告：不应在nvidia_p2p_get_pages()回调函数内部调用。

函数 nvidia_p2p_free_page_table()

int nvidia_p2p_free_page_table(struct nvidia_p2p_page_table *page_table);

此函数释放第三方P2P页表，旨在在nvidia_p2p_get_pages()回调执行期间调用。

函数 nvidia_p2p_dma_map_pages()

int nvidia_p2p_dma_map_pages(struct pci_dev *peer,
        struct nvidia_p2p_page_table *page_table,
        struct nvidia_p2p_dma_mapping **dma_mapping);

该函数使通过nvidia_p2p_get_pages()获取的物理页面可供第三方设备访问。

在那些用于PCIe点对点传输的PCIe资源的I/O地址与CPU访问这些相同资源所使用的物理地址不同的平台上，这是必需的。

在某些平台上，此功能依赖于Linux内核函数dma_map_resource()的正确实现。

函数 nvidia_p2p_dma_unmap_pages()

int nvidia_p2p_dma_unmap_pages(struct pci_dev *peer,
        struct nvidia_p2p_page_table *page_table,
        struct nvidia_p2p_dma_mapping *dma_mapping);

此函数用于解除之前通过nvidia_p2p_dma_map_pages()映射到第三方设备的物理页面。

不应从nvidia_p2p_get_pages()失效回调内部调用。

函数 nvidia_p2p_free_dma_mapping()

int nvidia_p2p_free_dma_mapping(struct nvidia_p2p_dma_mapping *dma_mapping);

此函数旨在从nvidia_p2p_get_pages()失效回调内部调用。

请注意，I/O映射的释放可能会被延迟，例如在从失效回调返回之后。

4.4. 移植到Tegra

GPUDirect RDMA功能在Jetson AGX Xavier平台上从CUDA 10.1开始支持，在基于Linux的DRIVE AGX Xavier平台上从CUDA 11.2开始支持，在Jetson Orin平台上从CUDA 11.4开始支持。从此刻起，本文档将统一把Jetson和Drive称为Tegra。由于Tegra在硬件和软件方面与Linux桌面系统存在特定差异，已开发的应用程序需要稍作修改才能移植到Tegra平台。以下子章节(4.4.1-4.4.3)将简要说明必要的修改内容。

4.4.1. 更改分配器

桌面版GPUDirect RDMA允许应用程序专门操作使用cudaMalloc()分配的GPU内存页。在Tegra平台上，应用程序需要将内存分配器从cudaMalloc()改为cudaHostAlloc()。应用程序可以选择以下方式：

将返回的指针视为设备指针，前提是iGPU支持UVA，或者当使用cudaDeviceGetAttribute()查询iGPU时，cudaDevAttrCanUseHostPointerForRegisteredMem设备属性值为非零。
获取与使用cudaHostGetDevicePointer()分配的主机内存对应的设备指针。应用程序获得设备指针后，适用于标准GPUDirect解决方案的所有规则同样适用于Tegra。

4.4.2. 内核API的修改

以下表格中Tegra API列下的声明可在NVIDIA驱动包中分发的nv-p2p.h头文件中找到。有关参数和返回值的详细说明，请参阅该头文件中的内联文档。下表展示了Tegra平台相较于桌面平台的Kernel API变更。

桌面端API	Tegra API
int nvidia_p2p_get_pages(uint64_t p2p_token, uint32_t va_space_token, uint64_t virtual_address, uint64_t length, struct nvidia_p2p_page_table *page_table, void ( free_callback)(void data), void data);	int nvidia_p2p_get_pages(u64 virtual_address, u64 length, struct nvidia_p2p_page_table *page_table, void (free_callback)(void data), void data);
int nvidia_p2p_put_pages(uint64_t p2p_token, uint32_t va_space_token, uint64_t virtual_address, struct nvidia_p2p_page_table *page_table);	int nvidia_p2p_put_pages(struct nvidia_p2p_page_table *page_table);
int nvidia_p2p_dma_map_pages(struct pci_dev peer, struct nvidia_p2p_page_table page_table, struct nvidia_p2p_dma_mapping **dma_mapping);	int nvidia_p2p_dma_map_pages(struct device dev, struct nvidia_p2p_page_table page_table, struct nvidia_p2p_dma_mapping **dma_mapping, enum dma_data_direction direction);
int nvidia_p2p_dma_unmap_pages(struct pci_dev peer, struct nvidia_p2p_page_table page_table, struct nvidia_p2p_dma_mapping *dma_mapping);	int nvidia_p2p_dma_unmap_pages(struct nvidia_p2p_dma_mapping *dma_mapping);
int nvidia_p2p_free_page_table(struct nvidia_p2p_page_table *page_table);	int nvidia_p2p_free_page_table(struct nvidia_p2p_page_table *page_table);
int nvidia_p2p_free_dma_mapping(struct nvidia_p2p_dma_mapping *dma_mapping);	int nvidia_p2p_free_dma_mapping(struct nvidia_p2p_dma_mapping *dma_mapping);

4.4.3. 其他亮点

请求映射的长度和基地址必须是4KB的倍数，否则会导致错误。
与桌面版不同，在nvidia_p2p_get_pages()注册的回调函数将在调用nvidia_p2p_put_pages()时始终被触发。内核驱动程序有责任通过调用nvidia_p2p_free_page_table()来释放分配的页表。请注意，与桌面版类似，在Unpin Callback中描述的场景下也会触发该回调。
由于cudaHostAlloc()可以通过cudaHostAllocWriteCombined标志或默认标志进行分配，应用程序在将内存映射到用户空间时需要谨慎，例如使用标准的Linuxmmap()函数。在这方面：
1. 当GPU内存被分配为writecombined时，用户空间映射也应通过将vm_area_struct的vm_page_prot成员传递给标准Linux接口`pgprot_writecombine() <https://elixir.bootlin.com/linux/latest/source/arch/arm64/include/asm/pgtable.h#L403>`__来实现writecombined映射。
2. 当GPU内存被分配为默认值时，不应修改vm_page_prot成员中的vm_area_struct。
映射和分配属性的不兼容组合将导致未定义行为。

5. 通知

5.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对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。

5.2. OpenCL

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

5.3. 商标

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