关于GPU一些笔记（SIMT方面)

GPU组成

《计算机组成原理 — GPU 图形处理器》已经大概说明出GPU一般都是由比CPU多的core组成，而每个core 相当于一个单独线程进行计算，并且可以同时触发执行相同的单一指令但是每个计算单元数据不同(称之为SIMD)的指令执行。在英伟达GPU中 core一般称之为之为cuda core，GPU内部一般集成了成千上万个cuda core。为了方便进行进行对这么多的核进行管理调度，GPU将按照一定数量的core(一般为32或者64）组成一个 SM(streaming multiprocessor 称之为流式多处理器），而在AMD存在同样的架构一般称之为CU(compute units 计算单元)。

A100 SM架构图

 硬件实现过程中，为了尽可能提高其硬件性能ALU会按照其计算数据类型一般分为 special function unit(SFU)、integer function unit（INT 32 整型计算单元）、float-point function unit(浮点型计算单元），其中浮点型计算单元可以分为FP32和FP64。除了上述单元之外还有用于专门加载和保存指令的处理单元LD/ST，专门用于从内存中加载和保存数据。故GPU内部一条指令一般都按照其操作的数据以及指令性质，会分为单独的硬件功能专门执行各自部分，以尽量提高执行效率效率。

由上图可知，A100中每个SM中含有64个INT32，64个FP32以及32个FP64，即相当于在计算INT32类型数据时一个SM 可以同时支持64个整型计算并发。

一个GPU一般由多个SM按照一定的组织形式组成。

SIMT

GPU 真正的调度执行模式是按照SIMT(single-instruction multiple-thread)模式，即每个SIMT包含一定数量的core用于执行（GPU中的core一定属于某个SIMRT），在NV中SIMT被称之为wrap, AMD中称之为wavefront，它是硬件调度的并行的最小单位，即一个SMIT 可以同时并行运行的线程数目。而SM一般由多个或者一个warp组成，一般程序开发人员感知不到warp的存在，只能感知到SM,一个通用别的GPU架构一般由下图组成：

A gengric modern GPU architecture

 

多个SIMT core 组成 SIMT Core Cluster即相当于 SM或者CU，SIMT Core Cluster进一步组成整个GPU，内部含有memory 控制器用于外界 off-line 内存。

整个GPU core的划分一般都遵循上述层次划分，经过层次划分之后才便于管理这个硬件的调度及执行状况。

与CPU不同的是， GPU的执行调度单位为warp或者wavefront，而CPU调度可以精细化为单个线程，这是因为CPU同时执行的线程数目比较少，而GPU要同时管理成千上万个线程。

GPU内存选择

GPU 由于同时有成千上万个线程同时访问内存，故其外挂内存一般采用高吞吐量（high throughput) 特点内存 以应对同时较大内存访问量，故一般采用GDDR内存。而CPU场景简单，同时只应对少量线程并发访问，要求延迟较低（low latency),其内存类型一般使用DRAM。

上述情况出现在CPU和GPU分离状况（discrete GPU)即 GPU和 CPU不在一个芯片中，而在集成的GPU中即CPU和GPU集成在一个芯片中例如AMD APU芯片，由于两个芯片被集成到一个芯片内，其内存类型一般采用DRAM，低延迟类型芯片，因为CPU和GPU之间的内存不需要来回进行传输拷贝。

 SIMT core架构

SIMT 内部core的一般通用架构设计主要如下图所示：

整个GPU按照流水线可以将SIMT core划分为SIMT front-end 和 SIMD back-end两个部分 ，整个流水线处理划分三级，每个级别为一个循环处理：

• 取指令循环（an instruction fetch loop, one loop)：主要包括 Fetch、I-Cache、Decode、I-Buffer几个部分。
• 发送指令循环（an instruction issue loop, two loop)：主要包括I-Buffer、Scoreboard、Issue、SIMT Stack几个部分
• 获取寄存器调度循环（a register access scheduling loop, three loop)：主要包括Operand Collector、ALU以及Memory几个部分。

上述几个部分主要功能模块：

