第二十六课：如何从零设计一个简化的 SIMT GPU

导语：这一课从 RTL 设计角度搭建一颗最小可运行的 SIMT GPU。重点是模块边界、状态机和实现顺序，而不是一开始复制完整商用 GPU。

前面已经分别学习了：

• Thread、Warp、Block；
• Warp Scheduler；
• Active Mask；
• Scoreboard；
• Register File；
• Shared Memory；
• Load/Store Unit；
• Block 调度；
• 分支发散；
• Barrier。

这一课把这些模块组合成一颗可以用 RTL 实现的简化 GPU。

目标不是复制 NVIDIA GPU，而是实现一个能够运行简单 Kernel 的教学型 SIMT 处理器。

---

一、先确定设计目标

第一版 GPU 不应追求：

• 32 Lane；
• 多个 SM；
• Cache；
• Tensor Core；
• 虚拟内存；
• 抢占；
• 完整 CUDA；
• 复杂乱序执行。

可以先设计一个小型 GPU：

1 个 SM
4 个 SIMD Lane
每个 Warp 4 个线程
最多驻留 4 个 Warp
每个线程 16 个通用寄存器
每个 Block 最多 16 个线程
1～4 KB Shared Memory
单发射、顺序执行
简单整数与 FP32 运算

基本参数示例：

parameter int LANES          = 4;
parameter int WARPS_PER_SM   = 4;
parameter int REGS_PER_THREAD = 16;
parameter int MAX_THREADS_BLOCK = 16;
parameter int SHARED_BYTES   = 4096;

这样可以先验证最重要的 SIMT 机制，同时控制 RTL 规模。

---

二、最小 GPU 应支持什么程序

第一阶段只要求运行以下 Kernel 类型：

C[i] = A[i] + B[i];

if (A[i] > 0)
    B[i] = A[i];
else
    B[i] = 0;

shared[threadIdx.x] = A[i];
__syncthreads();
B[i] = shared[blockDim.x - 1 - threadIdx.x];

也就是支持：

• 线程索引；
• 整数地址计算；
• Global Memory Load/Store；
• 基本算术；
• 条件分支；
• Shared Memory；
• Block Barrier；
• Kernel 结束。

暂时不需要支持：

• 动态并行；
• 原子操作；
• 多 Block 全局同步；
• Texture；
• Cache；
• 异常与缺页；
• 多进程上下文。

---

三、总体架构

最小 GPU 可以划分为：

Host / Testbench
       │
       ▼
Kernel Launch Interface
       │
       ▼
Block Dispatcher
       │
       ▼
┌─────────────────────────────┐
│             SM              │
│                             │
│  Warp Context Table         │
│  Warp Scheduler             │
│  Instruction Fetch/Decode   │
│  Scoreboard                 │
│  Vector Register File       │
│  SIMD ALU                   │
│  Load/Store Unit            │
│  Shared Memory              │
│  Barrier Unit               │
│  Divergence Stack           │
└─────────────────────────────┘
       │
       ▼
Global Memory Interface

RTL 顶层可以组织为：

simt_gpu_top
├── kernel_launch_unit
├── block_dispatcher
├── sm
│   ├── warp_context_table
│   ├── warp_scheduler
│   ├── instruction_fetch
│   ├── instruction_decoder
│   ├── scoreboard
│   ├── vector_register_file
│   ├── simd_alu
│   ├── load_store_unit
│   ├── shared_memory
│   ├── barrier_unit
│   └── divergence_unit
└── global_memory_model

---

四、最小指令集

为了实现简单，可以设计一个自定义32位 ISA。

例如：

指令	功能
MOV	寄存器复制
ADD	整数或浮点加法
SUB	减法
MUL	乘法
MAD	乘加
AND/OR/XOR	位运算
SHL/SHR	移位
SETP	比较并生成谓词
BRA	条件跳转
LDG	Global Memory Load
STG	Global Memory Store
LDS	Shared Memory Load
STS	Shared Memory Store
BAR	Block Barrier
EXIT	线程或 Warp 结束

还需要访问特殊寄存器：

threadIdx.x
blockIdx.x
blockDim.x
gridDim.x

可以设计：

S2R r0, THREAD_ID_X
S2R r1, BLOCK_ID_X
S2R r2, BLOCK_DIM_X

其中 S2R 表示 Special Register to General Register。

---

五、一条 SIMT 指令如何执行

假设当前 Warp 有4个线程，执行：

ADD r3, r1, r2

它并不是只做一次加法，而是：

Lane 0：r3[0] = r1[0] + r2[0]
Lane 1：r3[1] = r1[1] + r2[1]
Lane 2：r3[2] = r1[2] + r2[2]
Lane 3：r3[3] = r1[3] + r2[3]

但只有 Active Mask 为1的 Lane 真正更新结果。

