4

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

4

1 回答 1

12

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

在任何一种情况下（无论 P2P 是否可用/启用），典型的cuda 运行时 API 调用都是cudaMemcpyPeer/ cudaMemcpyPeerAsync，如 cuda p2pBandwidthLatencyTest示例代码中所示。

在 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
$

笔记：

  1. 传递任何命令行参数将禁用 P2P 的使用，即使它可用。
  2. 以上结果适用于可以进行 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 技术的其他变化之前编写的。对于任何给定的系统，P2P 在系统中任意 2 个 GPU 之间是否可用的最终仲裁器/决定因素是返回的结果cudaDeviceCanAccessPeer()。其他关于 NUMA 节点和其他系统拓扑特征的一般性陈述是次要的，不应用于做出最终决定。

于 2015-07-26T02:13:38.773 回答