• Fetch: 负责读取内存中的指令。
• I-Cache: 即instruction cache，为L1 缓存，从内存中读取的指令放入到L1缓存中以便下次能够快速获取到指令。
• Decode:负责解析读取到的指令，包含解析出指令需要使用到的源寄存器和目的寄存器 以及操作指令。
• I-Buffer: instruction buffer ，为L1缓存，用于缓存一些指令，主要用于识别缓存中这些指令之间是否有依赖，如果指令之间没有依赖则可以乱序执行，以最大可能提高效率。
• Score Board: 指令打分系统，通过该模块用于识别指令之间的依赖。
• SIMT-Stack: 用于解决同一个warp内分支问题，通过mask 在分支执行过程中，将不需要执行的线程通过mask 去掉。
• Issue:指令发送模块。
• Operand Collector:主要用于解决在cycle 切换时通过进行warp切换，寄存器过多而造成port过多问题，通过统一的bank来解决次问题。
• ALU：计算单元，该数目与warp 大小相等，按照其功能和数据划分INT32, FP32，FP64等
• MEM：GPU内存管理模块，包括与外部所接off-line模块。

ONE-LOOP APPROXIMATION

one-loop approximation可以认为是 instruction fetch loop。

上面提到GPU内部含有很多core，为了方便进行对这些core进行调度，将一定数量的core组成warps进行调度（AMD 称为wavefronts)，在同一个warp内的core使用相同的PC(program counter)指针即在同一个warp内的所有core同一时刻执行相同的指令。Fetch模块根据PC指向的指令从内存中获取到相应的指令（PC指针指向的指令是下一个将要执行的指令），之后SIMT将获取到的指令进行译码，并根据译码之后的结果从register file从获取到寄存器，与此同时 SIMT 进行mask处理，在接下来的处理流程中，根据SIMT mask结果，进行预测，哪些线程(core)将会执行。

该指令得到调度之后，SIMT中的core(图中的ALU)根据mask 结果按照SIMD(single instruction multiple data)模型进行执行，即warp内的每个core或者也可以称为ALU执行相同指令，但是数据不一样。同时GPU为了尽可能提高执行效率，将执行的指令按照功能或者执行的数据划分为专门的硬件单元进行执行，例如NV中将ALU划分为 SFU(specical function unit)、load/store unit、floating-point function unit、integer function unit，V100中还专门划分了Tensor 处理单元用于深度学习，即每个单元仅仅执行指令中的一部分功能。

SIMT 中的内硬件单元（即core）的数量 ？

SIMT内core的数量即warp 大小相等一般与lane相等， 各个GPU在时钟周期（clock cycles)内使用不同的warp 执行策略，可以通过提高时钟频率以提高warp内切换线程数目来取得较高的性能。而提高SIMT性能 有两种方法分为为增加流水线深度或者提高时钟频率（以增加单位时间内能够执行更多的指令）。

该循环为一个SIMT 流水线执行过程的大概逻辑过程，为整个SIMT整个执行流程。

SIMT STACK

SIMT执行模型是GPU中的一个重要特性，甚至可以说是GPU整个调度核心所在，SIMT可以使开发人员从功能角度来来各个线程之间是独立的，各个线程之间相互独立执行相同的指令。但是在实际开发过程中，不可避免的要使用到一些分支处理，即有些线程执行一个分支，而另外一些线程执行另外一个分支，这样就造成了线程之间执行指令并不相同，而SIMT模型中可以知道，在同一个warp内只有一个PC指针，如果在同一个warp内的线程内分别执行不同的分支，就会造成执行指令分化（称之为线程束分化，可以参考下面文章https://zhikunhuo.blog.csdn.net/article/details/105167996），线程束的分化要尽量避免，要解决线程束分化问题不仅仅要从上层应用开发解决，整个GPU硬件内部调度角度也需要解决该问题。

GPU SIMT执行模型为了解决上述问题，引入了SIMT STACK模块。

SIMT stack主要用来解决两个关键问题：

1：控制流嵌套问题（nested control flow): 当kernel存在多个分支，且其中一个分支依赖与另外一个分支，严重影响了线程独立性

2：如何跳过计算过程（skip computation)：由于分支的存在，造成了在同一个warp内的有些线程并不必要执行一些计算指令。

SIMT stack中使用了SIMT mask来解决上述问题，以下使用一个例子说明mask 如何来控制整个warp执行。