例如：

Active Mask = 1101

则：

Lane 0：执行
Lane 1：执行
Lane 2：屏蔽
Lane 3：执行

因此每条 Warp 指令都需要携带：

warp_id
PC
active_mask
opcode
source registers
destination register

---

六、线程状态与 Warp 状态

普通 CPU 为每个线程维护独立 PC。

传统 SIMT GPU 更常以 Warp 为调度单位，一个 Warp 共享：

• Program Counter；
• 指令流；
• Active Mask；
• Barrier 状态；
• Divergence Stack。

但每个 Lane 有独立：

• 通用寄存器值；
• Predicate；
• Thread ID；
• Load/Store 地址；
• 运算数据。

可以表示为：

Warp Context
├── valid
├── state
├── pc
├── active_mask
├── block_id
├── warp_in_block
├── barrier_id
└── reconvergence_stack

线程级状态主要位于：

Vector Register File
Predicate Register File

---

七、Warp Context Table

假设一个 SM 最多驻留4个 Warp，就维护4组上下文：

Warp 0 Context
Warp 1 Context
Warp 2 Context
Warp 3 Context

每个 Warp 可以处于以下状态：

typedef enum logic [2:0] {
    WARP_FREE,
    WARP_READY,
    WARP_WAIT_SCOREBOARD,
    WARP_WAIT_MEMORY,
    WARP_WAIT_BARRIER,
    WARP_FINISHED
} warp_state_t;

但在实际设计中，WAIT_SCOREBOARD 不一定必须存成固定状态，因为 Warp 可能每周期重新检查 Ready 条件。

更常见的方式是维护多个阻塞原因：

valid
finished
waiting_memory
waiting_barrier
instruction_ready
operand_ready

调度器综合判断 Warp 是否 Eligible。

---

八、Block Dispatcher

Block Dispatcher 负责把 Grid 中尚未执行的 Block 分配给 SM。

第一版只有一个 SM，因此分配过程较简单。

Kernel 启动描述符可以包含：

kernel_pc
grid_dim_x
block_dim_x
global_memory_base
shared_memory_bytes
参数地址

Block Dispatcher 需要完成：

1. 选择一个未执行 Block；
2. 计算该 Block 需要多少 Warp；
3. 检查是否有足够 Warp Slot；
4. 检查 Shared Memory 是否足够；
5. 初始化 Warp Context；
6. 设置线程特殊寄存器；
7. 设置初始 Active Mask；
8. 将 Warp 标记为 Ready。

---

九、Block 如何拆成 Warp

假设：

Warp Size = 4
Block Size = 10

则需要：

$$\left\lceil\frac{10}{4}\right\rceil=3$$

个 Warp。

映射为：

Warp 0：Thread 0、1、2、3
Warp 1：Thread 4、5、6、7
Warp 2：Thread 8、9、无效、无效

最后一个 Warp 的 Active Mask 为：

1100

或者根据 Lane 位编号定义为：

0011

关键是两个无效 Lane 从一开始就被屏蔽。

线程索引可以计算为：

$$threadIdx.x= warp\_in\_block\times WARP\_SIZE+lane\_id$$

全局线程索引为：

$$global\_id= blockIdx.x\times blockDim.x+threadIdx.x$$

---

十、特殊寄存器如何实现

特殊寄存器不一定真正存入普通 Register File。

例如执行：

S2R r0, THREAD_ID_X

硬件可以根据：

warp_in_block
lane_id

实时生成每个 Lane 的线程编号。

类似地：

BLOCK_ID_X
BLOCK_DIM_X
GRID_DIM_X
LANE_ID
WARP_ID

可以由控制寄存器或组合逻辑生成。

这样可以减少初始化时的大量寄存器写入。

---

十一、Vector Register File

每个线程都需要独立寄存器。

如果：

4 Warps
× 4 Lanes
× 16 Registers
× 32 bit

总容量为：

$$4\times4\times16\times32 =8192\text{ bit} =1\text{ KB}$$

地址可以表示为：

warp_id
lane_id
register_id

逻辑结构：

register_file[warp][lane][reg]

---

十二、Register File 如何组织更容易实现

不要一开始实现一个具有几十个读写端口的大型统一 Register File。

教学设计可以按 Lane 分 Bank：

Lane 0 Register Bank
Lane 1 Register Bank
Lane 2 Register Bank
Lane 3 Register Bank

每个 Lane Bank 存储所有 Warp 在该 Lane 上的寄存器：

reg_bank[lane][warp][reg]

当 Warp 执行一条双源指令时，每个 Bank 需要：

• 两个读端口；
• 一个写端口。

所有 Lane Bank 并行工作。

这种结构天然适配 SIMD：

同一个 reg_id
在不同 Lane Bank 中同时读取

