使用 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 上启用 GPUDirect RDMA 连接到 NVIDIA GPU 所需的步骤。

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 传输使用的 API 和控制流程非常相似。

有关更多硬件详细信息，请参阅 支持的系统 和 PCI BAR 大小。

1.2. 标准 DMA 传输

首先，我们概述一个从用户空间发起的标准 DMA 传输。 在这种情况下，存在以下组件

• 用户空间程序

• 用户空间通信库

• 对执行 DMA 传输的设备感兴趣的内核驱动程序

一般顺序如下

1. 用户空间程序通过用户空间通信库请求传输。 此操作采用指向数据的指针（虚拟地址）和大小（以字节为单位）。

2. 通信库必须确保与虚拟地址和大小对应的内存区域已准备好进行传输。 如果还不是这种情况，则必须由内核驱动程序处理（下一步）。

3. 内核驱动程序从用户空间通信库接收虚拟地址和大小。 然后，它要求内核将虚拟地址范围转换为物理页面列表，并确保这些页面已准备好进行传输。 我们将此操作称为固定内存。

4. 内核驱动程序使用页面列表来编程物理设备的 DMA 引擎。

5. 通信库启动传输。

6. 传输完成后，通信库最终应清理用于固定内存的任何资源。 我们将此操作称为取消固定内存。

1.3. GPUDirect RDMA 传输

为了使通信支持 GPUDirect RDMA 传输，必须对上述序列进行一些更改。 首先，存在两个新组件

• 用户空间 CUDA 库

• NVIDIA 内核驱动程序

如 UVA CUDA 内存管理基础知识 中所述，使用 CUDA 库的程序将其地址空间在 GPU 和 CPU 虚拟地址之间拆分，并且通信库必须为它们实现两个单独的路径。

用户空间 CUDA 库提供了一个函数，使通信库能够区分 CPU 地址和 GPU 地址。 此外，对于 GPU 地址，它返回额外的元数据，这些元数据是唯一标识地址表示的 GPU 内存所必需的。 有关详细信息，请参阅 用户空间 API。

CPU 地址和 GPU 地址路径之间的区别在于内存的固定和取消固定方式。 对于 CPU 内存，这由内置的 Linux 内核函数（get_user_pages() 和 put_page()）处理。 但是，在 GPU 内存的情况下，固定和取消固定必须由 NVIDIA 内核驱动程序提供的函数处理。 有关详细信息，请参阅 固定 GPU 内存 和 取消固定 GPU 内存。

一些硬件注意事项在 支持的系统 和 PCI BAR 大小 中进行了解释。

1.4. CUDA 6.0 中的变更

在本节中，我们简要列出 CUDA 6.0 中可用的更改

• CUDA 对等令牌不再是强制性的。 对于调用进程拥有的内存缓冲区（这是典型的），令牌可以在内核模式函数 nvidia_p2p_get_pages() 中替换为零 (0)。 此新功能旨在使现有第三方软件堆栈更容易采用 RDMA for GPUDirect。

• 作为上述更改的结果，引入了一个新的 API cuPointerSetAttribute()。 此 API 必须用于注册任何未使用对等令牌的缓冲区。 在对可能由 RDMA for GPUDirect 读取的内存进行操作时，必须确保 CUDA API 的正确同步行为。 在这些情况下未能使用它可能会导致数据损坏。 请参阅 令牌用法 中的更改。

• cuPointerGetAttribute() 已扩展为返回全局唯一的数字标识符，低级库可以使用该标识符来检测用户级代码中发生的缓冲区重新分配（请参阅 用户空间 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 中检查新缓冲区时。

• 由于 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 地址空间。 主要用例是在平台 I/O 地址的 PCIe 资源（用于 PCIe 对等事务）与 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 中的变更

Jetson AGX Xavier 平台上支持 GPUDirect RDMA。 有关详细信息，请参阅 移植到 Tegra 部分。

1.8. CUDA 11.2 中的变更

Drive AGX Xavier Linux 平台上支持 GPUDirect RDMA。 有关详细信息，请参阅 移植到 Tegra 部分。

1.9. CUDA 11.4 中的变更