do {
   t1 = tid*N; // A
   t2 = t1 + i;
   t3 = data1[t2];
   t4 = 0;
   if(t3 != t4) {
       t5 = data2[t2]; // B
       if( t5 != t4 ) {
           x += 1; // C
        } else {
           y += 2; // D
        }
                   //E
    } else {
        z += 3; // F
    }
    i++; // G
} while( i < N );

现在假设有上述一个 GPU端 kernel代码，在while循环中对应两个if else分支，分别为t3是否等于t4，以及t5是否等于t4两个判断。分别给上述代码标记几个关键指令，分别为A、B、C、D 、E、F、G等几个关键点。其中A指令为一个warp内所有线程都可以执行到，而B、C、D、F分别对应的四个分支各自处理指令， 而E和G处为指令位置称之为reconvergence point 位置，该位置为分支处理完毕之后，线程又重新聚合在一起进行执行，该点对于GPU对分支处理非常关键，其中E处是没有任何指令执行即 C和D两份分支处理完毕之后，聚合在一起的线程没有执行任何指令算法。对于上述例子，如果代码在同一个warp内执行需要处理4种分支情况即有多个线程分别执行不同的指令，出现线程束分化情况。

为了解决上述分支情况，SIMT stack模块使用mask来进行标记，其中置为1表示该线程需要执行，置0表示该线程不需要执行，每一个bit位代表了一个线程。假设一个GPU，warp大小为4 即同时支持4个线程并行运行，当上述kernel代码运行到A处时，所有的线程都运行，则标记位”A/1111“，其中mask bit位从低位到高位分别表示四个线程从第一个到第四个线程。当代码运行到if(t3 != t4) 时出现分支造成线程分化，其中第四个线程运行F分支，第一到三个线程运行B分支，则可分别表示"F/0001"和”B/1110"。处于B线程继续执行又遇到if( t5 != t4 )分支，其中第一个线程执行C，第二和三个线程执行D，mask可分别表示位"C/1000"和"D/0110"，当C和D执行完毕之后到E线程有重新聚合，即"E/1110"。而在G出 又出现了线程聚合，即"G/1111"。整个上述kernel执行控制流图（control flow graph)CFG，如下图所示：

GPU上述执行遇到上述分支时，只能采取顺序执行，即先执行其中一个分支然后再执行另外一个分支，当执行其中一个分支时 根据提前标记好的mask决定线程是否运行，如果在该分支有些线程mask 被置0，则只能进行等待状态。这时因为SIMT执行模式在同一个warp内只有一个PC执行即在warp内所有线程都只能执行这一条命令只能选择执行或者不执行，而不能在不执行时选择其他指令进行执行。即遇到分支时，warp只能按照顺序串行化进行执行，直到所有线程都执行完毕。上述kernel 执行时，在SIMT stack其他预测的执行顺序为：

 上述执行图种实心箭头代表线程执行，空箭头代表线程不执行被屏蔽掉，从上到下依次表示从第一个到第四个线程。

而在SIMT stack中为了指明整个序列号执行路径，按照一条条entry记录，使用类似如下图所示内容：

其中每个entry有三项分别为 reconvergence program counter(RPC或者Ret/Reconv PC，聚合点PC)、the address of the next instruction to execute(Next PC 要执行的下一条PC)、Active Mask(记录当前指令哪些线程执行，哪些线程被屏蔽），其中 TOS 即为 top of stack处于栈顶位置，指向的当前指令对应的entry，上述i 和ii 以及iii分别代表三个分支的entry。每个entry中如果Next PC 和Ret/Reconv PC不相等则说明存在分支情况（注意这是个人看法，暂时还没有从相关论文中找到这个明确说法）。

上述SIMT stack中在初始化c图部分，TOS指向下一个Next PC为B，即下一个要执行的指令为B出现分支，执行B指令的mask为1110。而B 指令的entry指向 d图的TOS，即指向完B之后 开始执行C 执行而B执行聚合点指令为E，即执行完C之后之后还要执行聚合点同样为E的指令 即D指令，只有聚合点都是E的指令执行完毕之后，才能进行往下执行，而 执行的指令则为E指令。此时同样E指令的聚合点指令为G，需要把所有聚合点位G的指令执行完毕之后 才能执行G 指令。

那么现在出现一个问题：出现分支之后，要先执行哪一个分支？具体采用什么样的准则？

AMD的官方文档中给出了建议，NV采用类似方法：

To reduce the maximum depth of the reconvergence stack to be logarithmicn in the number of threads in a warp it is best to put the entry with the most active threads on the stack first and then the entry with fewer active threads.

为了减少 reconvergence stack的深度，先执行活跃线程最多的分支，然后再执行活跃少的分支。

SIMT mask问题->死锁（deadlock)