---

十三、Predicate Register

条件比较可以生成每个 Lane 独立的谓词：

SETP.GT p0, r1, 0

结果可能是：

Lane 0：p0 = 1
Lane 1：p0 = 0
Lane 2：p0 = 1
Lane 3：p0 = 0

得到：

Predicate Mask = 1010

后续指令可以带谓词：

@p0 ADD r2, r2, 1

真实执行 Mask 为：

$$execute\_mask = active\_mask \land predicate\_mask$$

这样短条件可以不使用真正分支。

---

十四、Warp Scheduler

每个周期，Warp Scheduler 从多个 Warp 中选择一个发射。

候选 Warp 必须满足：

Warp 有效
尚未结束
不在等待内存
不在等待 Barrier
当前指令已取回
源操作数 Ready
目标执行单元可接收
Active Mask 非空

这样的 Warp 称为：

Eligible Warp。

---

十五、最简单的调度策略

第一版可以使用 Round-Robin：

周期 0：Warp 0
周期 1：Warp 1
周期 2：Warp 2
周期 3：Warp 3
周期 4：从 Warp 0 重新开始

但如果某个 Warp 不可发射，就跳过。

伪代码：

for offset = 0 to WARPS-1:
    candidate = round_robin_ptr + offset

    if warp[candidate].eligible:
        issue candidate
        round_robin_ptr = candidate + 1
        break

Round-Robin 简单、公平，适合第一版 RTL。

以后可以加入：

• Greedy-Then-Oldest；
• 优先继续同一个 Warp；
• 优先 Ready 时间最久的 Warp；
• 双 Warp Scheduler。

---

十六、Instruction Fetch

每个 Ready Warp 都有自己的 PC。

当调度器需要该 Warp 的指令时：

instruction = instruction_memory[warp_pc]

最简单设计可以：

• 使用单端口指令存储器；
• 每周期只取一条指令；
• 取指与发射串行；
• 不实现 Instruction Cache。

稍微改进后，可以给每个 Warp 配置一条 Instruction Buffer：

Warp PC
  ↓
Fetch
  ↓
Warp Instruction Buffer
  ↓
等待 Scheduler 发射

这样调度器不会因为当前周期取不到指令而停止。

---

十七、指令译码

Decoder 输出：

execution_unit
source_reg_0
source_reg_1
destination_reg
immediate
predicate
memory_space
branch_type
write_enable

例如：

ADD  → SIMD ALU
LDG  → Load/Store Unit
STS  → Shared Memory Unit
BAR  → Barrier Unit
BRA  → Branch Unit
EXIT → Warp Control

控制结构应明确区分：

• 指令是否写寄存器；
• 写回何时完成；
• 是否需要 Scoreboard；
• 是否改变 PC；
• 是否可能长期阻塞 Warp。

---

十八、Scoreboard

Scoreboard 用于记录寄存器结果是否尚未完成。

例如：

LDG r4, [r1]
ADD r5, r4, r2

LDG 发射后，数据可能很多周期后才能返回。

在返回前：

r4 = Busy

所以 ADD 不能发射。

---

十九、Scoreboard 的最小实现

可以为每个 Warp 的每个寄存器维护一个 Busy Bit：

busy[warp_id][reg_id]

发射会写目标寄存器的长延迟指令时：

busy[warp][dst] = 1

结果写回时：

busy[warp][dst] = 0

判断源操作数 Ready：

src0_ready = !busy[warp][src0]
src1_ready = !busy[warp][src1]

如果任一源寄存器 Busy，该 Warp 当前指令不能发射。

---

二十、所有指令都需要 Scoreboard 吗

可以分为两类。

固定短延迟指令

例如单周期整数加法。

如果保证写回在下一条同 Warp 指令读取前完成，可以简化处理。

长延迟指令

例如：

• Global Load；
• 乘法器；
• 浮点运算；
• Shared Memory 冲突访问。

这些指令更需要 Scoreboard。

为了设计统一，第一版也可以对所有写寄存器的指令设置 Busy，再在写回时清除。

这样控制更简单，但可能引入不必要的等待。

---

二十一、SIMD ALU

4-Lane GPU 可以使用4套相同 ALU：

Lane 0 ALU
Lane 1 ALU
Lane 2 ALU
Lane 3 ALU

每套执行：

ADD
SUB
AND
OR
XOR
SHIFT
COMPARE

输入为：

src0[lane]
src1[lane]

输出为：

result[lane]

写回使能为：

active_mask[lane]
&
predicate_mask[lane]

被屏蔽 Lane 不更新目标寄存器。

---

二十二、整数与浮点单元如何安排

第一版建议先只实现整数。

因为整数单元可以完成：

• 地址计算；
• 循环；
• 分支；
• 索引；
• 简单测试程序。