添加了一个新的内核模块 nvidia-peermem，它为基于 Mellanox InfiniBand 的 HCA（主机通道适配器）提供对 NVIDIA GPU 视频内存的直接对等读取和写入访问。 有关详细信息，请参阅 使用 nvidia-peermem。

Jetson Orin 平台上支持 GPUDirect RDMA。 有关详细信息，请参阅 移植到 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 module (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. 延迟取消固定优化

在 BAR 中固定 GPU 设备内存是一项昂贵的操作，最多需要毫秒。 因此，应用程序的设计应以最大限度地减少这种开销的方式进行。

使用 GPUDirect RDMA 的最直接的实现是在每次传输之前固定内存，并在传输完成后立即取消固定内存。 不幸的是，这通常会表现不佳，因为固定和取消固定内存是昂贵的操作。 但是，执行 RDMA 传输所需的其余步骤可以快速执行，而无需进入内核（DMA 列表可以缓存并使用 MMIO 寄存器/命令列表重放）。

因此，延迟取消固定内存是实现高性能 RDMA 的关键。 这意味着，即使传输完成，也要保持内存固定。 这利用了一个事实，即很可能同一内存区域将用于未来的 DMA 传输，因此延迟取消固定可以节省固定/取消固定操作。

延迟取消固定的示例实现将保留一组固定的内存区域，并且仅在区域的总大小达到某个阈值，或者由于 BAR 空间耗尽而固定新区域失败时（请参阅 PCI BAR 大小），才取消固定其中一些区域（例如，最近最少使用的区域）。

2.2. 注册缓存

通信中间件通常采用一种称为注册缓存或固定缓存的优化，以最大限度地减少固定开销。 通常，它已经存在于主机内存中，实现了延迟取消固定、LRU 取消注册等。 对于网络中间件，此类缓存通常在用户空间中实现，因为它们与能够进行用户模式消息注入的硬件结合使用。 CUDA UVA 内存地址布局使 GPU 内存固定能够与这些缓存一起工作，只需考虑一些设计注意事项。 在 CUDA 环境中，这一点甚至更重要，因为可以固定的内存量可能比主机内存受到更多限制。

由于 GPU BAR 空间通常使用 64KB 页面进行映射，因此维护四舍五入到 64KB 边界的区域缓存更具资源效率。 更重要的是，因为位于同一 64KB 边界内的两个内存区域将分配并返回相同的 BAR 映射。

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

• 检测所有 CUDA 分配和释放 API。

• 使用标签检查功能来跟踪释放和重新分配。 请参阅 注册缓存的缓冲区 ID 标签检查。

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

虽然拦截 CUDA API 超出了本文档的范围，但从 CUDA 6.0 开始，可以使用一种执行标签检查的方法。 它涉及在 cuPointerGetAttribute()（或 cuPointerGetAttributes()，如果需要更多属性）中使用 CU_POINTER_ATTRIBUTE_BUFFER_ID 属性来检测内存缓冲区释放或重新分配。 如果发生重新分配，API 将返回不同的 ID 值；如果缓冲区地址不再有效，则返回错误。 有关 API 用法，请参阅 用户空间 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，但未明确禁用。 它们在运行时可能无法工作。

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

IOMMU

GPUDirect RDMA 目前依赖于来自不同 PCI 设备角度的所有物理地址都相同。 这使其与执行除 1:1 之外的任何形式转换的 IOMMU 不兼容，因此必须禁用它们或配置为直通转换，GPUDirect RDMA 才能工作。

2.5. PCI BAR 大小

PCI 设备可以请求 OS/BIOS 将物理地址空间区域映射到它们。 这些区域通常称为 BAR。 NVIDIA GPU 当前公开多个 BAR，其中一些 BAR 可以支持任意设备内存，从而使 GPUDirect RDMA 成为可能。 GPUDirect RDMA 的最大可用 BAR 大小因 GPU 而异。 例如，目前 Kepler 类 GPU 上可用的最小 BAR 大小为 256 MB。 其中，32MB 目前保留供内部使用。 这些大小可能会更改。

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

2.6. 令牌用法

警告

从 CUDA 6.0 开始，令牌应被视为已弃用，但仍受支持。

正如在 用户空间 API 和 内核 API 中可以看到的那样，一种用于固定和取消固定内存的方法除了 GPU 虚拟地址之外还需要两个令牌。

这些令牌 p2pToken 和 vaSpaceToken 对于唯一标识 GPU VA 空间是必需的。 单独的进程标识符无法标识 GPU VA 空间。

令牌在单个 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 驱动程序将 内核 API 限制为拥有内存分配的进程。

2.7. 同步和内存排序

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

• 正确注册 CUDA 内存的 BAR 映射是必需的，以确保该映射与 CUDA API 对该内存的操作保持一致。

• 只有 CUDA 同步和工作提交 API 提供 GPUDirect RDMA 操作的内存排序。

CUDA API 一致性注册

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

此行为在每个分配粒度上启用，可以通过调用带有 CU_POINTER_ATTRIBUTE_SYNC_MEMOPS 属性的 cuPointerSetAttribute() 来实现，或者在使用旧版路径时为缓冲区检索 p2p 令牌。有关更多详细信息，请参阅 用户空间 API。

一个示例情况是 cuMemcpyDtoD() 和后续在副本目标上执行的 GPUDirect RDMA 读取操作之间的读后写依赖关系。作为一种优化，设备到设备内存复制通常在将复制排队到 GPU 调度器后异步返回给调用线程。但是，在这种情况下，这将导致通过 BAR 映射读取不一致的数据，因此禁用此优化，并在 CUDA API 返回之前完成复制。

用于内存排序的 CUDA API

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

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

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

3. 如何执行特定任务

3.1. 显示 GPU BAR 空间

从 CUDA 6.0 开始，NVIDIA SMI 实用程序提供了转储 BAR1 内存使用情况的功能。它可以用于了解 BAR 空间的应用程序使用情况，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 内存

1. 正确的行为需要对内存地址使用 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 path
           pass_to_kernel_driver(address, size);
       } else {
           // CPU path
           // ...
       }
   }
   

   这是必需的，以便 CUDA 驱动程序以特殊方式处理 GPU 内存缓冲区，从而保证 CUDA 内存传输始终与主机同步。有关 cuPointerSetAttribute() 的详细信息，请参阅 用户空间 API。

