跳到主要内容

使用 MCCL 与 MUSA Graphs

从 MCCL 2.9 开始,MCCL 操作可以被 MUSA 图形捕获。

MUSA 图形提供了一种将工作流定义为图形而不是单个操作的方法。它们可以通过单一 CPU 操作启动多个 GPU 操作来减少开销。有关 MUSA 图形的更多详细信息,可以在 MUSA 编程指南中找到。

MCCL 的集体、点对点和组操作都支持 MUSA 图形捕获。这种支持需要至少 11.3 版本的 MUSA。

操作启动是否被图形捕获被认为是该操作的集体属性,因此必须在所有参与启动的等级上统一(对于集体来说,这是通信器中的所有等级,对于点对点来说,是发送方和接收方)。包含被捕获的 MCCL 操作的图形(通过 musaGraphLaunch 等)的启动被认为是对于在捕获中出现的相同等级集的集体,这些等级中的每一个都必须使用从该集体捕获派生的图形。

以下示例代码显示了如何在 MUSA 图形中捕获计算内核和 MCCL 操作:

musaGraph_t graph;
musaStreamBeginCapture(stream);
kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
mcclAllreduce(..., stream);
kernel_C<<< ..., stream >>>(...);
musaStreamEndCapture(stream, &graph);

musaGraphExec_t instance;
musaGraphInstantiate(&instance, graph, NULL, NULL, 0);
musaGraphLaunch(instance, stream);
musaStreamSynchronize(stream);

从 MCCL 2.11 开始,当 MCCL 通信被捕获并且使用 CollNet 算法时,MCCL 允许通过用户缓冲区注册进一步提高性能。有关详细信息,请参见环境变量 MCCL_GRAPH_REGISTER

支持有多个未完成的 MCCL 操作,这些操作可以是图形捕获的或非捕获的任何组合。需要注意的是,MCCL 内部用来实现这一点的机制在多个通信器的图形从同一线程 musaGraphLaunch() 时,已被观察到会导致 MUSA 死锁。要禁用此机制,请参见环境变量 MCCL_GRAPH_MIXING_SUPPORT