之后再加入：

FP32 Add
FP32 Multiply
FP32 FMA

可以选择：

方案一：每个 Lane 一个浮点单元

吞吐率高，但面积较大。

方案二：多个 Lane 共享一个浮点流水线

面积小，但需要分时执行和更多调度。

教学 GPU 更适合方案一或直接使用 FPGA 浮点 IP。

---

二十三、执行流水线

一个最小顺序流水线可以是：

Fetch
  ↓
Decode
  ↓
Read Operands
  ↓
Issue / Execute
  ↓
Writeback

但由于不同 Warp 交错执行，每一级可能处理不同 Warp：

周期 0：Fetch Warp 0
周期 1：Decode Warp 0，Fetch Warp 1
周期 2：Execute Warp 0，Decode Warp 1，Fetch Warp 2

这需要携带：

warp_id
pc
active_mask
destination register

穿过流水线。

---

二十四、Warp 切换是否需要保存寄存器

不需要像 CPU 操作系统上下文切换那样保存到内存。

所有驻留 Warp 的状态都常驻在：

• Warp Context Table；
• Register File；
• Scoreboard；
• Divergence Stack。

从 Warp 0 切换到 Warp 1，只需要 Scheduler 选择不同的 warp_id。

因此 Warp 切换成本可以很低。

这正是 GPU 隐藏延迟的基础。

---

二十五、Load/Store Unit

Load/Store Unit 负责：

1. 读取每个 Lane 的地址基址；
2. 加上 Offset；
3. 应用 Active Mask；
4. 判断地址属于哪个空间；
5. 合并内存请求；
6. 发送到 Global 或 Shared Memory；
7. 等待返回；
8. 写回对应 Lane 的寄存器。

对于：

LDG r4, [r1 + offset]

每个 Lane 都会产生一个地址：

addr[0]
addr[1]
addr[2]
addr[3]

---

二十六、第一版是否需要实现内存合并

可以分阶段。

第一阶段：每个 Lane 独立请求

4 个 Lane
→ 最多 4 次内存访问

实现最简单，但性能较差。

第二阶段：合并连续地址

如果4个 Lane 访问：

0x1000
0x1004
0x1008
0x100C

可以合成一次16字节事务。

第三阶段：按 Cache Line 或 Memory Segment 合并

根据地址高位分组：

相同 Segment 的 Lane
合并为一个请求

教学设计可以先完成第一阶段，确保功能正确，再加入 Coalescer。

---

二十七、内存请求需要保存什么信息

Global Load 发出后，需要记住返回数据属于哪个 Warp、哪个寄存器、哪些 Lane。

Load Queue Entry 可以包含：

valid
warp_id
destination_reg
lane_mask
lane_addresses
transaction_id

返回时：

根据 transaction_id 找到请求
        ↓
把各 Lane 数据写入 Register File
        ↓
清除 Scoreboard Busy Bit
        ↓
Warp 重新变为 Eligible

如果只允许每个 Warp 同时存在一个未完成 Load，设计会简单很多。

---

二十八、第一版如何简化 Global Memory

可以使用一个固定延迟的 SRAM 模型：

请求发出
  ↓
等待 5～10 周期
  ↓
返回数据

暂时不实现：

• Cache；
• DRAM Bank；
• TLB；
• 虚拟地址；
• 请求重排序。

这样已经可以观察：

Warp 0 等待 Load
Scheduler 改发 Warp 1

验证延迟隐藏机制。

之后再替换为 AXI 接口或更真实的内存系统。

---

二十九、Shared Memory

Shared Memory 是 Block 内共享的低延迟 SRAM。

地址通常是：

shared_base_of_block
+
thread_calculated_offset

一个 SM 可以有一块物理 Shared Memory，由不同 Block 分配不同区域。

第一版只有一个驻留 Block 时，可以直接使用：

shared_mem[address]

以后支持多个 Block 时，需要加入：

block_shared_base

---

三十、Shared Memory Bank

可以把 Shared Memory 划分为4个 Bank，对应4个 Lane：

$$bank=\left(\frac{address}{4}\right)\bmod4$$

如果4个 Lane 访问：

Lane 0 → Bank 0
Lane 1 → Bank 1
Lane 2 → Bank 2
Lane 3 → Bank 3

可在一个周期完成。

如果多个 Lane 访问同一 Bank 的不同地址，则发生冲突，需要分多周期处理。

---

三十一、Bank Conflict 控制器

Shared Memory Unit 可以先计算：

bank_id[lane]

然后检测冲突。

最简单处理方式是：

每周期每个 Bank 只服务一个 Lane
未完成 Lane 保存在 Pending Mask 中
下一周期继续服务

例如：

4 个 Lane 都访问 Bank 0

则需要4个周期。

执行期间：

