如何统计CUDA流上所有核函数的实际总执行时间

原生CUDA事件的能力边界 #

你日常用的cudaEvent_t本身统计的就是GPU侧时间，不是CPU端的墙钟时间，在无高优先级流抢占的常规场景下完全可以满足需求，正确用法如下：

• 向目标流提交算法的第一个核函数前，调用cudaEventRecord(start_ev, target_stream)，将开始标记排入目标流的任务队列
• 正常提交算法所有逻辑，包括cub::DeviceRadixSort这类内部自动启动多个核的库调用，只要调用时传入的是同一个target_stream即可
• 所有核提交完成后，立刻调用cudaEventRecord(end_ev, target_stream)，将结束标记排入流队列尾部
• 等待end_ev同步完成后，调用cudaEventElapsedTime(&cost_ms, start_ev, end_ev)即可拿到耗时

注意：这个结果会包含目标流上相邻核的GPU调度间隙、以及目标流被更高优先级任务抢占时的GPU空闲时间，但不会包含CPU侧运行开销、其他流在两个标记点之外的执行耗时。如果你的生产环境没有开启高优先级流抢占，核之间的调度间隙通常只有几十纳秒，完全可以忽略，这个方案零额外开销，是首选。

如果你的场景存在多流抢占，需要完全排除调度间隙、抢占空闲的影响，精确统计所有核实际占用SM计算资源的纯执行总时长，原生CUDA Runtime API没有提供对应能力，这时候可以用CUPTI实现。

CUPTI 实现精确时长统计 #

CUPTI提供了两种可行的实现路径，根据你的场景选择即可：

• 低开销离线统计：用CUPTI Activity API
  这是生产环境首选的方案，额外开销通常低于1%。你只需要在程序初始化时配置CUPTI追踪核函数执行事件，指定要统计的目标流ID，运行过程中CUPTI会自动记录该流上每一个核函数实际在SM上开始执行、执行完成的GPU时间戳（不是核函数被提交到队列的时间）。等算法执行完成后，把该流下所有对应核的执行时长（结束时间减开始时间）累加，得到的就是纯执行总时长，完全排除了调度间隙、抢占空窗的干扰。
• 实时统计：用CUPTI 事件回调
  如果你需要在算法运行过程中实时拿到时长数据，可以注册CUPTI的核启动、核完成回调，在回调中过滤目标流的核事件，逐段累加每个核的实际执行时长即可。这个方案的回调逻辑在驱动层触发，额外开销在3%~5%左右，只适合调试场景，不建议生产环境使用。

避坑提醒：不要用cudaStreamSynchronize配合系统时钟取前后时间差的方式统计，这个结果会包含所有其他流执行、GPU空闲的时间，误差可以大到数倍，完全不具备参考价值。

内容的提问来源于stack exchange，提问作者Baxissimo