CUDA C++ Programming Guide——编程接口 CUDA Runtime Multi-Device System

分手后的思念是犯贱 2022-12-18 05:59 152阅读 0赞

### Multi-Device System ###

#### Device Enumeration ####

一个主机系统可以有多个设备。 下面的代码示例演示如何枚举这些设备，查询它们的属性以及确定启用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);
    }

#### Device Selection ####

主机线程可以随时通过调用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

#### Stream and Event Behavior ####

如果将内核启动发布到与当前设备不相关的流，则启动失败，如以下代码示例所示。

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()和cudaEventQuery()也会成功。即使输入流和输入事件关联到不同的设备，cudaStreamWaitEvent()也会成功。 因此，**cudaStreamWaitEvent()可以用于使多个设备彼此同步**。

每个设备都有其自己的默认流（请参阅默认流），因此，发布到设备的默认流的命令可能相对于发布到任何其他设备的默认流的命令无序或同时执行。

#### Peer-to-Peer Memory Access ####

取决于系统属性，特别是PCIe和/或NVLINK拓扑，设备能够寻址彼此的内存（即，在一个设备上执行的内核可以取消dereference对指向另一设备的内存的指针的引用）。 如果cudaDeviceCanAccessPeer()对这两个设备返回true，则在两个设备之间支持此**对等内存访问功能**。对等内存访问仅在64位应用程序中受支持，并且必须通过调用cudaDeviceEnablePeerAccess()在两个设备之间启用，如以下代码示例所示。 在未启用NVSwitch的系统上，每个设备最多可支持整个系统范围内的八个对等连接。

两个设备都使用统一的地址空间（请参阅[统一虚拟地址空间][Link 1]），因此可以使用相同的指针从两个设备中寻址内存，如下面的代码示例所示。

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

##### IOMMU on Linux #####

仅在Linux上，CUDA和显示驱动程序不支持启用IOMMU的裸机PCIe对等内存复制IOMMU-enabled bare-metal PCIe peer to peer memory copy。 但是，CUDA和显示驱动程序确实通过VM传递支持IOMMU。 因此，Linux上的用户在本机裸机系统上运行时应禁用IOMMU。 应该启用IOMMU，并将VFIO驱动程序用作虚拟机的PCIe直通。  
在Windows上，不存在上述限制。另请参阅[在64位平台上分配DMA缓冲区][64_DMA]。

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

两个不同设备的内存之间的复制（在隐式NULL流中）：在先前向任一设备发出的所有命令都已完成并且在复制到任一设备的复制开始之前发出的任何命令（请参见异步并发执行）之前运行完成之前，不会启动。does not start until all commands previously issued to either device have completed and runs to completion before any commands (see Asynchronous Concurrent Execution) issued after the copy to either device can start.

与流的正常行为一致，两个设备的内存之间的异步复制可能与另一个流中的复制或内核重叠。请注意，如果如[对等内存访问][Link 2]中所述通过cudaDeviceEnablePeerAccess()在两个设备之间启用了对等访问，则这两个设备之间的对等内存复制不再需要通过主机。

[Link 1]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#unified-virtual-address-space
[64_DMA]: https://download.nvidia.com/XFree86/Linux-x86_64/396.51/README/dma_issues.html
[Link 2]: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#peer-to-peer-memory-access