• 对应 Warp 标记为等待；
• 其他 Warp 可以继续执行；
• 所有 Lane 完成后统一写回；
• 清除目标寄存器 Busy。

这样可以真实体现 Bank Conflict 的性能影响。

---

三十二、Barrier Unit

__syncthreads() 的硬件语义是：

当前 Block 中所有尚未退出的线程都到达 Barrier 后，才能继续执行。

因为 Warp 是调度单位，可以按 Warp 记录到达情况。

对于一个 Block 有3个 Warp：

Barrier Arrived Mask = 000

Warp 0 到达：

001

Warp 1 到达：

011

Warp 2 到达：

111

全部到达后：

• 清除 Barrier 状态；
• 所有等待 Warp 恢复 Ready；
• PC 继续向后执行。

---

三十三、Barrier 需要知道哪些 Warp 属于同一个 Block

Block Context 可以记录：

block_valid
block_id
resident_warp_mask
finished_warp_mask
barrier_arrived_mask
shared_memory_base

Barrier 完成条件不是固定要求所有 Warp，而是：

$$arrived\_mask = resident\_warp\_mask \land \neg finished\_warp\_mask$$

已经退出的 Warp 不应继续阻止 Barrier。

不过第一版可以限制：

所有线程必须在同一路径上到达 Barrier，且不能提前退出。

这样实现更简单。

---

三十四、错误 Barrier 程序会怎样

例如：

if (threadIdx.x < 2) {
    __syncthreads();
}

只有部分线程到达 Barrier。

硬件会永远等待其他 Warp 或线程。

真实 GPU 中这属于未定义行为。

教学 GPU 可以加入超时或断言：

assert (!(barrier_wait_too_long));

但不能期望硬件自动修复错误同步程序。

---

三十五、分支发散

假设执行：

if (threadIdx.x < 2)
    x = 1;
else
    x = 2;

对于4-Lane Warp：

Lane 0、1：条件为真
Lane 2、3：条件为假

产生两个 Mask：

True Mask  = 0011
False Mask = 1100

硬件需要分别执行两条路径。

---

三十六、第一版如何处理分支

可以分三个阶段实现。

阶段一：只支持 Warp-Uniform Branch

要求同一 Warp 所有 Active Lane 的条件相同。

如果条件不同，报告错误。

这是最简单的第一版。

阶段二：使用谓词替代短分支

编译器把简单 if 转换为：

@p  指令
@!p 指令

无需 Divergence Stack。

阶段三：实现真正分支发散

维护 Active Mask 和 Reconvergence Stack，分别执行两个路径。

推荐按这个顺序开发。

---

三十七、最小 Divergence Stack

每个 Warp 可以维护一个栈，栈项包含：

resume_pc
resume_mask
reconverge_pc

遇到发散分支时：

1. 计算 True Mask；
2. 计算 False Mask；
3. 选择一条路径先执行；
4. 把另一条路径压入栈；
5. 当前 Active Mask 设置为第一条路径；
6. 到达汇合点后弹出另一条路径；
7. 两条路径都完成后恢复合并后的 Mask。

简化示例：

原 Active Mask：1111
True Mask：     0011
False Mask：    1100

先执行 True 路径：

Active Mask = 0011

栈中保存：

False PC
Mask = 1100
Reconverge PC

True 路径结束后，再执行 False 路径。

---

三十八、如何找到 Reconvergence Point

真实编译器通常根据控制流图计算分支的 Immediate Post-Dominator。

教学 ISA 可以简化为：

BRA.DIV predicate, true_target, reconverge_pc

由编译器直接把汇合地址编码到指令中。

例如：

BRA.DIV p0, TRUE_LABEL, END_IF

这样硬件不用自己分析控制流图。

这是软件与硬件协同简化设计的典型方式。

---

三十九、SIMT 分支执行示例

程序：

if p0:
    ADD r1, r1, 1
else:
    ADD r1, r1, 2
MUL r2, r1, 4

假设：

p0 Mask = 0101

执行过程：

1. 识别发散
2. Active Mask = 0101
3. 执行 ADD +1
4. Active Mask = 1010
5. 执行 ADD +2
6. 汇合，Active Mask = 1111
7. 所有 Lane 执行 MUL

两条分支路径被串行执行。

---

四十、EXIT 如何处理

如果所有 Lane 同时执行：

EXIT

则 Warp 结束。

如果部分 Lane 在条件下退出：

@p EXIT

则需要：

active_mask =
active_mask & ~predicate_mask

当：

active_mask == 0

Warp 才真正完成。

Block 中所有 Warp 都完成后：

• 释放 Warp Slot；
• 释放 Shared Memory；
• Block 标记完成；
• Dispatcher 可以分配下一个 Block。

---

四十一、Kernel 完成条件

Grid 中所有 Block 完成后：

