企业官网建设流程全解析

CUDA 多 GPU 点对点传输实战：从 cuda-samples simpleP2P 理解 P2P、UVA 与跨设备寻址

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples

本文以 NVIDIA cuda-samples 仓库中的 simpleP2P 示例 为核心，系统讲解多 GPU 场景下 Peer-to-Peer（P2P）内存拷贝、P2P 寻址（跨设备指针访问）与统一虚拟地址空间（UVA）三大 CUDA 特性的完整用法。通过阅读本文，你将掌握如何在两台以上支持 P2P 的 GPU 之间建立点对点连接、用cudaMemcpyDefault编写与设备无关的数据传输代码、让一个 GPU 上的 kernel 直接读写另一块 GPU 上的显存，并学会利用 CUDA Event 定量测出 P2P 拷贝带宽、验证跨设备计算结果的正确性——这些正是构建多 GPU 计算管线（如流水线并行、数据并行集群）的基础能力。

示例概览：它到底演示了什么

simpleP2P 是 CUDA Samples 中0_Introduction分类下的一个入门级多 GPU 示例，核心目标只有一个：用尽可能少的代码，完整走一遍多 GPU 点对点协作的官方推荐流程。

根据 示例 README 的描述，该应用演示了以下 CUDA 能力：

• P2P copies（点对点拷贝）：两块 GPU 的显存之间直接互相拷贝，不需要经过主机内存中转；
• P2P addressing（点对点寻址）：一块 GPU 上的 kernel 可以直接读写另一块 GPU 显存中的数据；
• Unified Virtual Memory Addressing（UVA，统一虚拟地址空间）：系统内所有 GPU 显存与主机固定内存共享同一个虚拟地址空间，运行时可根据指针自动判断数据的实际归属设备。

从源码结构看（simpleP2P.cu），整个程序是一条清晰的主流程：检查设备数量与 P2P 能力 → 建立双向 P2P 连接 → 分配显存与固定主机内存 → 用事件计时执行 100 次乒乓式 P2P 拷贝并测算带宽 → 跨设备启动 kernel → 回拷数据到主机并逐元素校验 → 拆除连接并清理资源。

值得一提的是，0_Introduction 目录索引 中对 simpleP2P 的定位说明指出：一般来说，P2P 支持存在于两块相同型号的 GPU 之间，但也有一些例外，例如部分 Tesla 与 Quadro 型号。这意味着程序必须在运行时动态探测 P2P 能力，而不能写死设备编号——这也是示例把能力检测放在最前面的原因。

前置条件：硬件、系统与依赖

支持的 GPU 架构（SM 版本）

README 列出该示例支持从 SM 5.0 到 SM 9.0 的广泛架构，覆盖 Maxwell 到 Hopper 之后的多代产品线：

SM 5.0 / 5.2 / 5.3 / 6.0 / 6.1 / 7.0 / 7.2 / 7.5 / 8.0 / 8.6 / 8.7 / 8.9 / 9.0

注意这里列出的 SM 版本仅代表示例代码的兼容范围（示例内核非常基础，没有使用特定架构特性），真正决定能否跑通 P2P 的是当前机器的具体 GPU 型号与拓扑，这一点由程序在运行时通过cudaDeviceCanAccessPeer动态确认。

操作系统与 CPU 架构

• 操作系统：Linux、Windows；
• CPU 架构：x86_64、ppc64le。

依赖项：仅限 64 位系统

README 的 Dependencies 一节引用了仓库根 README 中的 only-64-bit 说明：部分示例只能在 64 位操作系统上运行，simpleP2P 即属于这一类。原因在源码中写得很直白（simpleP2P.cu）：

inline bool IsAppBuiltAs64() { return sizeof(void *) == 8; } ... if (!IsAppBuiltAs64()) { printf("%s is only supported with on 64-bit OSs and the application must be " "built as a 64-bit target. Test is being waived.\n", argv[0]); exit(EXIT_WAIVED); }

程序在启动时检查指针宽度，若被编译为 32 位目标，会直接以EXIT_WAIVED（跳过测试）退出。这一设计的原因是 UVA 需要足够大的虚拟地址空间来统一映射多块显存与主机内存。

