Cuda中warp调度和warp上下文切换之间的关系

Question

据我了解，就绪的warp是可以在warp调度中执行的warp。等待扭曲正在等待获取或计算源操作数，因此无法执行。Warp 调度程序选择一个准备好的 warp 来执行“warp 调度”。

另一方面，当一个 warp 出现管道停顿或全局内存延迟较长时，另一个 warp 将被执行以隐藏延迟。这就是cuda中“warp上下文切换”的基本思想。

我的问题是：Cuda中的warp调度和warp上下文切换之间有什么关系。为了详细说明我的问题，下面是一个例子。

例如，当warp A 停止时，warp A 是等待获取全局内存的warp，一旦获取元素，warp A 将被调度或切换到就绪warp 池中。基于此，warp上下文切换是warp调度的一部分。这是对的吗？

任何人都可以提供有关 Cuda 中的 warp 上下文切换和 warp 调度的任何参考吗？英伟达似乎没有公开这些文件。

预先感谢您的回复。

Answer 1

就绪的扭曲是那些可以安排在下一个周期的扭曲。无法安排停滞的扭曲。

为了用一个极其简单的例子来回答有关延迟的问题，假设主内存的延迟是 8 个执行周期，并且让我们忽略机器是流水线的事实。我们假设如果数据准备好，所有指令都可以在一个周期内执行。

现在假设我有这样的 C 代码：

int idx = threadIdx.x+blockDim.x*blockIdx.x;

int myval = global_data[idx]*global_data[idx];

也就是说，myval当代码完成时，应该包含全局内存中某个项目的平方。这将被分解为一系列汇编语言指令。让我们假设它们看起来像这样：

I0: R0 = global_data[idx];
I1: R1 = R0 * R0;
I2: ...

每个线程都可以执行第一行代码（最初没有停顿）；目前还没有依赖性，并且读取本身不会导致停顿。然而，每个线程都可以继续执行第二行代码，现在 的值R0必须正确，因此会发生停顿，等待检索读取。如前所述，假设延迟为 8 个周期，并使用 32 个扭曲和 512 个线程块大小，我们总共有 16 个扭曲。为了简单起见，我们假设 Fermi SM 只有 32 个执行单元。该序列将如下所示：

cycle:     ready warps:   executing warp:     instruction executed:     Latency:
    0            1-16                   0            I0 -> I1 (stall)    --
    1            2-16                   1            I0 -> I1 (stall)     | --
    2            3-16                   2            I0 -> I1 (stall)     |  |
    3            4-16                   3            I0 -> I1 (stall)     |  |
    4            5-16                   4            I0 -> I1 (stall)     |  |
    5            6-16                   5            I0 -> I1 (stall)     |  |
    6            7-16                   6            I0 -> I1 (stall)     |  |
    7            8-16                   7            I0 -> I1 (stall)     |  |
    8          0,9-16                   8            I0 -> I1 (stall)    <-  |
    9          1,9-16                   0            I1 -> I2            <----

我们看到的是，在通过执行来自其他 warp 的指令来满足延迟后，之前“停滞”的 warp 将重新进入就绪的 warp 池，并且调度程序可以再次调度该 warp（即执行乘法运算）包含在I1) 中，在停顿条件消除后的下一个周期。

延迟隐藏和扭曲调度之间并不矛盾。它们一起工作，让代码有足够的工作要做，隐藏与各种操作相关的延迟，例如从全局内存中读取。

与实际行为相比，上面的示例是一个简化，但它充分代表了延迟隐藏和扭曲调度的概念，以演示扭曲调度如何在存在“足够的工作要做”的情况下隐藏延迟。