3. 事件与同步原语

多流并发真正难的地方，不是把 kernel 拆到多个流里，而是让这些流在正确的时间点互相“看见”对方的结果。CUDA 提供的同步原语可以分成两类：一类是事件（cudaEvent_t），它像插在流里的时间戳或信号旗，用来记录某个位置是否已经执行到；另一类是同步函数，它让主机或某个流停下来等待。理解它们的语义差别，才能避免“该等的不等、不该等的乱等”。

一、cudaEvent：记录与等待

事件对象由 cudaEventCreate 创建，默认带计时功能；如果只做同步，可以用 cudaEventCreateWithFlags(&ev, cudaEventDisableTiming) 降低开销。事件本身不执行计算，它只被“记录”到某条流的某个位置：

cudaEvent_t ev;
cudaEventCreateWithFlags(&ev, cudaEventDisableTiming);

// 在 streamA 中记录：ev 代表 streamA 中此前所有工作完成
cudaEventRecord(ev, streamA);

// 在 streamB 中等待：streamB 后续工作必须等 ev 就绪
cudaStreamWaitEvent(streamB, ev, 0);

这里有一个关键语义：cudaEventRecord 是异步的，它只是把“记录点”插入流中，主机调用后立即返回；事件真正“就绪”是在该流执行到记录点、且此前所有操作完成后。cudaEventSynchronize(ev) 则是主机阻塞等待事件就绪，常用于计时结束或需要读取结果前。计时场景的典型写法是：

cudaEventRecord(start, streamA);
kernelA<<<grid, block, 0, streamA>>>(...);
cudaEventRecord(stop, streamA);
cudaEventSynchronize(stop);
float ms = 0;
cudaEventElapsedTime(&ms, start, stop);

注意 cudaEventElapsedTime 要求两个事件都已完成，且默认计时事件才有意义。若事件被记录在两条不同流上，计时结果反映的是两条流之间的时间跨度，而不是某条流的纯执行时间，这一点在分析并发重叠时容易误判。

二、流内同步与跨流同步

流内同步是天然顺序：同一条流中的操作按提交顺序执行，前一个 kernel 完成后才会启动下一个。因此同一流内不需要额外事件来保证顺序。真正需要显式处理的是跨流同步：不同流之间默认没有任何顺序保证，streamB 可能在 streamA 的 kernel 还没写完数据时就开始读。

跨流同步最常用的手段就是 cudaStreamWaitEvent。它让目标流在事件就绪前不执行后续操作，但不阻塞主机，也不影响其它流。这使它成为构建依赖 DAG 的核心工具。例如“A 产生数据 → B 消费数据 → C 汇总”的流水线：

cudaEventRecord(evA, streamA);          // A 完成点
cudaStreamWaitEvent(streamB, evA, 0);   // B 等 A
kernelB<<<..., streamB>>>(...);
cudaEventRecord(evB, streamB);          // B 完成点
cudaStreamWaitEvent(streamC, evB, 0);   // C 等 B
kernelC<<<..., streamC>>>(...);

这样三条流可以并发提交，运行时按事件依赖调度，既保留了重叠机会，又保证了数据正确性。相比之下，如果错误地在每次 kernel 后调用 cudaDeviceSynchronize，所有流都会被拉平，并发度直接归零。

三、三种同步函数的适用场景

cudaStreamSynchronize(stream) 阻塞主机，直到指定流中此前所有操作完成。它适合“我只关心这一条流的结果，其它流可以继续跑”的场景，比如主计算流完成后要拷贝结果回主机，而辅助流仍在做日志或预取。

cudaDeviceSynchronize() 阻塞主机，直到设备上所有流的所有操作完成。它语义最强、开销也最大，适合程序收尾、统一错误检查、或确实需要全局一致性的时刻。在性能敏感的热路径中频繁调用它，会把多流并发退化成串行。

cudaEventSynchronize(ev) 只等待某个事件，粒度比前两者更细。它适合“等某个中间里程碑”而不是等整条流或整个设备，例如等待第一批数据就绪后立即启动主机侧后处理。

四、常见误用与排障

第一类误用是事件复用未重置：同一个事件被多次 record，后一次会覆盖前一次的状态，若在两次 record 之间插入 wait，可能等到的是旧记录点。第二类误用是在错误流上 record：事件只代表它被记录的那条流，跨流等待时必须确认 record 的流确实是生产方。第三类误用是用 deviceSynchronize 代替依赖：虽然结果正确，但并发被破坏，表现为“多流后总时间没有下降”。

排查时可以先看时间线：如果各流的总耗时接近串行之和，说明同步点过强；如果出现数据错误或随机结果，说明缺少必要的跨流依赖。一个实用的检查顺序是：确认每个消费者流是否对生产事件调用了 cudaStreamWaitEvent；确认事件是否在正确的流上 record；确认没有在循环体内反复调用 cudaDeviceSynchronize。

把上面的依赖关系落到代码里，一个可复用的模式是：为每个生产阶段创建一个事件，在每个消费流进入前调用一次 cudaStreamWaitEvent，最后只在需要读回结果时对目标流调用 cudaStreamSynchronize。这样既保证了正确性，又让运行时有机会把不同流的工作重叠起来。事件对象在程序结束前用 cudaEventDestroy 释放，避免资源泄漏。

总结来看，事件负责“表达依赖”，同步函数负责“让主机或流等待”。跨流依赖优先用 cudaStreamWaitEvent 表达，主机等待按粒度从细到粗依次选择 cudaEventSynchronize、cudaStreamSynchronize、cudaDeviceSynchronize。把这三者的语义和开销记清楚，多流程序才能在正确的前提下真正跑出并发收益。