completed_blocks == grid_dim_x

Kernel Launch Unit 产生：

kernel_done = 1

Host 或 Testbench 可以等待该信号。

第一版可以只允许：

一次只运行一个 Kernel

以后再加入命令队列和多个 Kernel。

---

四十二、一个 Vector Add Kernel 的执行过程

CUDA 风格代码：

int i = blockIdx.x * blockDim.x
      + threadIdx.x;

if (i < n)
    C[i] = A[i] + B[i];

可以编译为概念指令：

S2R   r0, BLOCK_ID_X
S2R   r1, BLOCK_DIM_X
S2R   r2, THREAD_ID_X

MUL   r3, r0, r1
ADD   r3, r3, r2

SETP.LT p0, r3, n

MUL   r4, r3, 4

@p0 LDG r5, [A + r4]
@p0 LDG r6, [B + r4]

@p0 ADD r7, r5, r6

@p0 STG [C + r4], r7

EXIT

---

四十三、Vector Add 的周期级行为

假设 Warp 0 执行到：

LDG r5, [A + r4]

Load 发出后：

busy[warp0][r5] = 1
warp0 waiting for memory

调度器切换到 Warp 1：

Warp 1 执行地址计算

然后 Warp 2、Warp 3。

若 Warp 0 数据返回：

写回 r5
清除 busy
Warp 0 恢复 Eligible

随后执行第二个 Load。

这就是 GPU 的延迟隐藏：

Warp 0 等待内存
期间执行 Warp 1、2、3

---

四十四、为什么第一版可以一次只允许一个未完成 Load

真实 GPU 允许大量在途请求，但实现很复杂。

第一版可以规定：

每个 Warp 同时最多一个 Memory Request

这样只需要每个 Warp 保存一个 Load Context。

缺点：

• Memory-Level Parallelism 很低；
• 同一 Warp 不能连续发射多个独立 Load；
• 性能不高。

但功能正确后，可以扩展为：

每 Warp 2～4 个未完成请求

并让 Scoreboard 分别跟踪不同目标寄存器。

---

四十五、最小硬件中的 RAW、WAR 和 WAW

由于每个 Warp 顺序发射：

RAW

LDG r1
ADD r2, r1, r3

需要 Scoreboard 阻止提前读取。

WAR

ADD r3, r1, r2
MOV r1, r4

如果执行严格顺序、写回顺序不乱，通常不会产生问题。

WAW

LDG r1, [A]
LDG r1, [B]

如果两个 Load 可以同时在途，后发请求可能先返回，造成错误。

第一版每 Warp 只允许一个在途 Load，可以避免该问题。

以后支持多个在途请求时，需要：

• 更严格 Scoreboard；
• 序号；
• 或禁止同一目标寄存器存在多个未完成写入。

---

四十六、是否需要 Bypass

假设：

ADD r1, r2, r3
MUL r4, r1, r5

如果 ADD 结果在执行级产生，但还没写回 Register File，MUL 可以通过 Bypass 直接读取。

第一版可以不实现 Bypass：

必须等待写回后下一条才能读

功能简单，但增加1～数个周期 Stall。

第二版再加入：

ALU Result
    ↓
Forwarding Mux
    ↓
下一条指令输入

Bypass 可以减少短依赖链延迟，但不会代替长延迟 Load 的 Scoreboard。

---

四十七、最小 GPU 的控制流水线

可以设计为：

Stage 0：Warp Select
Stage 1：Fetch
Stage 2：Decode / Scoreboard Check
Stage 3：Register Read
Stage 4：Execute
Stage 5：Writeback

但如果在 Stage 2 才发现源寄存器 Busy，前面已经取出的指令需要停顿或取消。

更简单的方法是：

每个 Warp 先有 Instruction Buffer
Scheduler 只从操作数 Ready 的 Warp 中选择

也就是：

Fetch 与 Issue 解耦

这更接近真实 GPU 的基本思想。

---

四十八、推荐的教学型模块划分

1. kernel_launch_unit

负责接收 Kernel 参数和启动命令。

2. block_dispatcher

负责创建 Block 和 Warp。

3. warp_context_table

保存 PC、Active Mask、Block ID 等。

4. warp_scheduler

选择 Eligible Warp。

5. instruction_fetch

按 Warp PC 读取指令。

6. decoder

生成控制信号。

7. scoreboard

追踪未完成寄存器结果。

8. vector_register_file

保存所有线程寄存器。

9. simd_alu

执行4-Lane 算术。

10. load_store_unit

处理 Global Memory。

11. shared_memory_unit

处理片上共享存储和 Bank Conflict。

12. barrier_unit

处理 Block 级同步。

13. divergence_unit

管理 Active Mask 和汇合栈。

---

四十九、推荐实现顺序

