CUDA编程接口：多设备系统详细解析-技术开发专区

CUDA编程接口：多设备系统详细解析

作者：yyfn风辰编辑：王玉圆 2011-12-22 00:05 来源：论坛

　　【IT168技术】本文将介绍CUDA中编程接口的多设备系统，主要包括以下内容：枚举设备、设备选择、流和事件行为、p2p存储器访问、p2p存储器复制。

　　相关阅读：

　　CUDA编程接口：异步并发执行的概念和API

　　CUDA编程接口：分页锁定主机存储器

　　CUDA编程接口：共享存储器实现矩阵相乘

　　CUDA编程接口：运行初始化与设备存储器

　　CUDA编程接口：使用nvcc编译器的兼容性

　　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

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

　　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

// This kernel launch will fail:
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0

　　如果两个输入事件关联到不同的设备，cudaEventElapsedTime()将会失败。即使输入事件关联的设备并非当前设备，cudaEventSynchronize()和cudaEventQuery()也会成功。即使输入流和输入事件关联到不同的设备，cudaStreamWaitEvent()也会成功。因此cduaStreamWaitEvent()可用于在不同的设备同步彼此。每个设备有自己的默认流(前文中讲述过，见“相关阅读”中的文章)，因此在一个设备上发射到默认流的一个命令会和发射到另一个设备上默认流中的命令并发执行。

　　4、 p2p存储器访问

　　当应用以64位进程运行时，以TCC模式在win7/Vista、在win XP或者在Linux上，计算能力2.0或以上，Tesla系列设备能够访问彼此的存储器(即运行在一个设备上的内核可以解引用指向另一个设备存储器的指针)。只要两个设备上的cudaDeviceCanAccessPeer()返回true，这种p2p的存储器访问特性在它们间得到支持。如下例所示，必须通过调用cudaDeviceEnablePeerAccess()启用两个设备间的p2p存储器访问支持。两个设备使用统一存储器地址(后文将讲述)，因为同一指针可用于访问两个设备的存储器，如下面的代码所示。

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);

　　5、 p2p存储器复制

　　可以在两个不同设备间的存储器上复制存储器内容。当两个设备使用统一存储器地址空间(后文将讲述)，使用前文提到的普通的存储器拷贝函数即可。否则使用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); // Set device 1 as current

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

　　两个不同设备之间的存储器复制：

　　①直到前面发射到任何一个设备的命令执行完，才开始执行

　　②任一存储器复制命令开始，后面发射的同步命令(参见3.2.5节)等待它们执行完。

　　注意如果通过如前文描述的cudaDeviceEnablePeerAccess()启用两个个设备间的p2p访问，两个设备间的p2p存储器拷贝就没有必要通过主机进行，因此更快。

更多内容请点击：

CUDA专区：http://cuda.it168.com/

CUDA论坛：http://cudabbs.it168.com/

关注我们