GPGPU的流式多处理器微架构原理解析.docx

《GPGPU的流式多处理器微架构原理解析.docx》由会员分享，可在线阅读，更多相关《GPGPU的流式多处理器微架构原理解析.docx（13页珍藏版）》请在第一文库网上搜索。

1、GPGPU的流式多处理器微架构原理解析作者陈巍博士：存算一体/GPU架构和博专家,高级职称。中关村云计算产业联盟,中国光学工程学会专家,国际计算机学会（KM）会员，中国计算机学会（CCF）专业会员。作者耿云川博士：资深SOC设计专家，软硬件协同设计专家,擅长人工智能加速芯片设计。流式多处理器（StreamMu1ti-processor,SM）是构建整个GPU的核心模块（执行整个Kerne1Grid）,一个流式薮:理器上一般同时运行多个线程块。每个流式多处理器可以视为具有较小结构的里，支持指令并行（多发射）。流式多处理器是线程块的运行载体，但一般不支持乱序执行。每个流式多处理器上的单个Warp以

2、SIMD方式执行相同指令。图3-1流式多处理器在GPU架构中的位置（以NVID1A在SIa架构为例,修改自NVIDIA）3.1整体微架构图3-3是流式多处理器（SM,幽称之为计算单元）微架构（根据公开文献和专利信息综合获得）。流式多处理器按照流水线可以分为SIMT前端和SIMD后端。整个流水线处理划分为六个阶段，包括取指、译码、发射、操作数传送、执行与写回。图3-2GPGPU的流式多处理器结构划分S1MD即单指令多数据，采用一个控制器来控制多组计算单元(或处理器)，同时对一组数据(向量)中的每一个数据分别执行相同的操作从而实现空间并行性计算的技术。SIMT即单指令多线程，多个线程对不同的数据集

3、执行相同指令。S1MT的的优势在于无须把数据整理为合适的矢量长度，并且SIMT允许每个线程有不同的逻辑分支。按照软件级别，S1MT层面，流式多处理器由线程块组成，每个线程块由多个线程束组成；SIMD层面，每个线程束内部在同一时间执行相同指令，对应不同数据，由统一的线程束调度器(Warpschedu1er)调度。一般意义上的CUDA核，对应于流处理器(SP),以计算单元和分发端口为主组成。线程块调度程序将线程块分派给SIMT前端，线程在流式多处理器上以Warp为单位并行执行。图3-3GPGPU的流式多处理器微架构流式多处理器中的主要模块包括：取指单元(I-Fetch)：负责将指令请求发送到指令缓

4、存。并将程序计数器(Pe)指向下一条指令。指令缓存(I-Cache)：如来自取指单元的请求在指令缓存中被命中，则将指令传送给译码单元，否则把请求保存在未命中状态保持寄存器(MSHR)中。译码单元(Decode)：将指令解码并转发至IfUffer。该单元还将源和目标寄存器信息转发到记分牌，并将指令类型、目标地址(用于分支)和其他控制流相关信息转发到SIMT堆栈。SIMT堆栈(S1MTStack)：SIMT堆栈负责管理控制流相关的指令和提供下一程序计数器相关的信息。记分牌(Scoreboard)：用于支持指令级并行。并行执行多条独立指令时，由记分牌跟踪挂起的寄存器写入状态避免重复写入。指令缓冲(I

5、-BUffer):保存所有WarP中解码后的指令信息。Warp的循环调度策略决定了指令发射到执行和写回阶段的顺序。后端执行单元：后端执行单元包括CUDA核心（相当于A1U）、特殊功能函数、1D/ST单元、张量核心（TenSorcore）o特殊功能单元的数量通常比较少，计算相对复杂且执行速度较慢。（例如，正弦、余弦、倒数、平方根）。共享存储：除了寄存器文件，流式多处理器也有共享存储，用于保存线程块不同线程经常使用的公共数据，以减少对全局内存的访问频率。3.2取指与译码read_operand图3-4GPU执行流程（修改自GPGPU-Sim）取指-译码-执行，是处理器运行指令所遵循的一般周期性操作

6、。取指一般是指按照当前存储在程序计数器（PrOgra1nCounter,PC）中的存储地址，取出下一条指令，并存储到指令寄存器中的过程。在取指操作结束时，PC指向将在下一个周期读取的下一条指令。译码一般是指将存储在指令寄存器中的指令解释为传输给执行单元的一系列控制信号。图3-5取指译码结构在GPGPU中，译码之后要对指令进行调度，以保证后继执行单元的充分利用。这一调度通过线程束调度器（WarpSchedu1er）实现。线程束是为了提高效率打包的线程集合（NVIDIA称之为WarPs,AMD称为Wavefronts）o在每一个循环中的调度单位是Warp,同一个WarP内每个线程在同一时刻执行相同

7、命令。取指与译码操作过程如下：取指模块（1Fetch）根据PC指向的指令，从内存中获取到相应的指令块。需要注意的是，在GPGPU中，一般没有CPU中常见的乱序执行。SIMT堆栈指令缓存图3-5取指模块指令缓存(IYache)读取固定数量的字节(对齐)，并将指令位存储到寄存器中。对I-CaChe的请求会导致命中、未命中或保留失败(ReSerVatiOnfai1)o保留失败发生于未命中保持寄存器(MSHR)已满或指令缓存中没有可替换的区块。不管命中或者未命中，循环取指都会移向下一Warp。在命中的情况下，获取的指令被发送到译码阶段。在未命中的情况下，指令缓存将生成请求。当接收到未命中响应时，新的指