构建工具链

• 安装对应平台的 CUDA Toolkit（文章以仓库当前内容为准，仓库要求 CMake 3.20 及以上）；
• 保证满足上述仅 64 位依赖。

另外从 simpleP2P 的 CMakeLists.txt 可以看到，该示例在aarch64平台上不会被构建（CMake 会打印提示Will not build sample simpleP2P - not supported on aarch64），这进一步印证了其支持的 CPU 架构列表（x86_64、ppc64le）是实际生效的约束。

用到的 CUDA API 全景

README 的 "CUDA APIs involved" 一节列出了示例涉及的全部 Runtime API，结合 simpleP2P.cu 源码，可以按职责归为五类：

其中三个 API 是整个示例的灵魂：cudaDeviceCanAccessPeer（探测）、cudaDeviceEnablePeerAccess（建立连接）、cudaMemcpy配合cudaMemcpyDefault（在 UVA 下自动路由数据路径）。

逐步拆解源码：P2P 应用的五个关键阶段

阶段一：设备发现与 P2P 能力矩阵探测

程序首先用cudaGetDeviceCount统计可用 GPU 数量，少于 2 块时直接EXIT_WAIVED跳过测试：

checkCudaErrors(cudaGetDeviceCount(&gpu_n)); if (gpu_n < 2) { printf("Two or more GPUs with Peer-to-Peer access capability are required for %s.\n", argv[0]); exit(EXIT_WAIVED); }

随后遍历所有设备读取属性，并用双重循环对**每一对 GPU（i → j）**调用cudaDeviceCanAccessPeer，打印完整的 P2P 能力矩阵：

for (int i = 0; i < gpu_n; i++) { for (int j = 0; j < gpu_n; j++) { if (i == j) continue; checkCudaErrors(cudaDeviceCanAccessPeer(&can_access_peer, i, j)); printf("> Peer access from %s (GPU%d) -> %s (GPU%d) : %s\n", prop[i].name, i, prop[j].name, j, can_access_peer ? "Yes" : "No"); ... } }

值得注意的两个工程细节：

1. P2P 并不总是对称的——i → j与j → i是两个独立查询，打印时也按方向分别输出Yes/No；
2. 程序只选取检测到的第一对（p2pCapableGPUs[0]与[1]）P2P 能力 GPU 用于后续实验；若全系统不存在任何 P2P 组合，则以EXIT_WAIVED结束。

这一阶段体现了多 GPU 编程的第一原则：永远在运行时探测能力，而不是假设硬件支持。

阶段二：建立双向 P2P 连接

选定 GPU 对后，示例在两个设备上双向调用cudaDeviceEnablePeerAccess，使双方都能访问对方显存：

checkCudaErrors(cudaSetDevice(gpuid[0])); checkCudaErrors(cudaDeviceEnablePeerAccess(gpuid[1], 0)); checkCudaErrors(cudaSetDevice(gpuid[1])); checkCudaErrors(cudaDeviceEnablePeerAccess(gpuid[0], 0));

cudaDeviceEnablePeerAccess的第二个参数是 flags，当前必须传0（保留位）。建立连接后，两块 GPU 的显存互相可见，这是后续 P2P 拷贝与跨设备寻址的前提。与之对称的是程序结尾的清理阶段，通过cudaDeviceDisablePeerAccess逐方向解除连接（同时为非 UVA 场景注销内存映射），见 simpleP2P.cu。

阶段三：UVA 下的内存分配与乒乓式 P2P 拷贝计时

示例在 GPU0 上分配g0、在 GPU1 上分配g1，并在主机侧用cudaMallocHost分配固定内存h0，三者尺寸均为 64MB（1024*1024*16*sizeof(float)）。源码注释特别强调：使用cudaMallocHost分配的内存借助 UVA 自动"可移植"（Automatically portable with UVA），也就是说任何设备都可以直接引用它。

接下来是示例的"表演时刻"——用事件计时，循环 100 次在 GPU0 与 GPU1 之间做乒乓式（ping-pong）双向拷贝：

