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