PROSAGA码农传奇-调度引擎-Cuda中的warp调度和warp上下文切换之间的关系

<div class =“post-text”itemprop =“text”>
  
    准备好的经线是可以在下一个周期安排的经线。停滞不前的经线无法安排。
  
  
    要通过极其简化的示例回答有关延迟的问题，假设主内存的延迟是8个执行周期，让我们忽略机器是流水线的事实。假设数据准备就绪，假设所有指令都可以在一个周期内执行。
  
  
    现在假设我有这样的C代码：
  
   <pre>
 <code>
 int idx = threadIdx.x+blockDim.x*blockIdx.x;

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

</code>
 </pre>
  
    那是，
     <code>
 myval
 </code>
     当代码完成时，应该包含全局内存中项的平方。这将被分解为汇编语言指令序列。让我们假设他们看起来像这样：
  
   <pre>
 <code>
 I0: R0 = global_data[idx];
I1: R1 = R0 * R0;
I2: ...

</code>
 </pre>
  
    每个线程都可以执行第一行代码（最初没有停顿）;还没有依赖，并且读取本身不会导致停顿。但是，每个线程都可以移动到第二行代码，现在是值的
     <code>
 R0
 </code>
     必须正确，因此发生停顿，等待检索读取。如前所述，假设延迟为8个周期，并且使用32的warp和512的线程块大小，我们总共有16个warp。让我们假设为了简单起见，我们有一个只有32个执行单元的Fermi SM。序列看起来像这样：
  
   <pre>
 <code>
 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 <----

</code>
 </pre>
  
    我们看到的是，在通过执行来自其他warp的指令完成延迟之后，先前“停滞”的warp将重新进入就绪warp池，并且调度程序可以再次调度该warp（即，进行乘法运算）包含在
     <code>
 I1
 </code>
    ）在失速状态消除后的下一个周期。
  
  
    延迟隐藏和warp调度之间没有矛盾。对于具有足够工作量的代码，它们协同工作，以隐藏与各种操作相关的延迟，例如从全局内存中读取。
  
  
    上面的示例是与实际行为相比的简化，但它充分代表了延迟隐藏和warp调度的概念，以演示在“足够的工作”存在的情况下，warp调度如何隐藏延迟。
  
</DIV>