SIMT mask可以解决warp内分支执行问题，通过串行执行完毕分支之后，线程在reconverge point时有重新聚合在一起以便最大提高其并行能力。但是对于一个程序来说，如果出现分支就意味这每个分支的指令和处理不一致，这就容易造成对一些共享的数据的一致性，失去一致性就意味着在同一个warp内的如果存在分支则线程之间不能够交互或者交互数据，在一些实际算法中为了进行数据交互 则需要使用lock机制，而mask恰恰会因为调度问题造成一个死锁deadlock问题.

The Pascal SIMT execution model maximizes efficiency by reducing the quantity of resources required to track thread state and by aggressively reconverging threads to maximize parallelism.Tracking thread state in aggregate for the whole warp, owever, means that when the execution pathway diverges, the threads which take different branches lose concurrency until they reconverge. This loss of concurrency means that threads from the same warp in divergent regions or different states of execution cannot signal each other or exchange data. This presents an inconsistency in which threads from different warps continue to run concurrently, but diverged threads from the same warp run sequentially until they reconverge. This means, for example,that algorithms requiring fine-grained sharing of data guarded by locks or mutexes can easily lead to deadlock, depending on which warp the contending threads come from. Therefore, on Pascal and earlier GPUs, programmers need to avoid fine-grained synchronization or rely on lock-free or warp-aware algorithms.

下面一个例子能够说明SIMT mask 调度引起的deadlock问题，该例子为一个cuda用例来源于：https://stackoverflow.com/questions/6426793/realistic-deadlock-example-in-cuda-opencl

__global__ kernel() {
  __shared__ int semaphore;
  semaphore=0;
  __syncthreads();
  while (true) {
    int prev=atomicCAS(&semaphore,0,1);
    if (prev==0) {
      //critical section
      semaphore=0;
      break;
    }
  }
}

 介绍该例子之前 首先需要明白cuda中原子操作函数系列函数，cuda kernel函数中 算法中经常会使用到共享全局变量，在warp内会经常遇到对同一个位置的共享变量进行操作（当然实际开发kernel函数要尽量避免这种情况），为了避免多线程对同一资源使用造成不一致问题，cuda提供了atomic原子操作一系列操作，避免对同一资源进行同时访问修改，atomic相当于一个自旋锁。

int atomicCAS(int* address, int compare, int val);

atomicCAS函数为将address地址的值与 compare比较，如果相等则将address地址的值设置为val，并返回address地址的旧值。

上述kernel函数中，针对一个处于共享内存的整型值semaphore，其值初始化为0，在同一个warp内的所有线程都可以访问到该值。进入到循环中首先为了防止内存一致性问题，使用atomicCAS函数，如果semaphore值为0，则将其设置为1，pre值为0然后如果此时其他线程也会调用int prev=atomicCAS(&semaphore,0,1); 指令但是此时semaphore被第一个首先获取到的线程设置为1，那么此时semaphore值为1，不等于0，pre返回的值为1。也即意味着如果warp size大小为32，则只能由一个线程进入到pre==0分支的流程中，而其他31个分支不能进入prev==0线程，一直在循环，此时出现两个分支其中一个分支激活线程（入prev=0）为1个，另外31个线程继续循环。 进入循环的第一个线程，则执行关键代码段，一般用来防止其他线程进行同时访问，以保护一些数据，运行完成之后将semaphore重新设置为0，并退出循环，其他31个线程又有一个线程能够进入到prev=0分支，剩余30个线程继续运行，如此进行反复运行，知道所有线程串行执行完毕。

