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

Question

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

Answer 1

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

在任何一种情况下（有或没有 P2P available/enabled），典型的 cuda runtime API call would be cudaMemcpyPeer/cudaMemcpyPeerAsync as demonstrated in the cuda p2pBandwidthLatencyTest sample code.

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

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

以下是一个完整的示例，展示了如何将数据从一个 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 link 连接，单向传输带宽约为 6GB/s。 P2P传输时间与此一致（32MB/5ms ~= 6GB/s）。非 P2P 传输时间更长，但不是两倍。这是因为对于传输 to/from 暂存缓冲区，在一些数据传输到暂存缓冲区后，传出传输就可以开始了。 driver/runtime 利用这一点来部分重叠数据传输。

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

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

关于 TCC 支持的说法是笼统的说法。并非所有 GPU 都受支持。在特定 GPU 上支持（或不支持）TCC 的最终决定因素是 nvidia-smi 工具。这里的任何内容都不应被解释为在您的特定 GPU 上支持 TCC 的保证。

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

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

How to copy memory between different gpus in cuda

cuda

multi-gpu