不要同时实现所有模块。

Phase 1：单 Warp、无分支

支持：

MOV
ADD
MUL
EXIT

只有一个 Warp，不需要调度器和 Scoreboard。

验证：

4 个 Lane 能执行同一指令
但拥有不同寄存器数据

---

Phase 2：特殊寄存器和多 Warp

加入：

• Thread ID；
• Warp Context；
• 多个 Warp；
• Round-Robin Scheduler。

先不支持内存长延迟。

验证：

多个 Warp 可以交错执行

---

Phase 3：Global Memory

加入：

• LDG；
• STG；
• 固定延迟存储器；
• 每 Warp 一个未完成请求；
• Scoreboard。

运行：

C[i] = A[i] + B[i];

这是第一颗真正有实用意义的简化 GPU。

---

Phase 4：Predicate 和边界判断

加入：

• SETP；
• 谓词指令；
• 部分 Lane 屏蔽。

支持：

if (i < n)

处理最后一个不完整 Warp。

---

Phase 5：Shared Memory 和 Barrier

加入：

• LDS；
• STS；
• Shared Memory；
• BAR；
• Block Context。

运行：

• 数组反转；
• Block Reduction；
• 简单矩阵 Tile 加载。

---

Phase 6：分支发散

加入：

• True/False Mask；
• Divergence Stack；
• Reconvergence PC；
• 部分线程退出。

运行：

if (threadIdx.x & 1)
    ...
else
    ...

---

Phase 7：性能增强

加入：

• 内存合并；
• Shared Memory Bank；
• 多个未完成 Load；
• Bypass；
• 指令缓冲；
• 流水化乘法；
• 多 Block 驻留。

---

五十、验证策略

GPU 控制状态很多，必须分层验证。

单元验证

分别测试：

ALU
Register File
Scoreboard
Scheduler
LSU
Shared Memory
Barrier
Divergence Stack

指令级验证

为每条指令编写短程序：

ADD
MUL
LDG
STG
SETP
BRA
BAR
EXIT

Kernel 级验证

运行：

1. Vector Add；
2. SAXPY；
3. 条件 ReLU；
4. Shared Memory Reverse；
5. Block Reduction；
6. 小型矩阵乘法。

随机验证

随机生成：

• 指令序列；
• Active Mask；
• 内存延迟；
• Warp Ready 状态。

与软件 SIMT 模型对拍。

---

五十一、为什么需要软件参考模型

建议先用 Python 或 C++ 编写一个指令级模拟器。

模型维护：

warp_pc
active_mask
registers[warp][lane][reg]
shared_memory
global_memory
barrier_state

每次执行一条 Warp 指令。

RTL 仿真时比较：

每次提交后的寄存器
内存写入
PC
Active Mask
Warp 状态

这样比只看最终输出更容易定位错误。

---

五十二、可以定义 Commit 接口

每当一条指令完成，可以输出：

commit_valid
commit_warp_id
commit_pc
commit_opcode
commit_active_mask
commit_dst_reg
commit_dst_data[LANES]
commit_memory_address[LANES]

软件模型和 RTL 可以逐条指令对拍。

这与 CPU 验证中的退休接口或 Trace 接口类似。

对于 GPU，还必须比较：

• Warp ID；
• Active Mask；
• 每 Lane 数据。

---

五十三、关键断言

可以加入以下 SystemVerilog Assertion：

不能向 Busy 目标寄存器发起冲突写入
不能读取未 Ready 的源寄存器
Barrier 释放前 Warp 不能继续
Free Warp 不能被调度
Finished Warp 不能再次发射
Memory Response 必须对应有效请求
Active Mask 为 0 的 Warp 应结束或切换路径
Register 0 若定义为常数则不能被修改

例如：

assert property (
    issue_valid |-> source_operands_ready
);

---

五十四、综合时最可能遇到的瓶颈

Register File

Lane 数和 Warp 数增加后，寄存器容量及端口压力很大。

Scheduler

需要同时检查多个 Warp 的 Ready 条件，优先选择逻辑可能形成长组合路径。

Scoreboard

需要大量寄存器 Busy Bit 和多路比较。

Shared Memory Bank

多 Lane 地址译码、冲突检测和返回路由可能成为关键路径。

Crossbar

多 Lane、多执行单元和多 Bank 之间的数据选择会迅速增加面积。

浮点单元

FP32 FMA 延迟和面积远高于整数 ALU。

---

五十五、为什么商用 GPU 会划分多个 Scheduler Partition

如果一个 SM 有大量 Warp 和大量执行单元，一个中央 Scheduler 和统一 Register File 很难扩展。

所以可以把 SM 划分为多个子分区：

SM
├── Partition 0
│   ├── Warp Scheduler
│   ├── Register File Bank
│   └── SIMD Units
├── Partition 1
├── Partition 2
└── Partition 3