2. 在内核驱动程序中，调用 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
       }
   }
   

   请注意，在调用分页锁定函数之前，起始地址如何对齐到 64KB 边界。

   如果函数成功，则内存已被分页锁定，并且 page_table 条目可用于编程设备的 DMA 引擎。有关 nvidia_p2p_get_pages() 的详细信息，请参阅 内核 API。

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() 的详细信息，请参阅 内核 API。

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

3.4. 处理释放回调

1. 如果 NVIDIA 内核驱动程序需要撤销映射，它将调用 nvidia_p2p_get_pages() 调用中指定的 free_callback(data)。有关详细信息，请参阅 内核 API 和 取消分页锁定回调。

2. 回调等待挂起的传输，然后清理页表分配。

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

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

3.5. 用于注册缓存的缓冲区 ID 标记检查

请记住，对于延迟敏感型实现，不建议使用围绕缓冲区 ID 标记检查构建的解决方案。相反，建议对 CUDA 分配和释放 API 进行检测，以向注册缓存提供回调，从而从关键路径中删除标记检查开销。

1. 首次遇到设备内存缓冲区并识别为尚未分页锁定时，将创建分页锁定的映射，并检索关联的缓冲区 ID 并将其一起存储在缓存条目中。cuMemGetAddressRange() 函数可用于获取整个分配的大小和起始地址，然后可用于分页锁定它。由于 nvidia_p2p_get_pages() 将需要对齐到 64K 的指针，因此直接对齐缓存地址很有用。此外，由于 BAR 空间当前以 64KB 的块进行映射，因此将整个分页锁定舍入到 64KB 更具资源效率。

   // struct buf represents an entry of the registration cache
   struct buf {
       CUdeviceptr pointer;
       size_t      size;
       CUdeviceptr aligned_pointer;
       size_t      aligned_size;
       int         is_pinned;
       uint64_t    id; // buffer id obtained right after pinning
   };
   

