【IT168技术】本文将介绍CUDA中编程接口的多设备系统，主要包括以下内容：枚举设备、设备选择、流和事件行为、p2p存储器访问、p2p存储器复制。
相关阅读：
1、枚举设备
主机系统上可以有多个设备。下面的代码展示了怎样枚举这些设备、查询他们的属性、确定有多少个支持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);
}
2、 设备选择
在任何时候，主机线程都可以使用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
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
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
3、 流和事件行为
如下面的例程所示，如果内核执行和存储器拷贝发射到非关联到当前设备的流，它们将会失败。如果输入事件和输入流关联到不同的设备，cudaEventRecord()将失败。
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
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
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0
如果两个输入事件关联到不同的设备，cudaEventElapsedTime()将会失败。 即使输入事件关联的设备并非当前设备，cudaEventSynchronize()和cudaEventQuery()也会成功。 即使输入流和输入事件关联到不同的设备，cudaStreamWaitEvent()也会成功。因此cduaStreamWaitEvent()可用于在不同的设备同步彼此。 每个设备有自己的默认流(前文中讲述过，见“相关阅读”中的文章)，因此在一个设备上发射到默认流的一个命令会和发射到另一个设备上默认流中的命令并发执行。
