xixi's blog

life is so beautiful!

0%

GPU SIMT Core

Posted on Edited on In ,
SIMT Core 简介

内容参考自：General-Purpose Graphics Processor Architecture

传统的图像渲染应用中，GPU 需要访问的数据集太大，很难在片上 cache 中被完全缓存，为提高性能，就需要足够大的片外带宽。
尽管每个线程的片上存储很小，但 cache 依旧可以通过空间局部性有效减少片外访存的数量(比如相邻像素之间的操作)

GPGPU Pipeline

gpgpu_pipeline

单个 SIMT Core 的前端是 SIMT ，后端是 SIMD。而且每条流水线包含3个调度循环：

  1. instruction fetch loop
    Fetch, I-Cache, Decode, I-Buffer
  2. instruction issue loop
    I-Buffer, Scoreboard, Issue, SIMT Stack
  3. register access scheduling loop
    Operand Collector, ALU, Memory

结构冒险的处理

  • 单线程的 in-order CPU pipeline 中，通常的解决方案是 stall 冒险及之后的指令
    GPU 不太可行：
    • 寄存器文件较大，以及在 GPU 所需的众多流水线阶段中，传播阻塞信号可能会对关键路径产生影响。流水线中的阻塞周期传播需要引入额外的缓冲，从而增加芯片面积
    • 其次，暂停一个 warp 中的指令可能导致其他 warp 中的指令被阻塞在其后面。如果这些指令并不需要导致阻塞的指令所占用的资源，那么吞吐量可能会受到影响
  • GPU 采用指令 replay 实现: 保持指令在 I-Buffer 中直到确认其执行完毕

SIMT 前端

  1. Fetch
    • 使用 warp 的 pc 访问 I-Cache
  2. Decode
    • 从 regfile 中取出源寄存器值
    • 获得 SIMT execution mask

I-Buffer

在 I-Cache 访问完成之后将指令放入 I-Buffer
I-Buffer 使用一个独立的调度器用于决定哪条指令应被下一个流水级执行

依赖检测：检测同一 warp 内指令之间的数据依赖性: in-order scoreboard

In-order Scoreboard


在 GPU 上存在的问题：

  1. 现代 GPU 中包含的寄存器数量非常庞大。每个 warp 最多有 128 个寄存器，每个核心最多包含 64 个 warp，因此每个核心的 scoreboard 需要 8192 位
  2. 遇到存在依赖关系的指令必须反复在 scoreboard 中查找其操作数，直到它所依赖的先前指令将结果写入寄存器文件
    • 在单线程设计中，不会引入太多复杂性
    • 在有序发射多线程处理器中，来自多个线程的指令可能都在等待早期指令完成。如果所有这些指令都必须查询 scoreboard ，则需要较多的读端口
      • 每个 warp 最多有 4 个操作数，若所有 warp 每周期都查询 scoreboard ，则需要 256 个读端口
      • 替代方法: 限制每周期可以查询 scoreboard 的 warp 数量
        也会限制可调度的 warp 数量
        如果所有检查的指令都存在依赖关系，即使其他未被检查的指令恰好没有依赖关系，也可能无法发出任何指令

解决策略: 为每个 warp 包含少量 (3 或 4 个) entry 而非为每个 warp 的每个寄存器设置一个 bit

  • 每个 entry 是寄存器号，该 entry 由一条已 issue 但尚未完成执行的指令写入
  • 这种 scoreboard 在当前指令被放入 I-Buffer 时以及指令将其结果写入寄存器文件时才被访问
    (常规的 in-order scoreboard 在存在指令 issue 和 writeback 时都会被访问)
  1. 一条指令从 I-Cache 中被取出并放置到 I-Buffer 时，会将该 warp 对应的 scoreboard entry 与该指令的源寄存器和目的寄存器进行比较，并将比较结果 (dependency bits) 记录到 I-Buffer
  2. 只有当 dependency bits 全部清零时指令调度器才会考虑调度该指令
  3. 当指令写回时会清零相应的 dependency bits

SIMT Execution Masking (SIMT Stack)

执行模型：在 GPGPU 中，各个线程完全独立执行

实现:

  1. 仅通过预测机制(传统 CPU)
  2. 通过传统预测机制与 SIMT Stack 的预测掩码堆栈相结合(当前 GPU)

SIMT Stack 有助于高效处理两个关键问题:

  1. 嵌套控制流: 在嵌套控制流中，一个分支的执行依赖于另一个分支
    • 传统支持预测的 CPU 通过使用多个预测寄存器来处理嵌套控制流
  2. 当一个 warp 中的所有线程都避开某个控制流路径时，可以完全跳过计算
    • 对于复杂的控制流来说，这可能带来显著的效率提升

