具有非默认流的 cudaElapsedTime

Question

我的问题是关于使用函数cudaEventElapsedTime来测量多流应用程序中的执行时间。根据CUDA文档

如果任一事件最后记录在非 NULL 流中，则结果时间可能会比预期时间长（即使两者使用相同的流句柄）。发生这种情况是因为 cudaEventRecord() 操作是异步发生的，并且不能保证测量的延迟实际上只是在两个事件之间。任意数量的其他不同的流操作可以在两个测量事件之间执行，从而以显着的方式改变时间。

我真的很难理解上面粗体的句子。看来，使用默认流来测量时间更准确。但我想明白为什么？如果我想测量流中的执行时间，我发现通过该流而不是默认流附加开始/停止事件更符合逻辑。请澄清一下吗？谢谢

Answer 1

首先让我们记住基本的CUDA 流语义：

1. 发布到同一流中的 CUDA 活动将始终按发布顺序执行。
2. 发布到单独流中的 CUDA 活动的执行顺序之间没有定义的关系。

CUDA默认流（假设我们没有覆盖默认的遗留行为）具有隐式同步的附加特征，这大致意味着发出到默认流的 CUDA 操作将不会开始执行，直到所有先前向该设备发出的 CUDA 活动都已完成。

因此，如果我们向遗留默认流发出 2 个 CUDA 事件（例如，开始和停止），我们可以确信在这两个问题点之间发出的任何和所有 CUDA 活动都将被计时（无论它们被发出到哪个流，或者它们是从哪个主机线程发出的）。我建议对于随意使用来说，这是直观的，并且不太可能被误解。此外，它应该产生一致的定时行为，运行到运行（假设主机线程行为相同，即以某种方式同步）。

OTOH，假设我们有一个多流应用程序。假设我们将内核发布到 2 个或更多非默认流中：

Stream1:  cudaEventRecord(start)|Kernel1|Kernel2|cudaEventRecord(stop)
Stream2:                                |Kernel3|

这些是从同一主机线程还是从单独的主机线程发出并不重要。例如，假设我们的单主机线程活动如下所示（精简）：

cudaEventRecord(start, Stream1);
Kernel1<<<..., Stream1>>>(...);
Kernel2<<<..., Stream1>>>(...);
Kernel3<<<..., Stream2>>>(...);
cudaEventRecord(stop, Stream1);

我们应该期待什么时间？将包含在和Kernel3之间经过的时间？startstop

事实上，答案是未知的，并且可能因运行而异，并且可能取决于在上述活动之前和期间设备上发生的其他情况。

对于上述发出命令，假设我们在设备上没有其他活动，我们可以假设在操作之后cudaEventRecord(start)，将立即Kernel1启动并开始执行。让我们假设它“填充设备”，以便没有其他内核可以同时执行。我们还假设 的持续时间比和Kernel1的启动延迟长得多。因此，当正在执行时， 和都排队等待执行。完成后，设备调度程序可以选择开始或。如果它选择，那么在完成时它可以将事件标记为已完成，这将把和之间的持续时间近似确定为和的持续时间。Kernel2Kernel3Kernel1Kernel2Kernel3Kernel1 Kernel2Kernel3Kernel2Kernel2stopstartstopKernel1Kernel2

Device Execution: event(start)|Kernel1|Kernel2|event(stop)|Kernel3|
                              |    Duration   |

Kernel3但是，如果调度程序选择在此之前开始Kernel2（基于流语义的完全合法且有效的选择），则事件在完成stop之前无法标记为完成，这意味着测量的持续时间现在将包括plus plusKernel2的持续时间。CUDA 编程模型中没有任何内容可以解决这个问题，这意味着测量的时序甚至可以在运行之间交替：Kernel1Kernel2Kernel3

Device Execution: event(start)|Kernel1|Kernel3|Kernel2|event(stop)|
                              |    Duration           |

此外，我们可以大大改变实际的发行顺序，将发行/启动放在第一个Kernel3 之前cudaEventRecord或最后一个之后cudaEventRecord，并且上述参数/可变性仍然成立。这就是调用的异步性质的意义cudaEventRecord所在。它不会阻塞 CPU 线程，但就像内核启动一样，它是异步的。因此，上述所有活动都可以在其实际开始在设备上执行之前发出。即使在Kernel3第一个 之前开始执行cudaEventRecord，它也会占用设备一段时间，从而延迟 的开始执行Kernel1，从而使测量的持续时间增加一定量。

如果Kernel3即使在最后一个之后也发出cudaEventRecord，因为所有这些发出操作都是异步的，所以完成Kernel3时可能仍会排队并准备就绪Kernel1，这意味着设备调度程序仍然可以选择要启动哪个，从而可能会发生变化定时。

当然还有其他类似的危险可以绘制出来。多流场景中的这种变化可能性引起了保守的建议，即避免尝试cudaEvent使用发布到非传统默认流中的事件来进行基于定时的操作。

当然，如果您使用可视化分析器，那么两个事件之间测量的内容应该相对较少有歧义（尽管它可能仍然因运行而异）。但是，如果您要使用可视化分析器，则可以直接从时间线视图中读取持续时间，而无需调用事件经过时间。

请注意，如果您覆盖默认流遗留行为，则默认流大致相当于“普通”流（特别是对于单线程主机应用程序）。在这种情况下，我们不能依赖默认的流语义来解决这个问题。一种可能的选择是在任何cudaEventRecord()呼叫之前进行cudaDeviceSynchronize()呼叫。我并不是建议这解决所有可能的情况，但对于单设备单主机线程应用程序，它应该相当于cudaEvent发布到默认旧流中的计时。

复杂的场景计时可能最好使用分析器来完成。许多人还完全放弃cudaEvent基于计时并恢复到高分辨率主机计时方法。无论如何，复杂的并发异步系统的计时都非常重要。保守的建议旨在避免临时使用时出现的一些问题。