CUDA C++ 编程指引：多设备系统 | CUDA

本系列参考自 CUDA C++ Programming Guide。

枚举设备

主机系统上可以有多个设备。下面的代码展示了怎样枚举这些设备、查询它们的属性、确定有多少个支持 CUDA 的设备。

int deviceCount;
cudaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device) {
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device);
    printf("Device %d has compute capability %d.%d.\n", device, deviceProp.major, deviceProp.minor);
}

设备指定

在任何时候，主机线程都可以使用 cudaSetDevice() 来设置操作的设备。设备内存分配和内核执行都作用在当前的设备上；流和事件关联当前设备。如果没有 cudaSetDevice() 调用，当前设备为 0 号设备。

下面的代码描述了设置当前设备如何影响内存分配和内核执行。

size_t size = 1024 * sizeof(float);
cudaSetDevice(0); // Set device 0 as current
float* p0;
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1

多设备流和事件行为

下面将讨论，如果对一个不属于当前设备的流或事件进行操作，哪些操作会成功，哪些操作会失败：

• 内核函数启动（失败）：如下面的例程所示，如果将内核压入不属于当前设备的流中，则内核会启动失败。也就是说，如果要向一个流中压入内核，必须先切换到流所在的设备。

  cudaSetDevice(0);   // Set device 0 as current
  cudaStream_t s0;
  cudaStreamCreate(&s0);  // Create stream s0 on device 0
  MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 0 in s0
  cudaSetDevice(1);   // Set device 1 as current
  cudaStream_t s1;
  cudaStreamCreate(&s1);  // Create stream s1 on device 1
  MyKernel<<<100, 64, 0, s1>>>(); // Launch kernel on device 1 in s1
  
  // This kernel launch will fail:
  MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0

• 内存拷贝（成功）：如果对一个不属于当前设备的流进行内存拷贝工作，内存拷贝会成功。
• cudaEventRecord()（失败）：必须现将设备上下文切换过去，再向流压入事件。
• cudaEventElapsedTime()（失败）：计算时间差前，必须先切换设备。
• cudaEventSynchronize() and cudaEventQuery()（成功）：即使处于不同的设备，事件同步和事件查询依然有效。
• cudaStreamWaitEvent()（成功）：比较特殊，即使函数输入的流和事件不在同一个设备上，也能成功执行。也就是说，可以让流等待另一个设备上的事件。这个函数可以用作多个设备间的同步。

每个设备有自己的默认流，因此在没有指定流的情况下，向不同设备分派的任务，实际上是压入了各个设备的默认流，他们之间是并行执行的。

对等设备内存访问（Peer-to-Peer Memory Access）

对等设备访问是指运行在一个设备上的内核可以解引用指向另一个设备内存的指针。只要两个设备上的 cudaDeviceCanAccessPeer() 返回true，就支持这种功能。如下例所示，必须通过调用 cudaDeviceEnablePeerAccess() 启用两个设备间的对等设备内存访问支持。

两个设备使用统一内存地址，因为同一指针可用于访问两个设备的内存，如下面的代码所示。

cudaSetDevice(0);   // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size);      // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0);    // Launch kernel on device 0
cudaSetDevice(1);               // Set device 1 as current
cudaDeviceEnablePeerAccess(0, 0);   // Enable peer-to-peer access with device 0

// Launch kernel on device 1
// This kernel launch can access memory on device 0 at address p0
MyKernel<<<1000, 128>>>(p0);

对等设备内存复制（Peer-to-Peer Memory Copy）

可以在两个对等设备间的内存上复制内容。当两个设备使用统一内存地址空间时，使用前面提到的普通的内存拷贝函数即可。否则使用 cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3Dpeer() 或者 cudaMemcpy3DpeerAsync()，如下面的代码所示。

cudaSetDevice(0);   // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size);  // Allocate memory on device 0
cudaSetDevice(1);
float* p1;
cudaMalloc(&p1, size);  // Allocate memory on device 1
cudaSetDevice(0);       // Set Device 0 as Current
MyKernel<<<1000, 128>>>(p0);    // Launch Kernel on Device 0
cudaSetDevice(1);               // Set Device 1 as Current
cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1
MyKernel<<<1000, 128>>>(p1);        // Launch Kernel on Device 1

对于两个不同设备之间的内存复制（在隐式的 NULL 流中）：

• 直到前面发射到任何一个设备的命令执行完，才开始执行。
• 只有在它们执行完之后，后面发射到两者中任一设备的异步命令可开始。

参考

CUDA C++ Programming Guide
《CUDA C Programming Guide》(《CUDA C 编程指南》)导读
《CUDA 并行程序设计：GPU 编程指南》