示例程序基本块：(以每个 warp 4 个线程为例)
branch_program

问题：由于 GPU 后端是 SIMD 执行，所以需要所有线程同时进入后端。

处理方式：对同一 warp 中沿着不同路径执行的线程进行序列化处理
re-convergence

  1. Initial State: 最初，每个线程都在执行基本块 B 。
  2. Divergent Branch: 在分支之后，前三个线程继续执行基本块 B 中的代码。此时，第四个线程被屏蔽
  3. Reconvergence: 为了保持 SIMD 执行，第四个线程在几个周期之后被调度执行基本块 F ，然后和前 3 个线程聚合一起执行基本块 G

序列化的实现: 使用 SIMT Stack
SIMT Stack 示例

SIMT Stack 在专利和指令集手册中存在多种实现方式（但都至少部分由特殊指令进行管理）

缺点：死锁

基于 SIMT Stack 的 SIMT 实现可能导致 “SIMT Deadlock”

1
2
3
4
A: *mutex = 0
B: while(!atomicCAS(mutex, 0, 1));
C: // critical section
   atomicExch(mutex, 0, 1);

atomicCAS: 对包含 mutex 的内存位置执行 compare-and-swap
编译器内建函数，会被转换为 atom.global.cas PTX 指令

  1. 所有线程都访问 mutex，只有一条线程可以离开 while 循环，而其他线程则留在循环中。
  2. 退出循环的线程将达到 Reconvergence Point ，因此将不存在于 SIMT Stack 中，而其他线程仍停在循环中未到达 Reconvergence Point, 从而无法执行 C 的 atomicExch 释放锁

Stackless SIMT Architectures: Independent Thread Scheduling

为避免 SIMT Stack 带来的死锁问题, Independent Thread Scheduling 被提出并采用。
(NVIDIA Volta GPU 架构)

Key Idea: 用每个 warp 的 convergence barrier 替代 SIMT Stack

新增硬件寄存器(由硬件 warp 调度器使用, 由软件管理，因为可能嵌套很深)：

  • Barrier Participation Mask: 跟踪给定 warp 中哪些线程参与了特定的 convergence barrier
    • 其位宽同 warp 中包含的线程数, 如果某一位被设置，则表示该 warp 中的对应线程参与此次收敛屏障
    • 对于一个给定的 warp ，可能有多个 Barrier Participation Mask (以支持嵌套控制流结构)
    • 在通常情况下，由某个 Barrier Participation Mask 跟踪的线程会在经历一个 divergent branch 后，等待其他线程到达程序中的共同点(WAIT 指令设定)，从而重新汇聚在一起 (功能同 SIMT Execution Mask)
  • Barrier State: 跟踪哪些线程已经到达了特定的 convergence barrier
  • Thread State: 记录 warp 中的每个线程的状态
    • ready : 是否准备好执行
    • block : 是否阻塞于 convergence barrier (如果是，阻塞到哪一个 barrier)
    • yielded: 是否已让出执行权。(用于在 SIMT 死锁的时候使能其他线程通过 barrier)
  • Thread Active: 表明该 warp 中的当前线程是否被屏蔽
  • Thread rPC: 对于 warp 中未被屏蔽的线程，记录 next PC

实现机制：

  1. Initial: 采用特殊的 ADD 指令
    当 warp 执行该 ADD 指令时，所有 active 的线程在由 ADD 指令指示的 Barrier Participation Mask 中被置位
  2. Divergent Branch: 调度器会选择一组具有相同 PC 的线程，并更新 Thread Active 以允许这些线程执行该 warp (这些线程子集称为“warp split”)
  3. Reconvergence: 采用特殊的 WAIT 指令
    • WAIT 指令用于在一个 warp split 达到 convergence barrier 时使其停止
      将 warp split 中的线程添加到该 barrier 的 Barrier State 寄存器中，并将线程的状态改为 blocked
    • WAIT 指令包含一个操作数，标识哪个 convergence barrier
    • 一旦所有参与该 barrier 的线程都执行了相应的 WAIT 指令，调度器就可以将原 warp 的所有线程切换为 active 继续执行(保持 SIMD 效率)

与基于栈的 SIMT 实现不同，使用 convergence barrier 实现时，调度器可以在分支发散后的线程组之间自由切换，使得当某些线程已获取锁而其他线程尚未获取时，能够实现 warp 内线程可以走出 barrier

基于栈的实现：