8、令块被加载到指令缓存中，然后WarP再次访问指令缓存。指令缓冲(!-Buffer)用于从IYache中获取指令后对译码后的指令进行缓冲。最近获取的指令被译码器译码并存储在I-Buffer中的相应条目中，等待发射。每个Warp都至少对应两个I-BUffe每个I-Buffer条目都有一个有效位(Va1id)就绪位(ReMy)和一个存于此Warp的已解码的指令。有效位表示在I-Buffer中的该已解码的指令还未发射，而就绪位则表示该WarP的已解码的指令已准备好发射到执行流水线。图3-4指令缓冲当WarP内的I-BUffer为空时，WarP以循环顺序访问指令缓存。（默认情况下，会获取两条连续的指令）

9、这时对应指令在Ifuffer中的有效位被激活，直到该Warp的所有提取的指令都被发送到执行流水线。当所有线程都已执行，且没有任何未完成的存储或对本地寄存器的挂起写入,则Warp完成执行且不再取指。当线程块中的所有WarP都执行完成且没有挂起的操作，标记线程块完成。所有线程块完成标记为内核己完成。相对于CPU,GPU的前端一般没有乱序发射，每个核心的尺寸就可以更小，算力更密集。3.3发射发射是指令就绪后，从指令缓冲进入到执行单元的过程。在（译码后的）指令发射阶段，指令循环仲裁选择一个WaP，将I-Buffer中的发射到流水线的后级，且每个周期可从同一Warp发射多条指令。所发射的有效指令应符合以

10、下条件:在Warp里未被设置为屏障等待状态；在I-Buffer中己被设置为有效指令（有效位被置为1）；已通过计分板（Scoreboard）检查；指令流水线的操作数访问阶段处于有效状态。在GPU中，不同的线程束的不同指令，经由S1MT堆栈和线程束调度，选择合适的就绪的指令发射。在发射阶段，存储相关指令（1oad、StOre等）被发送至存储流水线进行相关存储操作。其他指令被发送至后级SP（流处理器）进行相关计算。3.3.1SIMT堆栈SIMT堆栈用于在Warp前处理SIMT架构的分支分化的执行。一般采用后支配堆栈重收敛机制来减少分支分化对计算效率的负面影响。SIMT堆栈的条目代表不同的分化级别，每

11、个条目存储新分支的目标PC、后继的直接主要再收敛PC和分布到该分支的线程的活动掩码。在每个新的分化分支，一个新条目被推到栈顶；而当Warp到达其再收敛点时，栈顶条目则被弹出。每个Warp的SIMT堆栈在该Warp的每个指令发出后更新。线程束分化从功能角度来看，虽然SIMT架构下每个线程独立执行，但在实际的计算过程中会遇到一些分支的处理，即有些线程执行一个分支，而另外的线程则执行其他分支。如果在同一个WarP内不同的线程执行不同的分支，就会造成线程束分化，导致后继S1MD计算的效率降低。因此应尽量避免线程束的分化。ifCthreadIdx.xandStack(S)Iiiiiiiiiiiiiiii

12、iiiiiiiiiiiiiiii32threadwarpVo1taSSSSSSSSSSSSSSSSSSSSSSSSSSSSSSSS4kConvergenceOptimizerUUMUUOUOooukjuoowoououuuuuOMVUuodddddddddddddddddddddddddddddddd32threadwarpwithindependentschedu1ing图3-8V1OOWarP调度对比图2解决死锁的方法如下：NVIDIA为V1OO中Warp内的每个线程都分配了一个PC指针和堆栈，将PC指针的颗粒度细化到每一个线程中去，保障数据交换避免死锁。(图3-5)为避免细粒度的PC指针

13、和堆栈与GPU的SIMT执行模型产生冲突，硬件仍以Warp为单位来进行线程调度。使用了SChedU1eOptimizer(调度优化器)硬件模块来决定哪些线程可以在一个WarP内进行调度，将相同的指令重新进行组织排布到一个WarP内，并执行S1MD模型，以保证利用效率最大化2.3.3.2线程束调度与记分牌进行线程束(Warp)调度的目的是充分利用内存等待时间，选择合适的线程束来发射，提升执行单元计算效率。在理想的计算情况下，GPU内每个WarP内的线程访问内存延迟都相等，那么可以通过在WarP内不断切换线程来隐藏内存访问的延迟。GPU将不同类型的指令分配给不同的单元执行，1D/ST硬件单元用于读取内存,而执行计算指令可能使用INT32或者FP32硬件单元，且不同硬件单元的执行周期数一般不同。这样，在同一个WarP内，执行的内存读取指令可以采用异步执行的方式，即在读取内存等待期间，下一刻切换线程其他指令做并行执行，使得GPU可以一边进行读取内存指令，一边执行计算指令动作，通过循环调用(RoundRobin)隐藏内存延迟问题，提升计算效率。在理想状态下，可以通过这种循环调用方式完全隐藏掉内存延迟。但在实际计算流程中，内存延迟还取决于内核访问的内存位置，以及每个线程对内存的访问数量。内存延迟问题影响着Warp调度，需要通过合理的Warp调度来隐藏掉内存延迟问题。1）指令顺