MUSA Event 相关接口使用
2024-5-10
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);
kernel3<<<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生态中是一个不可或缺的组件