2. 创建后，每次使用注册缓存条目时，都必须首先检查其有效性。执行此操作的一种方法是使用 CUDA 提供的缓冲区 ID 作为标记来检查释放或重新分配。

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

   当缓冲区标识符更改时，相应的内存缓冲区已被重新分配，因此相应的内核空间页表将不再有效。在这种情况下，内核空间 nvidia_p2p_get_pages() 回调将被调用。因此，缓冲区 ID 提供了一个标记，以使分页锁定缓存与内核空间页表保持一致，而无需内核驱动程序向上调用到用户空间。

   如果 cuPointerGetAttribute() 返回 CUDA_ERROR_INVALID_VALUE，则程序应假定内存缓冲区已被释放，或者不是有效的 GPU 内存缓冲区。

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

   // in the registration cache code
   if (buf->is_pinned && !buf_is_gpu_pinning_valid(buf)) {
     regcache_invalidate_entry(buf);
     pin_buffer(buf);
   }
   

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

1. 运行提取脚本

   ./NVIDIA-Linux-x86_64-<version>.run -x
   

   这会提取 NVIDIA 驱动程序和内核包装器。

2. 导航到输出目录

   cd <output directory>/kernel/
   

3. 在此目录中，为您的内核构建 NVIDIA 模块

   make module
   

   完成此操作后，您的内核构建目录下的 Module.symvers 文件包含 nvidia.ko 的符号信息。

4. 使用以下行修改您的内核模块构建过程

   KBUILD_EXTRA_SYMBOLS := <path to kernel build directory>/Module.symvers
   

3.7. 使用 nvidia-peermem

NVIDIA GPU 驱动程序包提供了一个内核模块 nvidia-peermem，它为基于 NVIDIA InfiniBand 的 HCA（主机通道适配器）提供对 NVIDIA GPU 视频内存的直接对等读取和写入访问。它允许基于 GPUDirect RDMA 的应用程序将 GPU 计算能力与 RDMA 互连一起使用，而无需将数据复制到主机内存。

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

NVIDIA OFED（开放光纤企业分发）或 MLNX_OFED 在 InfiniBand Core 和对等内存客户端（例如 NVIDIA GPU）之间引入了一个 API。nvidia-peermem 模块通过使用 NVIDIA GPU 驱动程序提供的对等 API 向 InfiniBand 子系统注册 NVIDIA GPU。

内核必须具有对 RDMA 对等内存的必要支持，可以通过内核的附加补丁或通过 MLNX_OFED 作为加载和使用 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 驱动程序，以确保使用 MLNX_OFED 提供的 RDMA API 编译 nvidia-peermem。

4. 参考资料

4.1. UVA CUDA 内存管理的基础知识

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

CUDA VA 空间寻址

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

GPU

由 GPU 内存支持的页面。主机无法访问此页面，并且所讨论的 VA 永远不会在主机上具有物理支持。从 CPU 取消引用指向 GPU VA 的指针将触发段错误。

CPU

由 CPU 内存支持的页面。主机和 GPU 都可以通过相同的 VA 访问此页面。

FREE

这些 VA 由 CUDA 保留，用于将来的分配。

这种分区允许 CUDA 运行时通过内存对象在保留的 CUDA VA 空间内的指针值来确定内存对象的物理位置。

地址在页面粒度上细分为这些类别；页面内的所有内存类型都相同。请注意，GPU 页面的大小可能与 CPU 页面的大小不同。CPU 页面通常为 4KB，而 Kepler 类 GPU 上的 GPU 页面为 64KB。GPUDirect RDMA 专门在 CUDA VA 空间内的 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

如果 pointer 指向 GPU 内存，并且 CUDA 驱动程序能够为整个设备内存分配设置新行为。

任何其他值

如果 pointer 指向 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);

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

参数

data [out]

带有两个令牌的结构 CUDA_POINTER_ATTRIBUTE_P2P_TOKENS。

attribute [in]

在 GPUDirect RDMA 范围内应始终为 CU_POINTER_ATTRIBUTE_P2P_TOKENS。

pointer [in]

一个指针。

返回值

CUDA_SUCCESS

如果 pointer 指向 GPU 内存。

任何其他值

如果 pointer 指向 CPU 内存。

