cudaEventRecord更适合测时，因其异步打点不阻塞主机线程，配合cudaEventElapsedTime可精确获取GPU真实执行耗时；而cudaEventQuery轮询会引入CPU开销与计时抖动，破坏微秒级精度。
cudaEventRecord 为什么比 cudaEventQuery 更适合测时
因为 cudaEventRecord 是异步打点，不阻塞主机线程，而 cudaEventSynchronize（或 cudaEventElapsedTime 内部调用的同步逻辑）才真正等待事件完成。如果用 cudaEventQuery 轮询判断完成，不仅浪费 CPU，还引入计时抖动——尤其在多核/高负载下，轮询间隔不可控，微秒级精度就崩了。
实操建议：
- 始终用
cudaEventRecord记录起止点，再用cudaEventElapsedTime一次性读取毫秒值（注意：它返回的是 float，单位是毫秒，需乘以 1000 得到微秒） - 避免在记录事件后立刻调用
cudaEventSynchronize再手动算差值，这多一次同步，反而增加不确定性 - 确保两个事件属于同一 CUDA stream；跨 stream 默认不保证顺序，测出的时间可能为 0 或负值
为什么 cudaEventElapsedTime 返回值常是 0.0 或极小浮点数
这是最常踩的坑：cudaEventElapsedTime 要求两个事件都已记录（recorded），且第二个事件必须发生在第一个之后（按 CUDA 时间线）。如果核函数执行太快（比如只有几十纳秒），而 GPU 时钟分辨率有限（常见为 ~0.5–1 μs），cudaEventElapsedTime 就会向下取整为 0.0 ms → 0 μs。
解决办法：
- 用循环重复 launch 同一 kernel N 次，测总耗时再除以 N（N ≥ 100，推荐 1000）
- 确认 GPU 支持高精度事件：运行
nvidia-smi -q -d CLOCK查看graphics_clock频率，高频卡（如 A100/H100）事件分辨率更高 - 不要用
cudaEventCreate，改用cudaEventCreateWithFlags并传入cudaEventBlockingSync—— 这不会提升精度，但能排除因驱动异步调度导致的“未记录”误判
正确测时的最小完整代码结构
以下是最简可靠模板，省略错误检查但保留关键约束：
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
<p>// 注意：必须在同一 stream，显式指定比默认 stream 更可控
cudaStream_t stream;
cudaStreamCreate(&stream);</p><p>cudaEventRecord(start, stream);
my_kernel>>();
cudaEventRecord(stop, stream);</p><p>// 此处才真正同步并计算，不提前干预
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("time: %.3f μs\n", milliseconds * 1000.0f);</p><p>cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaStreamDestroy(stream);
</p>
关键点：
- 所有
cudaEventRecord必须带 stream 参数，哪怕用0（即默认 stream），也得写清楚 - 核函数 launch 和
cudaEventRecord(stop)之间不能插入任何可能触发隐式同步的操作（如cudaMemcpy、cudaDeviceSynchronize） - 如果 kernel 有参数依赖（比如输入数据刚从 host memcpy 过来），必须确保 memcpy 也提交到同一 stream，并在
cudaEventRecord(start)之前完成
实际微秒级测量仍不准？先查这三个地方
即使代码结构正确，真实环境里仍可能看到波动 >1 μs，大概率是这些干扰项：
-
nvcc编译时没加-use_fast_math或-Xptxas -v，导致 kernel 实际指令数浮动，影响执行稳定性 - GPU 处于动态调频状态（
nvidia-smi -q -d POWER查看power draw是否稳定），建议用nvidia-smi -r重置，再用nvidia-smi -pl XXX锁定功耗墙 - host 端用了
std::chrono::high_resolution_clock混合测量（比如测 launch 到 event 完成的总开销），它和 CUDA event 不同源，不能交叉比对
微秒级精度本质是硬件+驱动+代码三者对齐的结果，少一个环节，测出来的就只是“大致快慢”，不是“精确耗时”。
C++免费学习笔记（深入）：立即使用
在学习笔记中，你将探索 C++ 的入门与实战技巧！
