Q4：CUDA event 复用时，wait 等的是哪一次 record？
修正先 wait 后 record 的顺序，并用一个 event 为多批数据建立跨 stream 依赖。
假设 e 有效但从未被 record，producer 和 consumer 是同一 GPU 上两个不同的 nonblocking stream；produce 写入 x，consume 读取 x。缓冲区有效，API 调用成功。以下均为教学片段，省略初始化和错误检查，未经 GPU 实测：
cudaStreamWaitEvent(consumer, e, 0);
produce<<<grid, block, 0, producer>>>(x);
cudaEventRecord(e, producer);
consume<<<grid, block, 0, consumer>>>(x);
cudaStreamSynchronize(consumer);Q: 这段代码能保证先写后读吗？如何最小修复，并在多批数据之间复用同一个 event？
A: 不能。从未 record 的 event 表示空工作集；这里的 wait 不会等待未来的 record，因此 produce 和 consume 可能竞争。wait 使用的是 调用 wait API 时最近一次 record 捕获的状态，后续 record 不会改变已经提交的 wait。参见 CUDA 事件文档。
最小修复：把 wait 移到 record 之后。
produce<<<grid, block, 0, producer>>>(x);
cudaEventRecord(e, producer);
cudaStreamWaitEvent(consumer, e, 0);
consume<<<grid, block, 0, consumer>>>(x);
cudaStreamSynchronize(consumer);这样，consumer 后续的 consume 会等事件所覆盖的 produce 完成。末尾同步只负责等待 consumer 完成，不能补上原代码缺失的跨 stream 依赖。参见 CUDA stream 文档。
多批处理时，可以在循环外只创建一次 event。假设由一个 CPU 线程按顺序提交，各个 x[i] 是互不重叠的有效缓冲区，并一直存活到使用它们的工作完成：
cudaEvent_t e;
cudaEventCreateWithFlags(&e, cudaEventDisableTiming);
for (int i = 0; i < batches; ++i) {
produce<<<grid, block, 0, producer>>>(x[i]);
cudaEventRecord(e, producer);
cudaStreamWaitEvent(consumer, e, 0);
consume<<<grid, block, 0, consumer>>>(x[i]);
}
cudaStreamSynchronize(consumer);
cudaEventDestroy(e);每批的 wait 都保留调用时捕获的状态，下一批重新 record 不会把前一批的 wait 改成等待下一批。这复用了 event 资源；event 可以复用，不代表仍在使用的数据缓冲区可以改写或释放。
如果改成先 record 第 0 批、再 record 第 1 批、最后才 wait，wait 引用的就是第 1 批的记录。两次记录都在同一个 producer stream 时，第 1 批记录所覆盖的前序工作也包含第 0 批。参见 cudaEventRecord 文档。