Independent Thread Scheduling
Volta架构被设计成比以前的gpu更容易编程，使用户能够在更复杂和多样化的应用程序上高效地工作。Volta GV100是第一个支持独立线程调度的GPU，它支持程序中并行线程之间更细粒度的同步和协作。Volta的主要设计目标之一是减少在GPU上运行程序所需的工作量，并在线程协作方面提供更大的灵活性，从而提高细粒度并行算法的效率。
Prior Nvidia GPU SIMT Models
Pascal及之前的 NVIDIA GPU 以 SIMT(单指令多线程)的方式执行由32个线程组成的warp。Pascal warp在所有的32个thread中共享一个 PC(程序计数器 program counter),同时使用一个 active mask来指定在任意给定的时间内哪些thread处于active状态、这意味着不同的执行分支使一些线程处于 inactive 状态，对warp的不同部分进行序列化执行，如图1所示。原始的mask一直保存，直到warp在分支路径的末尾重新聚合，这时mask被恢复，并且线程再次一起执行。
图1：基于Pascal和早期NVIDIA gpu的SIMT warp执行模型的线程调度。大写字母表示程序伪代码中的语句。warp中的发散分支被序列化，以便在执行另一侧的任何语句之前，分支一侧的所有语句一起执行完成。在else语句之后，warp的线程通常会重新聚合
Pascal SIMT执行模型通过减少跟踪线程状态所需的资源数量和积极地重新聚合线程以最大化并行性来最大化效率。
然而，对整个 warp 的线程状态进行聚合跟踪意味着，当执行路径发生分歧时，执行不同分支的线程将失去并发性，直到它们重新聚合。这种并发性的丧失意味着来自不同分支区域或不同执行状态的同一warp线程不能相互发送信号或交换数据。这就产生了一种不一致性，即来自不同 warp 的线程继续并发运行，而来自同一warp的不同线程则串行化运行，直到它们重新聚合。这意味着，例如，需要细粒度共享由锁或互斥锁保护的数据的算法很容易导致死锁，这取决于竞争线程来自来自哪个 warp。因此，在Pascal和早期的gpu上，程序员必须避免细粒度的同步，或者依赖于无锁或 warp感知的算法。
Volta SIMT Model
Volta 架构通过在所有线程之间启用相同的并发性来改变这一情况，而不管warp。它通过维护每个线程的执行状态(包括程序计数器和调用栈 pc和call stack)来做到这一点，如图2所示。
图2：Volta(下)与Pascal和更早的架构(上)相比的独立线程调度架构框图。Volta维护每个线程的调度资源，比如program counter (PC)和call stack (S)，而早期的架构为每个 warp 维护这些资源。
Volta的独立线程调度允许GPU生成任何线程的执行，或者更好地利用执行资源，或者允许一个线程等待另一个线程产生数据。为了最大限度地提高并行效率，Volta包含了一个调度优化器，它决定如何将来自相同warp的 active 线程分组到SIMT单元中。这保留了与之前的NVIDIA gpu一样的SIMT执行的高吞吐量，但具有更大的灵活性:线程现在可以分散并以 sub-warp 粒度重新收敛，Volta仍然会将执行相同代码的线程分组在一起并并行运行它们。
图1中的代码示例在Volta上的执行看起来有些不同。程序中的if和else分支中的语句现在可以在时间上交替执行，如图3所示。注意，仍然是以SIMT形式执行的:在任何给定的时钟周期，CUDA cores对同一warp中的所有active线程执行相同的指令，就像以前一样，保持了之前架构的执行效率。重要的是，Volta独立调度warp中线程的能力使它能够以一种更自然的方式实现复杂、细粒度的算法和数据结构。尽管调度器支持线程的独立执行，但它对非同步代码进行了优化，以保持尽可能多的收敛，以获得最大的SIMT效率。
图3：Volta独立线程调度允许来自不同分支的语句交错执行。这使得细粒度并行算法的执行成为可能，同一个个warp中的线程可以同步和通信。(这可以达到类似CPU中线程切换执行的效果)
有趣的是，图3并没有显示所有线程同时执行statement Z。这是因为调度器必须保守地假设 Z 可能产生其他发散的执行分支所需的数据，在这种情况下，自动强制重新收敛是不安全的。在A、B、X和Y不包含同步操作的常见情况下，调度器可以确定warp在Z上自然收敛是安全的，就像在以前的架构上一样。
程序可以调用新的CUDA 9 warp同步函数__syncwarp()来强制重新收敛，如图13所示。在这种情况下，warp的发散部分可能不会一起执行Z，但warp中的所有线程的执行路径将在任何线程到达__syncwarp()之后的语句之前完成。类似地，在执行Z之前调用__syncwarp()将强制在执行Z之前重新收敛，如果开发者知道这对他们的应用程序是安全的，可能会提高SIMT效率。
Starvation-Free Algorithms
无饥饿算法是独立线程调度支持的一种关键模式。只要系统确保所有线程都有足够的权限访问竞争的资源，那么Starvation-free algorithms 就是可以保证正确执行的并发计算算法。例如，如果一个线程试图获取互斥锁，并且保证最终会成功，那么互斥锁(或锁)可以用于无饥饿算法中。在不支持饥饿自由的系统中，一个或多个线程可能会反复获取和释放互斥锁，而另一个线程却无法成功获取互斥锁。
考虑一个由Volta的独立线程调度支持的简化示例:在多线程应用程序中将节点插入到一个双链接列表中。
__device__ void insert_after(Node *a, Node *b)
{
Node *c;
lock(a); lock(a->next);
c = a->next;
a->next = b;
b->prev = a;
b->next = c;
c->prev = b;
unlock(c); unlock(a);
}
在本例中，双向链表的每个元素至少有三个组件:一个next指针、一个previous指针和一个提供所有者独占访问以更新节点的锁。图5显示了在节点A之后插入节点B，并更新了A和C的下一个和前一个指针。
图5：带有细粒度锁的双向链表。在将节点B插入到链表(右)之前，获取每个节点的锁(左)。
Volta中的独立线程调度确保了即使一个线程T0当前持有节点a的锁，另一个处于相同waro的线程T1也可以成功地等待锁可用，而不会阻碍线程T0的进程。但是，请注意，由于warp中的active线程一起执行，在锁上自旋的线程可能会降低持有锁的线程的性能。
同样需要注意的是，在上面的例子中使用每个节点的锁对于GPU的性能是至关重要的。传统的双向链表实现可以使用粗粒度的锁，该锁提供对整个结构的独占访问，而不是单独保护单个节点。这种方法通常会导致具有许多线程的应用程序的性能较差——Volta可能有多达163,840个并发线程，——这是由于对锁的高度争用造成的。通过在每个节点上使用细粒度锁，大列表中的平均每个节点争用通常会很低，除非在某些病态节点插入模式下。
这个带有细粒度锁的双链接列表是一个简单的例子，但它展示了独立线程调度如何让开发人员能够以自然的方式在GPU上实现熟悉的算法和数据结构。
