cudaEvent_t 是一个 GPU 上（或者说 Stream 上比较好？）的“标记”或“信号”

// 事件与执行状态查询
 
// 查询事件状态（是否完成）非阻塞函数
cudaEventQuery(cudaEvent_t event);
cudaError_t status = cudaEventQuery(event);
if (status == cudaSuccess) {
    // 已完成
} else if (status == cudaErrorNotReady) {
    // 还没完成
}
 
// 等待事件完成
cudaEventSynchronize(cudaEvent_t event);
 
cudaEvent_t event;
cudaEventCreate(&event);      // 创建
cudaEventDestroy(event);      // 销毁
cudaEventCreateWithFlags(&event, cudaEventDisableTiming); // 创建但不记录时间
 
// 在指定 stream 中记录事件（表示当前之前的操作已完成）
// 这个【Record】很有误导性：不是Record Now，而是【把Event插入Stream的末尾】
// 插入一个Event，就是为了等待其在未来发生，如果插入在Now其实也没什么意义
// 另外【Record】并不“记录数据”，它是在 stream 上插一个“标志位”，也不记录错误信息
// 如果 stream 为空，默认是默认流（stream 0）
cudaEventRecord(event, stream);
 
cudaEventSynchronize(event);  // 阻塞 CPU，直到这个 event 所在 stream 的之前操作全部完成
 
float ms; // 注意 ms 是浮点数
cudaEventElapsedTime(&ms, start, stop);  // 计算 start 到 stop 之间的毫秒数
 
// 控制两个 stream 顺序执行
cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);
 
cudaEvent_t event;
cudaEventCreate(&event);
 
someKernel<<<... , ..., 0, s1>>>(...);
cudaEventRecord(event, s1);               // 在 s1 中记录事件
 
cudaStreamWaitEvent(s2, event, 0);        // s2 等待 s1 到达 event 之后再开始执行
 
anotherKernel<<<... , ..., 0, s2>>>(...); // 现在这个 kernel 会等 s1 完成后执行
 
// 自定义的stream之间默认没有绑定
 
// event用于debug
myKernel<<<-1, -1, 0, stream>>>(...);  // 明显非法
cudaEventRecord(event, stream); // 这个 record 会执行失败，event 不会被插入 stream
 
myKernel<<<-1, -1, 0, stream>>>(...);  // 非法，返回 launch error
cudaError_t err = cudaGetLastError();   // 捕获并清除这个错误
printf("%s\n", cudaGetErrorString(err)); // 会打印错误信息
 
cudaEventRecord(event, stream);         // 这个操作与上面的错误状态无关，能继续执行
cudaEventQuery(event);                  // 因为 stream 没有任务，这个 event 会立刻完成
 
myKernel<<<..., ..., 0, stream>>>(...);  // 假设这里有断言失败
cudaEventRecord(event, stream); // event 插入成功，但是永远不会 complete