checkCudaErrors(cudaEventCreateWithFlags(&start_event, cudaEventBlockingSync)); ... checkCudaErrors(cudaEventRecord(start_event, 0)); for (int i = 0; i < 100; i++) { // With UVA we don't need to specify source and target devices, the // runtime figures this out by itself from the pointers if (i % 2 == 0) { checkCudaErrors(cudaMemcpy(g1, g0, buf_size, cudaMemcpyDefault)); } else { checkCudaErrors(cudaMemcpy(g0, g1, buf_size, cudaMemcpyDefault)); } } checkCudaErrors(cudaEventRecord(stop_event, 0)); checkCudaErrors(cudaEventSynchronize(stop_event)); checkCudaErrors(cudaEventElapsedTime(&time_memcpy, start_event, stop_event));

这段代码有三个教学价值极高的点：

• cudaMemcpyDefault是 UVA 的标志性用法：不需要显式指定cudaMemcpyDeviceToDevice还是别的方向，运行时根据源/目的指针所属的虚拟地址区间自动推断数据路径（是同一设备内拷贝、跨设备 P2P 拷贝，还是主机-设备拷贝）。这正是 UVA "统一寻址" 的体现。
• P2P 拷贝不需要切换当前设备：整个循环始终运行在默认设备上下文上，cudaMemcpy自己处理跨设备传输。
• 带宽计算：完成 100 次拷贝（每次 64MB）后，用cudaEventElapsedTime得到总耗时，并按总字节数 / 总时间折算为 GB/s 打印：

printf("cudaMemcpyPeer / cudaMemcpy between GPU%d and GPU%d: %.2fGB/s\n", gpuid[0], gpuid[1], (1.0f / (time_memcpy / 1000.0f)) * ((100.0f * buf_size)) / 1024.0f / 1024.0f / 1024.0f);

事件使用cudaEventBlockingSync标志创建（simpleP2P.cu），使cudaEventSynchronize在等待时让出 CPU 而非忙等——这是多 GPU 场景下避免占满一个 CPU 核的常用技巧。

阶段四：P2P 寻址——kernel 跨设备读写显存

这是示例最核心的演示：GPU1 上运行的 kernel，直接读取 GPU0 上的缓冲g0，并把结果写入 GPU1 自身的缓冲g1；紧接着反转，在 GPU0 上运行的 kernel 读取g1写入g0：

// Run kernel on GPU 1, reading input from the GPU 0 buffer, writing // output to the GPU 1 buffer checkCudaErrors(cudaSetDevice(gpuid[1])); SimpleKernel<<<blocks, threads>>>(g0, g1); checkCudaErrors(cudaDeviceSynchronize()); // Run kernel on GPU 0, reading input from the GPU 1 buffer, writing // output to the GPU 0 buffer checkCudaErrors(cudaSetDevice(gpuid[0])); SimpleKernel<<<blocks, threads>>>(g1, g0); checkCudaErrors(cudaDeviceSynchronize());

内核本身极简（simpleP2P.cu），只是逐元素做dst[idx] = src[idx] * 2.0f，但正因为简单，它干净地证明了：在建立 P2P 连接后，设备端代码可以像访问本地显存一样访问远端 GPU 的显存，无需任何特殊语法或拷贝步骤。这就是 README 中所说的 "Peer-To-Peer (P2P) addressing"。

启动配置为threads(512, 1)、blocks按总元素数除以 512 计算，两次 kernel 之间以cudaDeviceSynchronize保证顺序。

阶段五：回拷、验证与资源清理

两次 kernel 各乘了一次 2.0f，因此 GPU0 上g0中的最终值应为原始数据的 4 倍。程序把结果cudaMemcpy回主机缓冲h0，并逐元素校验：

if (h0[i] != float(i % 4096) * 2.0f * 2.0f) { printf("Verification error @ element %i: val = %f, ref = %f\n", ...); if (error_count++ > 10) break; }

一旦错误元素超过 10 个即提前终止打印，最终依据error_count输出Test failed或Test passed，并设置对应退出码。清理阶段依次：拆除双向 P2P 连接 → 销毁事件 → 释放两块显存与固定主机内存 → 遍历所有设备调用cudaSetDevice(i)恢复现场。

这种"生成输入 → 跨设备计算 → 回拷 → 逐元素校验"的闭环，是 CUDA 示例验证正确性的标准范式，值得在自己的多 GPU 代码中复用。

错误处理与 "Waived" 退出码机制

示例全程使用checkCudaErrors宏包裹所有 CUDA 调用（来自仓库公共头文件 Common/helper_cuda.h）。该宏在调用返回非成功错误码时，打印文件:行号 + 错误枚举 + 调用表达式并以EXIT_FAILURE退出：

#define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__)

同时，示例引入了公共头文件定义的EXIT_WAIVED（Common/helper_cuda.h，值为 2）作为第三态退出码：当环境不满足条件（非 64 位、GPU 少于 2 块、无 P2P 能力组合）时，程序不视为失败，而是以 "跳过测试" 的方式退出。这一机制与仓库根 README 中描述的测试框架（run_tests.py批量运行样本并收集结果）配合使用，避免在无多 GPU 的机器上把示例误报为构建或运行失败。

构建与运行指南

simpleP2P 使用 CMake 构建。仓库根 README.md 给出了标准流程，这里以 Linux 为例（需 CMake 3.20+ 与 CUDA Toolkit）：

mkdir build && cd build cmake .. make -j$(nproc)

示例的 CMakeLists.txt 本身还有几个值得留意的配置点：

• 通过set(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120)预设了一批目标架构，覆盖 Ampere 到 Blackwell 之后的主流算力，可保证生成的 cubin 在当前硬件上直接可用；
• include_directories(../../../Common)引入仓库公共头文件目录，即checkCudaErrors、EXIT_WAIVED等辅助设施的来源；
• 在aarch64系统上该示例会被跳过（CMake 层同样做了平台判断）；
• 若在配置时开启ENABLE_CUDA_DEBUG，会追加-G编译选项以支持 cuda-gdb 调试（默认关闭，默认追加-lineinfo）。

构建完成后，可执行文件位于build/bin/${TARGET_ARCH}/${TARGET_OS}/${BUILD_TYPE}下（Linux x86_64 Release 对应build/bin/x64/linux/release），或直接在 build 目录对应位置运行：

./simpleP2P

典型输出依次为：设备数量统计 → 全对 P2P 能力矩阵（Peer access from ... : Yes/No）→ 建立连接的提示 → 分配缓冲提示 →cudaMemcpy ... between GPUx and GPUy: xx.xxGB/s带宽结果 → 两次跨设备 kernel 运行提示 → 回拷与验证 →Disabling peer access...→ 最终Test passed。

小结：从 simpleP2P 提炼的多 GPU 编程清单

综合 README 与 simpleP2P.cu 源码，可以提炼出一份可直接复用的多 GPU P2P 开发清单：

1. 探测先行：用cudaGetDeviceCount+cudaDeviceCanAccessPeer在运行时确认 P2P 能力，切勿假设硬件支持；
2. 双向建连：在两个设备上分别cudaDeviceEnablePeerAccess，才能实现互相读写；
3. 依赖 UVA：用cudaMallocHost分配可移植主机内存，用cudaMemcpyDefault让运行时自动路由拷贝方向；
4. 跨设备寻址：kernel 直接以远端设备指针作为参数即可完成 P2P 寻址访问；
5. 事件量化：用cudaEventCreateWithFlags(cudaEventBlockingSync)+cudaEventElapsedTime精确测量 P2P 带宽；
6. 闭环验证：回拷主机后逐元素比对，确保跨设备计算正确；
7. 对称清理：结束时cudaDeviceDisablePeerAccess拆除连接，并释放全部资源；
8. 优雅降级：环境不满足时以EXIT_WAIVED退出，把"跳过"与"失败"区分开。

这套流程是理解 CUDA 多 GPU 生态（如streamOrderedAllocationP2P、simpleMultiGPU等其他样本）的基础，掌握了 simpleP2P，就等于掌握了多 GPU 数据通路的第一块基石。

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples