如何在cuda中的不同gpu之间复制内存

Posted 2023-04-15

【中文标题】如何在cuda中的不同gpu之间复制内存

【英文标题】：How to copy memory between different gpus in cuda

【发布时间】：2015-10-16 03:50:17

【问题描述】：

目前我正在使用两个 gtx 650 。我的程序类似于简单的客户端/服务器结构。我将工作线程分布在两个 gpus 上。服务器线程需要从客户端线程收集结果向量，所以我需要在两个 gpu 之间复制内存。不幸的是，cuda 示例中的简单 P2P 程序无法正常工作，因为我的卡没有 TCC 驱动程序。花了两个小时在谷歌和 SO 上搜索，我找不到答案。一些消息来源说我应该使用 cudaMemcpyPeer ，而其他一些消息来源说我应该使用 cudaMemcpy 和 cudaMemcpyDefault。有没有一些简单的方法可以得到我的工作除了复制到主机然后复制到设备之外完成。我知道它一定已经记录在某个地方，但我找不到它。谢谢你的帮助。

【问题讨论】：

我很确定答案是你不能。您要么需要支持 GPU（Telsa 或 Quadro）的 TTC 驱动程序模式，要么切换到 64 位 linux

所以我唯一的选择是通过 cudaMallocHost 维护一个 cpu 缓冲区，然后从中复制并复制回来，对吗？如果我切换到linux，不同gpus之间复制的正确方法是什么？

p2pBandwidthLatencyTest cuda 示例代码演示了如何在 2 个设备之间进行 memcpy 操作。此代码旨在使用 P2P 机制（如果可用），否则使用“回退”路径（如果不可用）。它不需要您显式维护 CPU 缓冲区，但如果驱动程序使用回退路径，则会“在后台”创建这样的缓冲区。所以@talonmies 给出的评论是正确的。

请注意，该示例代码使用的唯一cudaMemcpy 类型操作是cudaMemcpyPeerAsync。此调用可以使用 P2P 路径（如果可用且已启用），否则使用“回退”路径（如果不可用）。正如@talonmies 所指出的，您需要一个适当的 P2P 环境才能直接从一台设备复制到另一台设备。没有它，副本将遍历主机内存（尽管通过查看 cudaMemcpyPeerAsync 调用并不明显 - 它由驱动程序在后台处理。）

OK，允许客户端和服务器之间交换的数据只有 10000 float/double ，cudaMemcpyPeerAsync 就足够了。

【参考方案1】：

将数据从一个 GPU 传输到另一个 GPU 通常需要通过主机内存“暂存”。例外情况是 GPU 和系统拓扑支持点对点 (P2P) 访问并且已明确启用 P2P。在这种情况下，数据传输可以直接通过 PCIE 总线从一个 GPU 流向另一个。

在任何一种情况下（无论 P2P 是否可用/启用），典型的 cuda runtime API call 将是 cudaMemcpyPeer/cudaMemcpyPeerAsync，如 cuda p2pBandwidthLatencyTest sample code 中所示。

在 Windows 上，P2P 的要求之一是两个设备都由 TCC 模式下的驱动程序支持。在大多数情况下，TCC 模式不是 GeForce GPU 的可用选项（最近，使用 CUDA 7.5RC 工具包中提供的驱动程序和运行时的 GeForce Titan 系列 GPU 例外。）

因此，在 Windows 上，这些 GPU 将无法利用直接 P2P 传输。然而，可以使用几乎相同的序列来传输数据。 CUDA 运行时将检测传输的性质，并“在后台”执行分配以创建暂存缓冲区。然后，传输将分两部分完成：从源设备到暂存缓冲区的传输，以及从暂存缓冲区到目标设备的传输。

以下是一个完整的示例，展示了如何将数据从一个 GPU 传输到另一个 GPU，同时利用 P2P 访问（如果可用）：

$ cat t850.cu
#include <stdio.h>
#include <math.h>

#define SRC_DEV 0
#define DST_DEV 1

#define DSIZE (8*1048576)

#define cudaCheckErrors(msg) \
    do  \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess)  \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
         \
     while (0)