上述方法 一般时用来执行一些只能由一个线程运行的代码段，以防止产生数据不一致现象，咋一看似乎没有问题，但是上述一个简单场景恰似一个经典的死锁deadlock问题。下面是对这个问题的详细描述：

The atomicCAS instruction ensures that exaclty one thread gets 0 assigned to prev, while all others get 1. When that one thread finishes its critical section, it sets the semaphore back to 0 so that other threads have a chance to enter the critical section.

The problem is, that while 1 thread gets prev=0, 31 threads, belonging to the same SIMD unit get a value 1. At the if-statement CUDA scheduler puts that single thread on-hold (masks it out) and let other 31-threads continue their work. In normal circumstances it is a good strategy, but in this particular case you end up with 1 critical-section thread that is never executed and 31 threads waiting for infinity. Deadlock.

Also note, the existence of break which leads the control flow outside of the while loop. If you ommit the break instruction and have some more operations after the if-block that are supposed to be executed by all threads, it may actually help the scheduler avoid the deadlock.

造成上述死锁deadlock问题原因恰是SIMT mask调度引起的。上述代码逻辑在CPU执行过程中没有问题，但是在GPU执行过程中由于出现分支，SIMT stack的调度原则是首先执行活跃数线程最多的分支，上述例子是首先执行分支为31个线程的情况，而一直循环，而另外一个prev返回0的分支一直由于上述31个线程一直处于循环之中而造成得不到调度，这样会一直处于一个dealock状态，而无法继续向前执行。

造成上述问题的根本原因就是在warp内的所有线程只有一个PC指针，无法同时执行其他命令，而这种SIMD模型正式GPU的核心所在，看似这种问题是属于GPU SIMT执行模型的天生问题无法解决

Ahmed ElTantawy and Tor M. Aamodt等人在《MIMD Synchronization on SIMT Architectures》论文中明确指出了天生的deadlock问题，该论证中使用的模型如下：

和上述例子中使用的基本一样。

deadlock问题长期依赖困扰着 GPU应用开发人员，长期依赖于开发人员的细心度。NV一直到 V100 GPU 中才得到解决，《NVIDIA TESLA V100 GPU ARCHITECTURE》中，NV将解决deadlock问题作为一个亮点，V100白皮书中指出了对上述问题对SIMT stack模块进行了改进：

 V100 中NV 为warp内每个线程都分配了一个PC指针和Stack，这样将PC指针的颗粒度细化到了每一个线程中去，但是这样就于GPU 的根基SIMT执行模型又有明显的冲突，因为这样每个线程都有自己的PC，岂不是和CPU没什么本质上的差别。为了解决上述问题，在V100内部调用中，硬件还是使用的warp这一单位进行调度线程，V100内部中使用了a schedule optimizer硬件模块决定哪些线程可以在一个warp内进行调度(这样就涉及到另外一种技术rearrange thread 稍后再讲)，讲相同的指令重新进行组织排布到一个warp内，执行SIMD模型，以保证最大利用效率。下面为V100白皮书中的原话：

Volta’s independent thread scheduling allows the GPU to yield execution of any thread, either to make better use of execution resources or to allow one thread to wait for data to be produced by another. To maximize parallel efficiency, Volta includes a schedule optimizer which determines how to group active threads from the same warp together into SIMT units. This retains the high throughput of SIMT execution as in prior NVIDIA GPUs, but with much more flexibility: threads can now diverge and econverge at sub-warp granularity, while the convergence optimizer in Volta will still group together threads which are executing the same code and run them in parallel for maximum efficiency

V100可以将分支组织成一个sub-warp来保证处于同一分支的在同一个sub-warp内。

上述例子为V100中官方例子，其中有两个分支A和B ，X和Y，warp按照串行执行两个分支。上图中有一个明显改进就是 执行A之外 可以切换到执行X，这样做的好处可以隐藏A内存操作延迟，这样交替执行更能提高硬件利用率。同时也可以发现两个分支的聚合点Z，并没有等待所有分支执行完毕之后再一起聚合执行，这是因为无法识别到Z是否依赖各自的结果，所以只能各自执行以便提高效率。如果有些算法需要聚合之后再同步执行Z，NV提供了另外一个函数__syncwarp()，可以是Z之后的代码进行聚合执行，如下提u所示：

 

