我的问题是关于使用函数 cudaEventElapsedTime 来测量多流应用程序中的执行时间。 根据CUDA文档
如果任一事件最后记录在非 NULL 流中，结果时间可能比预期的长（即使两者使用相同的流句柄）。发生这种情况是因为 cudaEventRecord() 操作异步发生，不能保证测量的延迟实际上只是在两个事件之间。许多其他不同的流操作可以在两个测量的事件之间执行，从而改变时间安排很重要。
我真的很难理解上面的粗体句子。看来，使用默认流来测量时间更准确。但我想明白为什么？如果我想测量流中的执行时间，我发现通过该流而不是默认流附加启动/停止事件更合乎逻辑。请澄清一下？谢谢
【问题讨论】：
首先让我们记住基本的CUDA stream 语义：
CUDA default stream（假设我们没有覆盖默认的legacy行为）还有一个隐式同步的特性，这大致意味着一个CUDA操作发出到默认流将不会开始执行，直到所有之前发布的 CUDA 活动到该设备完成。
因此，如果我们将 2 个 CUDA 事件（例如，启动和停止）发布到旧版默认流中，我们可以确信在这两个发布点之间发布的任何和所有 CUDA 活动都会被计时（无论它们是从哪个流中发出的，或者它们是从哪个主机线程发出的）。我建议随意使用，这是直观的，不太可能被误解。此外，它应该产生一致的时序行为，run-to-run（假设主机线程行为相同，即以某种方式同步）。
OTOH，假设我们有一个多流应用程序。假设我们将内核发布到 2 个或更多非默认流中：
Stream1: cudaEventRecord(start)|Kernel1|Kernel2|cudaEventRecord(stop)
Stream2: |Kernel3|
这些是从同一个主机线程还是从不同的主机线程发出的并不重要。例如，假设我们的单主机线程活动看起来像这样（浓缩）：
cudaEventRecord(start, Stream1);
Kernel1<<<..., Stream1>>>(...);
Kernel2<<<..., Stream1>>>(...);
Kernel3<<<..., Stream2>>>(...);
cudaEventRecord(stop, Stream1);
我们应该期待什么时间？ Kernel3 是否会包含在 start 和 stop 之间的经过时间中？
事实上，答案是未知的，并且可能因运行而异，并且可能取决于在上述活动之前和期间设备上发生的其他情况。
对于上述发布顺序，假设我们在设备上没有其他活动，我们可以假设在cudaEventRecord(start) 操作之后，Kernel1 将立即启动并开始执行。让我们假设它“填充设备”，因此没有其他内核可以同时执行。我们还假设Kernel1 的持续时间比Kernel2 和Kernel3 的启动延迟要长得多。因此，当Kernel1 正在执行时，Kernel2 和Kernel3 都在排队等待执行。在Kernel1 完成时，设备调度程序可以选择开始任一 Kernel2 或Kernel3。如果它选择Kernel2，那么在Kernel2完成时，它可以将stop事件标记为已完成，这将确定start和stop之间的持续时间为Kernel1和Kernel2的持续时间， 大约。
Device Execution: event(start)|Kernel1|Kernel2|event(stop)|Kernel3|
| Duration |
但是，如果调度程序选择在 Kernel2 之前开始 Kernel3（基于流语义的完全合法且有效的选择），则在 Kernel2 完成之前，stop 事件无法标记为完成，这意味着测量的持续时间现在将包括Kernel1 加上Kernel2 加上Kernel3 的持续时间。 CUDA 编程模型中没有任何东西可以解决这个问题，这意味着测量的时序甚至可以交替运行：
Device Execution: event(start)|Kernel1|Kernel3|Kernel2|event(stop)|
| Duration |
此外，我们可以大大改变实际的发布顺序，将Kernel3 的发布/启动在 第一个cudaEventRecord 或在 最后一个cudaEventRecord ，而上述论点/变量仍然成立。这就是cudaEventRecord调用的异步性质的意义所在。它不会阻塞CPU线程，但就像内核启动一样，它是异步的。因此，上述所有活动都可以在它实际开始在设备上执行之前发出。即使Kernel3 在第一个cudaEventRecord 之前开始执行，它也会占用设备一段时间，延迟Kernel1 的开始执行，因此会增加一定量的测量持续时间。
如果Kernel3即使在最后一个cudaEventRecord之后发出，因为所有这些发出操作都是异步的，Kernel3可能仍然排队等待Kernel1完成时准备去，这意味着设备调度程序仍然可以选择启动哪个，从而可能产生可变时间。
当然还有其他类似的危险可以被绘制出来。这种在多流场景中发生变化的可能性导致了保守的建议，即避免尝试使用发布到非传统默认流中的事件进行基于 cudaEvent 的计时。
当然，例如，如果您使用visual profiler，那么在两个事件之间测量的内容应该没有什么歧义（尽管它可能仍会因运行而异）。但是，如果您要使用可视化分析器，则可以直接从时间线视图中读取持续时间，而无需调用事件经过时间。
请注意，如果您覆盖默认流遗留行为，则默认流大致等同于“普通”流（尤其是对于单线程主机应用程序）。在这种情况下，我们不能依赖默认的流语义来解决这个问题。一种可能的选择是在任何cudaEventRecord() 调用之前使用cudaDeviceSynchronize() 调用。我并不是建议对所有可能的情况进行分类，但对于单设备单主机线程应用程序，它应该等同于将cudaEvent 定时发布到默认遗留流中。
可能最好使用分析器来完成复杂场景的时间安排。许多人还完全放弃了基于cudaEvent 的计时方法，转而使用高分辨率主机计时方法。无论如何，复杂的并发异步系统的时间安排是不平凡的。保守的建议是为了避免随意使用其中一些问题。
【讨论】：