int main(int argc, char *argv[])

  int disablePeer = 0;
  if (argc > 1) disablePeer = 1;
  int devcount;
  cudaGetDeviceCount(&devcount);
  cudaCheckErrors("cuda failure");
  int srcdev = SRC_DEV;
  int dstdev = DST_DEV;
  if (devcount <= max(srcdev,dstdev)) printf("not enough cuda devices for the requested operation\n"); return 1;
  int *d_s, *d_d, *h;
  int dsize = DSIZE*sizeof(int);
  h = (int *)malloc(dsize);
  if (h == NULL) printf("malloc fail\n"); return 1;
  for (int i = 0; i < DSIZE; i++) h[i] = i;
  int canAccessPeer = 0;
  if (!disablePeer) cudaDeviceCanAccessPeer(&canAccessPeer, srcdev, dstdev);
  cudaSetDevice(srcdev);
  cudaMalloc(&d_s, dsize);
  cudaMemcpy(d_s, h, dsize, cudaMemcpyHostToDevice);
  if (canAccessPeer) cudaDeviceEnablePeerAccess(dstdev,0);
  cudaSetDevice(dstdev);
  cudaMalloc(&d_d, dsize);
  cudaMemset(d_d, 0, dsize);
  if (canAccessPeer) cudaDeviceEnablePeerAccess(srcdev,0);
  cudaCheckErrors("cudaMalloc/cudaMemset fail");
  if (canAccessPeer) printf("Timing P2P transfer");
  else printf("Timing ordinary transfer");
  printf(" of %d bytes\n", dsize);
  cudaEvent_t start, stop;
  cudaEventCreate(&start); cudaEventCreate(&stop);
  cudaEventRecord(start);
  cudaMemcpyPeer(d_d, dstdev, d_s, srcdev, dsize);
  cudaCheckErrors("cudaMemcpyPeer fail");
  cudaEventRecord(stop);
  cudaEventSynchronize(stop);
  float et;
  cudaEventElapsedTime(&et, start, stop);
  cudaSetDevice(dstdev);
  cudaMemcpy(h, d_d, dsize, cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy fail");
  for (int i = 0; i < DSIZE; i++) if (h[i] != i) printf("transfer failure\n"); return 1;
  printf("transfer took %fms\n", et);
  return 0;


$ nvcc -arch=sm_20 -o t850 t850.cu
$ ./t850
Timing P2P transfer of 33554432 bytes
transfer took 5.135680ms
$ ./t850 disable
Timing ordinary transfer of 33554432 bytes
transfer took 7.274336ms
$

注意事项：

传递任何命令行参数将禁用 P2P 的使用，即使它可用。 上述结果适用于可以进行 P2P 访问的系统，两个 GPU 通过 PCIE Gen2 链路连接，单向传输带宽约为 6GB/s。 P2P传输时间与此一致（32MB/5ms ~= 6GB/s）。非P2P传输时间更长，但不会翻倍。这是因为对于进出暂存缓冲区的传输，在将一些数据传输到暂存缓冲区之后，可以开始传出传输。驱动程序/运行时利用这一点来部分重叠数据传输。

请注意，一般来说，P2P 支持可能因 GPU 或 GPU 系列而异。在一种 GPU 类型或 GPU 系列上运行 P2P 的能力并不一定表明它可以在另一种 GPU 类型或系列上运行，即使在相同的系统/设置中也是如此。 GPU P2P 支持的最终决定因素是通过cudaDeviceCanAccessPeer 查询运行时提供的工具。 P2P 支持也会因系统和其他因素而异。此处的任何陈述均不保证任何特定设置中的任何特定 GPU 都支持 P2P。

注意：Windows 中的 TCC 驱动程序要求已在最近的驱动程序中放宽。使用最近的驱动程序，只要满足其余要求，就应该可以在 WDDM 模式下的设备之间交换 P2P 数据。

关于 TCC 支持的声明是一般性声明。并非所有 GPU 都受支持。在特定 GPU 上支持（或不支持）TCC 的最终决定因素是nvidia-smi 工具。此处的任何内容都不应被解释为保证在您的特定 GPU 上支持 TCC。

此答案是在 NVLINK 出现以及 CPU 和 GPU 技术的其他变化之前编写的。对于任何给定的系统，系统中任意 2 个 GPU 之间是否可用 P2P 的最终仲裁器/决定因素是 cudaDeviceCanAccessPeer() 返回的结果。其他关于 NUMA 节点和其他系统拓扑特征的一般性陈述是次要的，不应作为最终决定。