每个分区只管理一部分 Warp 和执行单元。

这样可以降低：

• 调度器扇入；
• Register File 端口；
• Crossbar 复杂度；
• 物理连线长度。

这与前面学习的现代 SM 分区思想一致。

---

五十六、从4 Lane扩展到32 Lane会发生什么

功能上只是：

LANES = 4
→
LANES = 32

但硬件复杂度不会只增加8倍。

还会带来：

• 32套数据通路；
• 更宽寄存器写回；
• 更复杂内存合并；
• 更复杂 Bank Conflict 检测；
• 更宽 Active Mask；
• 更复杂 Shuffle Network；
• 更大的旁路网络；
• 更高布线压力。

因此实际 GPU 常采用分区化和流水化，避免一个超大单周期32-Lane Crossbar。

---

五十七、从一个 SM 扩展到多个 SM

需要新增全局模块：

Grid / Block Dispatcher
多个 SM
L2 Cache 或共享内存系统
全局片上互连
Memory Controller
Atomic Unit

Block Dispatcher 需要选择哪个 SM 有足够：

• Warp Slot；
• Register；
• Shared Memory；
• Block Slot。

一个 Block 被分配后，其 Warp 必须全部位于同一 SM，才能共享 Shared Memory 和 Barrier。

---

五十八、多 SM 后最重要的新问题

1. Block 分配

不同 SM 何时接收新 Block。

2. Memory Arbitration

多个 SM 同时访问 Global Memory。

3. Cache

是否加入私有 L1 和共享 L2。

4. 原子操作

多个 SM 更新同一地址。

5. 全局完成检测

所有 Block 是否完成。

6. 公平性

某个 SM 或 Warp 是否长期得不到服务。

这些问题可以在单 SM 功能正确后再处理。

---

五十九、最小 GPU 和 NPU 的区别

最小 SIMT GPU 的核心是：

指令驱动
大量线程上下文
Warp 调度
Active Mask
通用 Load/Store

NPU 常更强调：

张量指令
规则数据流
DMA
片上 Buffer
矩阵阵列
较少线程级控制

GPU 用硬件线程调度隐藏不规则延迟。

NPU 更倾向提前规划数据搬运和计算流水。

因此 GPU 更灵活，但控制和寄存器成本更高。

---

六十、本课核心结论

第一，一颗教学型 SIMT GPU 可以从1个 SM、4个 Lane、4个 Warp 开始。

第二，Warp 是指令调度单位，Lane 是实际执行数据通路。

第三，一个 Warp 共享 PC 和 Active Mask，但每个 Lane 拥有独立寄存器数据。

第四，Block Dispatcher 负责把 Block 拆成 Warp，并初始化线程编号和 Active Mask。

第五，Vector Register File 可以按 Lane 分 Bank，避免设计超大统一多端口存储器。

第六，Warp Scheduler 从 Eligible Warp 中选择一个发射，Round-Robin 是合适的第一版策略。

第七，Scoreboard 通过 Busy Bit 阻止线程读取尚未写回的寄存器结果。

第八，SIMD ALU 对所有 Active Lane 执行同一操作，被屏蔽 Lane 不写回。

第九，Load/Store Unit 需要保存 Warp ID、目标寄存器、Lane Mask 和请求标识。

第十，Global Memory 长延迟可以通过切换到其他 Warp 隐藏。

第十一，Shared Memory 可以按 Bank 组织，冲突访问需要拆成多个周期。

第十二，Barrier Unit 按 Block 记录各 Warp 是否到达同步点。

第十三，分支发散需要分别执行 True Mask 和 False Mask，并在汇合点恢复线程集合。

第十四，第一版可以只支持 Uniform Branch，再逐步加入谓词和 Divergence Stack。

第十五，GPU RTL 应按照：

单 Warp
→ 多 Warp
→ Global Memory
→ Scoreboard
→ Predicate
→ Shared Memory
→ Barrier
→ Divergence
→ 性能优化

逐步实现。

第十六，软件指令级模型、Commit Trace 和模块级断言是验证 SIMT 控制逻辑的关键。

第十七，GPU 扩展的核心困难通常不是增加 ALU，而是：

$$\boxed{ \text{register ports} + \text{scheduling} + \text{memory system} + \text{synchronization} + \text{on-chip interconnect} }$$

下一课是本系列最后一课：

第二十七课：GPU 架构学习总结与统一分析框架

将最终整理：

看到 GPU 架构图应该先看什么
看到 CUDA Kernel 应该如何拆解
如何判断计算、带宽和延迟瓶颈
如何从算法映射到 GPU
GPU、NPU、FPGA、ASIC 如何选择
接下来适合阅读哪些文档
如何开始真正的 GPU RTL 项目