MUSA Event 相关接口使用

   更多内容请访问摩尔线程博客中心 

1. 引言

近年来随着科学计算和AI大模型训练的算力需求不断增长，市场对并行计算的算力持续增长，GPU以其高并发、高带宽、易互联的优势在市场中扮演着最重要的角色，摩尔线程（Moore Threads）是一家国产全功能GPU芯片厂商，MUSA是摩尔线程设计推出的统一系统架构，兼容CUDA，依托摩尔线程GPU，使用户更简单高效发挥出GPU的算力。

Event（事件）是MUSA提供的一个便捷的工具，我们可以利用event来统计耗时、同步多stream等

2. 如何使用 MUSA Event

我们先来列举一些MUSA Event常见的runtime接口：

::musaEventCreate
::musaEventCreateWithFlags
::musaEventDestroy
::musaEventRecord
::musaEventQuery
::musaEventSynchronize
::musaEventElapsedTime

首先我们需要使用musaEventCreate来创建一个MUSA event对象，如果有一些特殊需求如IPC等，也可以使用musaEventCreateWithFlags接口传递指定的flag，以创建你需要的event类型；

当我们创建的event使用完毕，我们需要调用musaEventDestroy销毁event，以避免资源泄露。

2.1. 统计设备端耗时

当我们想要统计设备端的耗时，我们可以在创建event之后，在适当的位置调用musaEventRecord接口，注意create event时要enable timing（默认是enable的）。

注意这个接口是异步的，所以它会很快结束返回，如果此时你想查询这个event是否完成了，你可以调用musaEventQuery接口，如果它的返回值是musaSuccess表明结束了，而如果是musaErrorNotReady，即表示还没有完成；

如果你想要等待这个event完成再做其他的操作，那你有两种选择，一是不停调用musaEventQuery直至返回musaSuccess；另一种是调用musaEventSynchronize接口，它是同步接口，会在event执行结束后返回。

当我们想要测量一个kernel或musaMemcpy的耗时时，我们可以创建两个event，分别在它前后record，这样同步后就可以测量出它的耗时了

下图我们示例使用eventSynchronize来同步event测量一个kernel的耗时的原理：

你可以这样编码：

musaEventRecord(event1);
kernel<<<1, 1>>>();
musaEventRecord(event2);
musaEventSynchronize(event2);
float time_ms;
musaEventElapsedTime(&time_ms, event1, event2);

2.2. 同步 stream

当我们使用musa的multiple stream时，如果需要同步不同的stream，我们可以选择使用event来做为同步点，

例如我们创建stream1和stream2，stream1需要等待stream2的一个kernel/memcpy完成后再继续执行后续的task，我们可以选择在stream2上插入一个event同步点（即调用musaEventRecord插入一个event），在stream2上wait这个event（调用musaStreamWaitEvent）

这样我们就无需在host侧去等待有依赖的两条stream同步，这种异步依赖效率会更高，也减少了代码复杂度。

对应我们可以这样设计musa程序：

kernel1<<<1, 1, 0, stream1>>>();
kernel2<<<1, 1, 0, stream2>>>();
musaEventRecord(event1, stream2);
musaStreamWaitEvent(stream1, event1);
kernel<<<1, 1, 0, stream1>>>();

2.3. IPC Event

当你需求多个进程共享MUSA Event时，你可以使用MUSA的IPC Event接口；

A进程使用musaEventCreateWithFlags创建IPC类型的event，之后通过musaIpcGetEventHandle来获取这个event的handle，此时我们就可以通过任何主机侧IPC手段将handle传递给进程B；

B进程收到handle后，可以通过musaIpcOpenEventHandle打开对应的event，如果一切顺利，你就可以在进程B中操作这个event了，比如可以做前面讲到的同步某个stream、query event状态，不过需要注意的是，如果你使用了IPC类型的event，就无法通过这个event获取时间戳了。

进程A

musaEventCreateWithFlags(&event, musaEventDisableTiming | musaEventInterprocess);
musaIpcGetEventHandle(&handle, event);
memcpy(IPC_Shared, &handle, sizeof(musaIpcEventHandle_t));

进程B

memcpy(&ipc_handle, IPC_Shared, sizeof(musaIpcEventHandle_t));
musaIpcOpenEventHandle(&ipc_event, ipc_handle);
musaStreamWaitEvent(stream, ipc_event);

总结

我们看到MUSA Event的用法多种多样，我们可以通过这个工具实现很多DIY的MUSA程序逻辑，在AI大模型训练中我们常能看到多stream使用event做同步和测量耗时的用法，在MUSA生态中是一个不可或缺的组件