NV已经将上述方法申请了专利 《Execution of divergent threads using a convergence barrier》最后可能和真正商用稍有不同，但是其基本思想基本一致，该专利中使用a convergence barrier方法来解决上述死锁问题，就是如果出现上述分支依赖情况，就需要创建一个a convergence barrier，来保证第一分支先执行，后一个分支后执行，而为了保存维护这些信息，每个线程都需要维护各自的entry信息，如下：

Barrier Participation Mask是用来跟踪维护哪些线程归属于哪个convergence barrier，因为在一个kernel中可能会存在多个分支，意味着会出现多个 convergence barrier， 线程会一直根据Participation Mask进行等待 直到都达到了common point(类似于reconvergence point ).

Barrier State用于跟踪哪些线程达到了convergence barrier状态。

Thread State用来记录线程状态，处于active状态还是处于block状态等。

Thread rPC: 线程要执行的下一个指令

Thread Active:记录线程是否处于激活状态。

上述各个entry大小于warp相关，如果warp为32个，则各自都有32个信息。

该专利中给出了明确执行流程，通过 convergence barrier产生一个额外的等待分支，知道 convergence barrier执行之后才进行下一个分支执行。

在调度开始时，首先确认在出现分支时的第一条指令，如果处于多数激活的线程执行是由少数或者特殊的线程来启动，则再调度单元中表明需要一个convergerence barrier。即第一步确认是否需要一个convergerence barrier。

如果需要一个convergerence barrier，则需要第一条执行路径需要执行少数或者特殊的那个线程。

少数或者特殊的那个线程执行第一个指令执行完毕之后，会触发执行第二个分支指令第一条指令。

然后执行特殊进程的第二个指令。

然后是否查看少数或者特殊的那个线程 是否达到了convergerence barrier？如果没有则执行额外的一条路径指令，特殊进程的第三条指令继续执行，不断重复，知道convergerence barrier 达到。

执行完成之后清除掉convergerence barrier。

WARP SCHEDULING

warp调度为GPU执行过程中非常关键的一部分，直接决定了每个时钟周期哪些线程在warp得到调度运行，不同的GPU其调度算法不同，各式各样。但是都有一个前提是在每个warp执行时，每个warp都只能同时执行相同指令，只有前一个指令执行完毕之后，warp才会通过调度执行下一个指令。

在一种理想GPU状态下，GPU内的每个warp内的线程访问内存延迟都相等，那么可以通过在warp不断切换线程可以隐藏内存访问的延迟，比如在同一个warp内，此时执行的内存读取指令，那么可以采用异步方式，在读取内存等待过程中，下一刻切换线程其他指令进行并行执行，这样GPU就可以一边进行读取内存指令，一边执行计算指令动作。该方法主要时因为GPU 将不同类型的指令分配给不同的单元进行执行，读取内存使用LD/ST硬件单元，而执行计算指令可能使用INT32或者FP32硬件单元。这样就可以通过循环调用（round robin）隐藏内存延迟问题。在理想状态下，可以完全通过这种循环调用方式完全隐藏掉内存延迟问题。

使用循环调用的调度方式时，就需要在每个时钟周期内不断切换线程，每个线程都需要有自己的专用寄存器保存私有相关信息，随着warp 切换数量不断增加，warp其所需要的寄存器会不断增多，同样会造成芯片面积不断增大。同样在固定面积大小的芯片上，随着warp数量增加，core的数量将不断减少，同样core数量减少将造成性能下降，反回来会不足够完成隐藏掉内存延迟，这本身时称为一个矛盾问题。

在实际上内存延迟问题还取决于 应用程序访问的内存位置以及每个线程对off-chip 内存的访问数量。

内存延迟问题 影响着 warp 调度，可以通过合理的warp 调度隐藏掉内存延迟问题。

TWO-LOOP APPROXIMATION

two-loop 即instruction issue loop，主要解决即将执行的指令，将需要执行的指令发送到相应的ALU中进行执行。

