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

Question

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

Answer 1

将数据从一个 GPU 传输到另一个 GPU 通常需要通过主机内存进行“分段”。异常(exception)情况是 GPU 和系统拓扑支持对等 (P2P) 访问并且已明确启用 P2P。在这种情况下，数据传输可以直接通过 PCIE 总线从一个 GPU 流向另一个 GPU。
在任何一种情况下(有或没有 P2P 可用/启用)，典型的 cuda runtime API call 将是 cudaMemcpyPeer/cudaMemcpyPeerAsync，如 cuda p2pBandwidthLatencyTest sample code 所示。
在 Windows 上，P2P 的要求之一是两个设备都在 TCC 模式下由驱动程序支持。在大多数情况下，TCC 模式不是 GeForce GPU 的可用选项(最近，使用 CUDA 7.5RC 工具包中可用的驱动程序和运行时的 GeForce Titan 系列 GPU 异常(exception)。)
因此，在 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 节点和其他系统拓扑特征的一般性陈述是次要的，不应用于做出最终决定。