可以随时调用此函数，包括在 CUDA 初始化之前，并且它具有 CUDA 同步行为，如 CU_POINTER_ATTRIBUTE_SYNC_MEMOPS 中所示，因此应被认为是昂贵的，并且每个缓冲区应仅调用一次。

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

请注意，出于安全原因，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]

指向将存储缓冲区 ID 的 64 位变量的指针。

attribute [in]

CU_POINTER_ATTRIBUTE_BUFFER_ID 枚举器。

pointer [in]

指向 GPU 内存的指针。

返回值

CUDA_SUCCESS

如果 pointer 指向 GPU 内存。

任何其他值

如果 pointer 指向 CPU 内存。

以下是一些一般性说明

• cuPointerGetAttribute() 和 cuPointerSetAttribute() 仅是 CUDA 驱动程序 API 函数。

• 特别是，cuPointerGetAttribute() 不等同于 cudaPointerGetAttributes()，因为所需的功能仅存在于前一个函数中。这绝不会限制 GPUDirect RDMA 可能使用的范围，因为 cuPointerGetAttribute() 与 CUDA 运行时 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 虚拟内存范围下的页面。

警告

这是一个昂贵的操作，应尽可能少地执行 - 请参阅 延迟取消分页锁定优化。

函数 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 资源的 I/O 地址（用于 PCIe 对等事务）与 CPU 用于访问这些相同资源的物理地址不同的平台上，这是必需的。

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

函数 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 在 CUDA 10.1 开始的 Jetson AGX Xavier 平台、CUDA 11.2 开始的基于 DRIVE AGX Xavier Linux 的平台以及 CUDA 11.4 开始的 Jetson Orin 平台上受支持。从那时起，本文档将 Jetson 和 Drive 统称为 Tegra。由于 Tegra 相对于 Linux-Desktop 的硬件和软件特定差异，已开发的应用程序需要稍作修改才能将其移植到 Tegra。以下小节 (4.4.1-4.4.3) 简要介绍了必要的更改。

4.4.1. 更改分配器

Desktop 上的 GPUDirect RDMA 允许应用程序专门在使用 cudaMalloc() 分配的 GPU 页面上运行。在 Tegra 上，应用程序将必须将内存分配器从 cudaMalloc() 更改为 cudaHostAlloc()。应用程序可以

1. 将返回的指针视为设备指针，前提是 iGPU 支持 UVA 或 cudaDevAttrCanUseHostPointerForRegisteredMem 设备属性在使用 cudaDeviceGetAttribute() 查询 iGPU 时为非零值。

2. 获取使用 cudaHostGetDevicePointer() 分配的 host 内存对应的设备指针。一旦应用程序拥有设备指针，则适用于标准 GPUDirect 解决方案的所有规则也适用于 Tegra。

4.4.2. 内核 API 的修改

下表 Tegra API 列下的声明可以在 NVIDIA 驱动程序包中分发的 nv-p2p.h 头文件中找到。有关参数和返回值的详细描述，请参阅该头文件中包含的内联文档。下表表示 Tegra 相对于 Desktop 的内核 API 更改。

Desktop 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. 其他亮点

1. 请求的映射长度和基地址必须是 4KB 的倍数，否则会导致错误。

2. 与 Desktop 版本不同，当调用 nvidia_p2p_put_pages() 时，始终会触发在 nvidia_p2p_get_pages() 处注册的回调。内核驱动程序有责任通过调用 nvidia_p2p_free_page_table() 来释放 page_table 分配。请注意，与 Desktop 版本类似，在 取消分页锁定回调 中解释的场景中也会触发回调。

3. 由于可以使用 cudaHostAllocWriteCombined 标志或默认标志分配 cudaHostAlloc()，因此当将内存映射到用户空间时（例如使用标准 linux mmap()），应用程序应谨慎行事。在这方面

   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 内存分配为 default 时，不应对 vm_area_struct 的 vm_page_prot 成员进行任何修改。

   map 和分配属性的不兼容组合将导致未定义的行为。

5. 声明

5.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 对客户因本文所述产品而承担的总体和累积责任应根据产品的销售条款进行限制。

5.2. OpenCL

OpenCL 是 Apple Inc. 的商标，已授权 Khronos Group Inc. 使用。

5.3. 商标

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