在上篇文章中最后warp 调度中提到为了解决内存操作延迟，通过在时钟周期内不断切换线程隐藏掉内存延迟。但是当时间执行的线程不够多或者内存延迟不够时，仅仅依靠warp调度的层次不能完全解决掉该问题，还要在其他方面同时进行解决。本节主要是说明在同一个warp内 在单个线程内通过调正发送到ALU将要执行的指令顺序，同样可以隐藏掉一部分内存延迟问题。

通过调整发送的指令思路比较简单：假如有两个指令 一个指令为内存读取指令，另外一个指令为加法指令。而前一个指令执行读取内存指令（读取指令可能会消耗几个或者几十个甚至上百个时钟周期），在数据读取完成之前，其实core什么事情都干不了，只能等待读取指令周期执行完毕之后，再往下执行。既然读取指令和加法指令使用的是不同的硬件单元，那么再第一个时钟周期执行内存读取指令之后，下一个时钟周期不必等待读取内存指令而是执行加法指令，从而实现一个一边计算一边读取并行的执行，从而提高整个运行效率。

然而在实际情况中，后一个指令是依赖于前一个指令的读取结果。为了解决该办法，就需要就GPU提前进行对指令之间的依赖关系进行预测，解析出指令之间的独立性，是否对其他指令有依赖关系。

为了解析要该指令集是否对上一个指令有依赖关系，GPU需要提前从内粗中读取一些指令，而不是只读取一个指令。为了存储这些指令，SIMT中增加了instruction buffer(I-buffer)，将这些指令存储起来。I-buffer在实际GPU设计的过程中一般采用的是L1 cache，以便提高读取效率。

For this purpose, GPUs implement an instruction buffer were instruction are placed affter caches access.

这些指令被从内存读取出来放入到I-buffer中，但是这样做还不够，还无法识别出这些指令的依赖。

那么如何解决这些指令的依赖性？可以借鉴CPU中的设计，在CPU中 为了识别指令集的依赖关系一般采用 scoreboard 和reservation stations两种方法，其中 reservation stations 为了识别出指令集之间的依赖关系，需要创建出一种associative logic关联逻辑关系，而创建该关系不仅仅会增加芯片面积还会增加芯片消耗，显然对于GPU来说该方法不合适。scoreboard 方法与reservation stations相比相对要简单许多，在CPU中为了解析指令之间的独立性，为每个寄存器都增加一个bit位，用于表示该寄存器是会被写，如果置1则表示该寄存器被写，此时如果另外一个指令中操作的源或者目的寄存器发现该寄存器bit位被置1，则会处于一直等待状态（说明该指令依赖与前一个指令），一直到该寄存器的bit位被清零（表明之前写寄存器操作完成）。这样就可以防止前面一个命令对该寄存器写之后，另外一个指令同时对该寄存器写造成出现数据不一致。同样也可以防止对该寄存器入write-after-read操作顺序（先write后读）变成read-after-write。

each register is represented in the scoreboard with a single bit that is set whenever an instruction issues that will write to that register.Any instruction that wants to read or write to a register that has its corresponding bit set in the scoreboard  is stalled until the bit s cleared by the instruction writing to the register.This prevent both read-after-writer and writer-after-writer hazards.when combined with in-order instruction issue this simple scoredboard can prevent write-after-read hazards provided that reading of the register file is constrained to occur in-order which is typically the case in in order CPU designs.

但是将scoreboard应用到GPU还存在一定两个问题需要解决：

• 由于GPU寄存器要比CPU多的多，如果为每个一个寄存器增加一位那么将需要更多的额外内存，比如假设一个GPU 每个core有64 个warps，每个warp有128个寄存器，那么为每个core增加8192bits大小内存。
• 在GPU一旦一个指令由于依赖性被堵塞，那么将会一直进行轮询查看scoreboard中的状态，直到前一个指令集执行完成将该bit位被清零。在GPU中，由于同时会有很多个线程在相同时刻执行相同指令，一旦其执行的指令被堵塞，那么将会有很多线程同时访问scoreboard，将会造成很大压力，同时i-buffer中有没有依赖的指令，也无法发出去。

