CUDA多GPU编程要点

2022-03-02

总结一下多GPU的相关要点。

1、使用多gpu的原因：

• 单一gpu内存难以加载庞大的数据集
• 提升吞吐量和效率

2、常见的gpu间通信模式

• 各个gpu间没有数据交换，各自执行
• 有数据交换，各gpu间需要冗余数据存储 在这种情况下，必须考虑数据如何在设备之间实现最优移动。要避免通过主机内存中转数据。重要的是要注意有多少数据被传输了和发生了多少次传输。

3、单一节点内遍历并选择设备的代码

cudaError_t cudaGetDeviceCount(int *count);
cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device);
cudaError_t cudaSetDevice(int device);

4、应用程序为64位运行在计算能力4.0以上，连接在同一个PCIe根节点上的gpu，可以通过cuda p2p api直接进行通信，允许两种通信模式：

• 点对点访问：在cuda内核和gpu间直接加载和存储地址
• 点对点传输：在gpu间直接复制数据

gpu连接在不同的PCIe根节点上，仍然可以使用cuda p2p api 进行点对点传输，驱动器将通过主机内存透明地传输数据。 32位应用程序不支持点对点访问

5、启用点对点访问的代码

//检查设备是否支持点对点访问
cudaError_t cudaDeviceCanAccessPeer(int* canAccessPeer, int device, int peerDevice);
//显式启用点对点访问
cudaError_t cudaDeviceEnablePeerAccess(int peerDevice, unsigned int flag);
//关闭点对点访问
cudaError_t cudaDeviceDisablePeerAccess(int peerDevice);

点对点访问的授权是单向的。

6、点对点内存复制

cudaError_t cudaMemcpyPeerAsync(void* dst, int dstDev, void* src, int srcDev, size_t nBytes, cudaStream_t stream);

7、多gpu间的同步 流和事件也适用于多gpu应用程序。典型的工作流程如下所示：

• 选择应用程序将适用的gpu集
• 为每个设备创建流和事件
• 为每个设备分配设备资源
• 通过流在每个gpu上启动任务
• 用流和事件来查询和等待任务完成
• 清空所有设备的资源

只有与流相关联的设备是当前设备时，在流中才能启动内核。只有与该流相关联的设备是当前设备时，才可以在流中记录事件。

这里要格外关注一个概念，即当前设备的概念。

任何时间都可以在任何流中进行内存拷贝，无论该流与什么设备相关联或当前设备是什么。即使流或事件与当前设备不相关，也可以查询或同步它们。

8、一旦选定当前设备，如下cuda运算将被应用到那个设备上：

• 任何从主线程中分配来的设备内存将完全地常驻于该设备上
• 任何由cuda运行时函数分配的主机内存都会有与该设备相关的生存时间
• 任何由主机线程创建的流或事件都会与该设备相关
• 任何由主机线程启动的内核都会在该设备上执行

9、跨设备的内存复制不需要显式地设置当前设备。如果在内存复制前指定了设备，也不会影响它的行为。

10、每个流和事件都与单一设备相联系，但同步流和事件时，并不需要该设备为当前设备。

Refs：

《cuda c编程 权威指南》