1
2
3
4
5
6
7
8
if (threadIdx.x < 4) { 
  A; 
  B; 
} else { 
  X; 
  Y; 
}  
Z;

基于栈实现 Reconvergence

基于 convergence barrier 的实现：

1
2
3
4
5
6
7
8
9
10
if (threadIdx.x < 4) { 
  A; 
  B; 
} else { 
  X; 
  Y; 
}  
Z;
__syncwarp();

Volta GPU 实现 Reconvergence

Issue

warp 调度策略：

  1. 轮询调度 (round robin)
    • 能让每个发出的指令大致获得相等的时间来完成执行
    • 局部性特性可能有利于或阻碍轮询调度
      • 有利于：当不同线程在其执行过程中于相似的阶段共享数据时，让线程均匀地推进是有益的(可以增加 cache hit)
      • 有利于：当地址空间中相邻的位置在时间上也相邻被访问时，访问 DRAM 会更高效
      • 阻碍：当线程主要访问互不重叠的数据时，让某个线程反复被调度以最大化局部性可能是有益的

trade-off:

  • 如果核心中 warp 的数量乘以每个 warp 的发射时间超过内存延迟，那么核心中的执行单元将始终处于忙碌状态。从理论上讲，增加线程块数量直到这一临界点可以提升每个核心的吞吐量。
  • 每个核心中增加 warp 的数量会提高芯片面积中用于寄存器文件存储的比例
    • 为了使不同的warp在每个周期内都能发出指令，必须让每个线程拥有自己的寄存器（这避免了在寄存器和内存之间复制和恢复寄存器状态的需要）
    • 对于固定的芯片面积，增加每个核心的warp数量将减少每片芯片中的核心总数

SIMD 后端

Register File

GPU 每个核心需要支持大量 warps 。为了实现 warps 之间的逐周期切换，需要一个很大的寄存器文件包含每个正在执行的 warps 的单独的物理寄存器

  1. 容量大: 在 NVIDIA Kepler, Maxwell 和 Pascal 等架构中，寄存器文件容量可达 256 KB
  2. 端口多

减少寄存器文件面积的方式：分 Bank 设计, 通过多个单端口存储器的分组来模拟大量端口
两种实现方式：

  1. 将这些 Bank 暴露给 ISA
  2. 采用 Operand Collector (更透明)

Banked Register

  • RR: Register Read Stage
  • Register File 由 4 个(可能更多)单端口逻辑 bank 组成
  • 每个逻辑 bank 可能会进一步分解为更多的物理 bank
  • RR/EX 流水线寄存器用于缓冲源操作数，然后将其传递到 SIMD 执行单元
  • 仲裁器控制对各个 Bank 的访问，并通过 CrossBar 将结果路由到合适的 RR/EX 缓冲寄存器

Operand Collector

Operand Collector 微架构：
Operand Collector

主要架构：将 RR/EX 流水段寄存器替换为 Collector Units

  • 当指令进入读寄存器(RR)阶段时，每个指令都会被分配一个 Collector Unit
  • 由于存在多个 Collector Units，多个指令可以重叠读取源操作数，这在源操作数之间存在 Bank 冲突的情况下有助于提高吞吐量
  • 每个 Collector Units 包含执行指令所需的所有源操作数的缓冲空间
  • 多个指令的源操作数数量较多，仲裁器更有可能实现更高的 Bank 级并行，从而允许同时访问多个 Register File Bank

可能的问题： WAR 冒险
发生在同一 warp 的 2 条指令之间，第一条读，第二条写同一寄存器。当第一条指令遭遇重复的 Bank 冲突时，第二条指令可能在第一条指令读入旧值之前写入一个新值

防止 WAR 冒险的方法： 要求来自同一 warp 的指令按照程序顺序离开 Operand Collector 并进入执行单元

  1. 方案一: release-on-commit warpboard: 允许每个 warp 每次仅有一个指令在执行(会损失性能)
  2. 方案二: release-on-read warpboard: 允许每个 warp 每次仅有一个指令在 Operand Collector 中收集操作数(性能损失降低)
  3. 方案三: bloomboard: 使用一个小的 bloom filter 跟踪待处理的寄存器读取操作(性能损失最小)

NVIDIA Maxwell GPU 采用 read dependency barrier 控制指令避免特定指令的 RAW 冒险

Function Units

  • 每个线程在与某条 lane 相关联的功能单元上执行( SIMT 执行掩码已被设置)
  • 功能单元通常是异构的，即每个功能单元仅支持部分指令
    例如，NVIDIA GPU 包含 SFU, LSU, FPU, ALU，Tensor Core (Volta 架构引入)