为了解决上述问题，Coon等人在《Tracking Register Usage During Multithreaded Processing Using a Scorebard having Separate Memory Regions and Storing Sequential Register Size Indicators》提出了动态解决方案，在本文中提出为每个warp创建一个表，表中的每个entry记录被指令做写操作的寄存器。这样当一个指令集从内存中读取出来放入到I-buffer时将该指令集中的源寄存器和目的寄存器与entry做比较，是否有其他指令集已经对该寄存器再做写操作，如果有则返回一个bit vector与该寄存器一起写入到I-buffer中。如果该指令集的写操作完成，将会刷新I-buffer中的该指令集寄存器去的bit vector，将bit vector清除掉。另外如果一个指令集做写操作，需要将该寄存器放入的entry中，但是此entry已经满了，那么该指令集将会一直等待 或者被丢弃过一定时钟周期被重新获取再次查看entry是否满。

if all entries are used up for a given warp then either fetch stalls for all warps or the instruction is discarded and must be fetched again.When an instruction that has executed is ready to write to the register file it clears the emtry that was allocated to it in the scoreboard and also clears the corresponding dependency bit for any instruction from the same warp that are stored in the instruction buffer.

 The second loop selects an instruction in the instruction buffer that has no outstanding dependencies and issue it to execution units。

THREE-LOOP APPROXIMATION

通过再warp内不在每个时钟周期不断切换线程来隐藏内存访问延迟问题，但是切换线程就意味着需要大量register file来保存上下文信息，比如在NV(Kepler,MAXwell等）架构涨其register file达到了256KB，理想情况下每个时钟周期每个指令操作的寄存器都需要一个port，这样就需要大量port来解决访问问题，显然在实际中不太可能。为了减少port数量，Coon等人通过 mutiple banks of single-ported memories来解决问题，设计除了著名的operand collector结构。

 上图为operand collector原型图，该图中拥有4个bank 通过crossbar链接链接到 Pipleline Register，pipleline Register 为stage register，用于临时存储从regiester file读取过来的指令/数据，stage register最后将得到的数据/指令发送到SIMD执行单元中。

上图一个寄存器一个为4个bank的布局，其中r0表示的是r0寄存器， w0代表warp 0，那么w0:r0代表来自于w0的r0寄存器分布在bank0中，来自于w0的r1寄存器分布在bank1中，依次进行排布，当寄存器的数量大于bank时，那边将轮转重新排布，即w:r4排布到bank 0寄存器中。依次w1:r0重新从bank0开始排布。那边如果在一个warp0和wapr1中如果warp 0读取的是r0寄存器 warp1读取的是r1寄存器，分布到两个不同的bank内，即可以同时进去读取。如果是warp0和warp读取 r0寄存器那么将会产生一个bank冲突，造成无法并行，bank冲突会严重影响性能。下面有个例子将会说明bank冲突将会怎样影响性能调度

 假设有两条指令，其中i1为mad 指令（乘加指令）以及i2 为add(加指令）。i1指令使用到源寄存器位r5、r4、r6分别使用bank1、bank0、bank2，目的寄存器位r2，使用的是bank2。i2 指令使用到的源寄存器为r5和r1，都是用到的bank1，目的寄存器为r5。

假设有这样一个时钟执行周期顺序，第0个时钟周期w3执行i1指令，第1个时钟周期w0执行i2指令，第4个时钟周期w1执行i2指令。在开始执行第一个w3:i1指令时，首先需要从源数据中读取数据，由于r5,r4,r6分别位于不同的bank则可以同时进行读取这三个寄存器数据分别占用bank0，bank1,bank2。在下一个时钟周期切换到w0:i2中 由于r5和r1使用相同bank 1,故只能一次读取一个寄存器 ，首先读取r1寄存器。接着下个时钟周期，w3计算完毕需要将结果保存到r2中，需要占用bank1。同时执行w0：i2指令继续执行读取r5值。到第4个时钟周期 切换到指令w1::i2指令，同样首先读取i2的值。第5个时钟周期将w0:i2指令结果保存到r5，同时w1要读取r5寄存器的值指令由于bank被占用 无法读取，只能在下一个周期执行。

由上面执行可以看到由于r5和r1寄存器的bank冲突，只能使程序串行，无法完全并行。

现代GPU设计中，如何考虑避免bank冲突是需要着重考虑的一点。同样可以通过应用程序的调整合理避免bank冲突，可以参考《CUDA bank 及bank conflict》