第六课：GPU 中的数据依赖、数据冒险与同步

导语：这一课讲 GPU 中的数据依赖与数据冒险，重点区分“依赖本身”和“流水线造成的冲突”，并说明 Scoreboard、原子操作和内存屏障如何保证正确性。

这一部分需要先区分两个容易混淆的概念：

数据依赖是程序本身的关系；数据冒险是硬件并行执行这些指令时可能发生的问题。

例如：

a = b + c;
d = a * e;

第二条语句必须使用第一条语句产生的 a，因此二者存在数据依赖。

即使在一个完全非流水线、一次只执行一条指令的处理器中，这种依赖仍然存在。但因为第二条指令一定在第一条完成后才开始，所以不会形成流水线数据冒险。

---

一、为什么流水线会产生数据冒险

假设加法需要四个流水级：

取指 → 译码 → 执行 → 写回

现在有两条指令：

I1：ADD R1, R2, R3
I2：MUL R4, R1, R5

I2 需要读取 I1 产生的 R1。

在流水线中，它们可能重叠：

周期    I1                I2
1       取指
2       译码              取指
3       执行              译码
4       写回 R1           准备读取 R1

如果 I2 在 R1 写回之前就读取，就可能得到旧值。

这就是数据冒险。

所以更准确地说：

数据依赖
    +
指令重叠执行
    ↓
可能产生数据冒险

---

二、三种常见数据依赖

1. RAW：Read After Write

读发生在写之后，也称为真依赖。

a = b + c;
d = a * e;

对应：

I1：写 a
I2：读 a

I2 必须等待 I1。

这是最重要、最常见的一类依赖。

---

2. WAR：Write After Read

写发生在读之后，也称为反依赖。

d = a + b;
a = e + f;

对应：

I1：读 a
I2：写 a

程序要求 I1 先读取旧的 a，然后 I2 才能覆盖它。

如果 I2 提前写入，I1 可能读到新值。

WAR 不表示数据真正从 I1 流向 I2，主要是因为二者使用了同一个名字 a。

---

3. WAW：Write After Write

两条指令写同一个位置，也称为输出依赖。

a = b + c;
a = d * e;

最终必须保留第二次写入的结果。

如果硬件因为执行延迟不同，导致第一条指令反而最后写回，就会产生错误。

---

三种依赖的总结

类型	前一条指令	后一条指令	性质
RAW	写	读	真依赖
WAR	读	写	名字依赖
WAW	写	写	名字依赖

在 GPU 中，RAW 是最常见、最直接需要硬件等待的依赖。

WAR 和 WAW 往往可以通过以下方式减少：

• 编译器重新分配寄存器；
• 使用不同的物理寄存器；
• 保证 Warp 内指令按顺序发射；
• 硬件依赖跟踪。

---

三、GPU 中的寄存器依赖

考虑一个 GPU 线程中的代码：

float a = x + y;
float b = a * z;
float c = p + q;

简化指令可能是：

I0：ADD R1, R2, R3
I1：MUL R4, R1, R5
I2：ADD R6, R7, R8

其中：

I1 依赖 I0
I2 不依赖 I0 和 I1

当 Warp 发射 I0 后，R1 的结果还未产生。

Scoreboard 会记录：

R1：Not Ready

因此 I1 不能立即发射。

但 I2 使用的是其他寄存器，理论上可以执行。

实际是否能够越过 I1 发射 I2，取决于具体 GPU 架构、编译结果和指令调度方式。通常 GPU 更依赖编译器提前安排独立指令，而不是采用 CPU 那样复杂的动态乱序执行。

---

四、Scoreboard 如何处理 RAW 依赖

Scoreboard 是 GPU 内部的依赖跟踪结构。

假设 Warp 的指令为：

LOAD R1, [A]
ADD  R2, R1, R3

当 LOAD 发射后：

R1 = Pending

ADD 需要读取 R1，因此不能发射。

Scoreboard 的工作过程可以简化为：

检查下一条指令的源寄存器
          ↓
源寄存器是否全部 Ready
          ↓
是：Warp 可以参与调度
否：Warp 暂时停顿

当内存数据返回：

R1 = Ready

Warp 才重新成为可发射 Warp。

---

五、Load-Use Hazard

GPU 中很常见的一种 RAW 冒险是：

float x = A[i];
float y = x + 1.0f;

简化后：

LOAD R1, [A+i]
ADD  R2, R1, 1

ADD 必须等待 LOAD 返回。

而 Global Memory 访问可能有较长延迟，所以这个等待可能远长于普通加法依赖。

执行过程可能是：

图示：五、Load-Use Hazard

这就是多 Warp 延迟隐藏的典型场景。

---

六、连续依赖链为什么影响性能

考虑：

a = b + c;
d = a + e;
f = d + g;
h = f + i;

形成一条长依赖链：

b,c → a → d → f → h

每一条指令都必须等待上一条完成。

即使执行流水线吞吐率很高，也无法让这一条依赖链中的指令完全重叠。

这种情况称为较低的指令级并行性。

相对而言：

a = b + c;
d = e + f;
g = h + i;
j = k + l;

这些指令相互独立，可以更好地填充流水线。

所以优化 GPU 程序时，有时会采用多个独立累加器。

较差：

sum += a0;
sum += a1;
sum += a2;
sum += a3;

形成连续依赖：

sum0 → sum1 → sum2 → sum3

较好的思路：

sum0 += a0;
sum1 += a1;
sum2 += a2;
sum3 += a3;

sum = sum0 + sum1 + sum2 + sum3;

这样可以增加独立指令数量。

---

七、同一 Warp 中不同线程的寄存器是否会冲突

通常不会。

每个线程都有自己的逻辑寄存器状态。

例如同一个 Warp 中执行：

x = x + 1;

虽然所有线程的源代码都叫 x，但实际是：

Thread 0 的 x
Thread 1 的 x
Thread 2 的 x
……

它们属于不同线程的寄存器上下文。

所以：

Thread 0 写自己的 R1

不会直接覆盖：

Thread 1 的 R1

寄存器依赖通常首先是同一个线程指令流内部的依赖。

---

八、不同 Warp 之间有寄存器依赖吗

通常没有直接寄存器依赖。

因为：

• 每个线程有自己的寄存器；
• 每个 Warp 由不同线程组成；
• Warp 之间不能直接读取彼此的寄存器。

例如：

Warp 0 的 Thread 0 的 R1

与：

Warp 1 的 Thread 0 的 R1

是不同的物理状态。

不同 Warp 如果需要交换数据，通常必须通过：

• Shared Memory；
• Global Memory；
• 原子操作；
• Warp/Block 同步机制。

因此，跨 Warp 的依赖主要是内存依赖，而不是寄存器依赖。

---

九、内存依赖比寄存器依赖更复杂

考虑两个线程：

// Thread 0
shared_data[0] = 100;

// Thread 1
int x = shared_data[0];

Thread 1 希望读取 Thread 0 写入的结果。

这里存在跨线程的 RAW 依赖：

Thread 0 写 shared_data[0]
Thread 1 读 shared_data[0]

但 GPU 的寄存器 Scoreboard 通常不会自动替程序员解决这种跨线程关系。

硬件并不知道程序员是否希望：

• Thread 1 读取旧值；
• Thread 1 读取新值；
• 两个线程谁先执行。

因此需要程序显式同步。

---

十、什么是数据竞争

考虑：

__shared__ int x;

if (threadIdx.x == 0) {
    x = 100;
}

if (threadIdx.x == 1) {
    int y = x;
}

这里没有保证 Thread 0 一定先写完。

可能出现：

情况一：
Thread 0 先写
Thread 1 后读
→ y = 100

也可能：

情况二：
Thread 1 先读
Thread 0 后写
→ y 可能是旧值

这种结果取决于执行时序的情况称为：

Data Race，数据竞争。

数据竞争通常包含：

1. 多个线程访问同一个存储位置；
2. 至少一个访问是写操作；
3. 缺少正确的同步或原子操作；
4. 执行结果依赖实际时序。

---

十一、使用 __syncthreads() 解决 Block 内依赖

正确写法可以是：

__shared__ int x;

if (threadIdx.x == 0) {
    x = 100;
}

__syncthreads();

if (threadIdx.x == 1) {
    int y = x;
}

执行过程：

图示：十一、使用 __syncthreads() 解决 Block 内依赖

__syncthreads() 同时提供两个重要功能：

1. 执行同步：Block 内参与线程都到达后才继续；
2. 内存可见性：同步点前的相关内存操作按规则对 Block 内线程可见。

---

十二、Barrier 和 Fence 不是同一回事

这两个概念容易混淆。

Barrier：屏障

屏障要求线程相互等待。

例如：

__syncthreads();

可以理解为：

所有线程都到齐
才能继续向后执行

Fence：内存栅栏

Fence 主要约束内存操作的顺序和可见性。

它不一定要求其他线程也停下来等待。

可以理解为：

保证我之前的写操作
在我之后的某些操作之前完成或可见

但它不代表其他线程已经执行到某个位置。

因此：

Barrier = 线程执行同步
Fence   = 内存顺序约束

有些同步原语会同时具有两种效果，但概念上需要区分。

---

十三、常见 Memory Fence

CUDA 中常见的内存栅栏概念包括：

__threadfence_block();
__threadfence();
__threadfence_system();

可以粗略理解为：

操作	主要约束范围
__threadfence_block()	当前 Block
__threadfence()	当前 GPU
__threadfence_system()	更大的系统范围，包括主机等参与者

它们主要保证当前线程的内存访问顺序。

但需要特别记住：

Fence 本身通常不会让其他线程等待。

例如：

data = 100;
__threadfence();
flag = 1;

其意图是让观察到 flag == 1 的线程，有机会按照正确协议确认 data 已经写入。

但读取方仍需要正确的原子操作、可见性规则或同步协议。

---

十四、原子操作是什么

考虑多个线程同时执行：

counter = counter + 1;

这条语句实际包含：

1. 读取 counter
2. 加 1
3. 写回 counter

假设初始值为 0。

Thread 0 和 Thread 1 同时执行：

Thread 0 读取 0
Thread 1 读取 0

Thread 0 计算 1
Thread 1 计算 1

Thread 0 写入 1
Thread 1 写入 1

最终结果是 1，而不是 2。

这叫做丢失更新。

使用原子操作：

atomicAdd(&counter, 1);

硬件保证这个读—改—写过程不可被其他线程插入破坏。

执行结果相当于：

Thread 0：counter 0 → 1
Thread 1：counter 1 → 2

---

十五、原子操作解决什么问题

原子操作常用于：

• 全局计数器；
• 直方图；
• 工作队列；
• 锁；
• 引用计数；
• 并行归约的部分阶段；
• 多线程更新同一数据。

常见原子操作包括：

atomicAdd
atomicSub
atomicMax
atomicMin
atomicCAS
atomicExch

其中 atomicCAS 是 Compare-And-Swap：

如果内存中的旧值等于期望值
就写入新值
否则不写入

它可以用来构建锁和其他同步结构。

---

十六、原子操作为什么可能很慢

如果许多线程同时更新同一个地址：

atomicAdd(&counter, 1);

这些操作必须以某种方式序列化。

例如 32 个线程访问同一个计数器：

Thread 0 更新
Thread 1 更新
Thread 2 更新
……

即使硬件能够优化，热点地址仍可能成为瓶颈。

因此常见优化方式是：

每个 Warp 或 Block 先局部累加
             ↓
只让少数线程执行全局原子更新

例如：

32 个线程
    ↓
Warp 内归约
    ↓
1 个线程 atomicAdd

这样可以把 32 次原子操作减少为 1 次。

---

十七、原子操作不等于全局同步

假设执行：

atomicAdd(&counter, 1);

它只保证该原子操作本身具有规定的原子性和内存语义。

它不表示：

• 所有线程都执行到这里了；
• 所有线程的其他计算都完成了；
• 所有 Block 都同步了；
• 整个 Kernel 到达统一屏障。

所以：

Atomic ≠ Barrier

原子操作用于保护共享数据更新，屏障用于协调执行阶段。

---

十八、不同 Block 能否使用 __syncthreads() 同步

不能。

__syncthreads() 只作用于一个 Block。

原因是不同 Block：

• 可能位于不同 SM；
• 可能不是同时驻留；
• 某些 Block 甚至还没有开始执行。

假设 GPU 一次只能运行部分 Block：

Block 0～7：正在执行
Block 8～15：尚未调度

如果前面的 Block 等待后面的 Block 到达屏障，就可能永远无法继续。

因此普通 Kernel 内通常不能依赖所有 Block 同时执行。

跨 Block 的全局同步通常采用：

1. 结束当前 Kernel；
2. 由 CPU 或运行时启动下一个 Kernel。

Kernel 边界天然形成一个全局执行阶段边界。

也可以使用协作组等特殊机制，但需要硬件和启动方式支持。

---

十九、Kernel 边界为什么可以用于全局同步

例如：

Kernel 1：所有 Block 生成中间结果
Kernel 1 完成
Kernel 2：读取中间结果继续计算

只有 Kernel 1 的全部 Block 完成后，依赖它的 Kernel 2 才会开始，前提是它们位于同一有序执行流并正确建立依赖。

因此：

Kernel 1 完成
      ↓
全局结果已经形成
      ↓
Kernel 2 开始

这是 GPU 编程中很常见的阶段划分方式。

---

二十、同一个 Warp 内是否需要同步

过去常假设同一 Warp 中线程严格锁步执行。

但现代 NVIDIA GPU 支持更灵活的线程调度，因此当 Warp 内线程通过 Shared Memory 或其他方式交换数据时，不应盲目依赖隐式锁步。

例如：

shared[lane] = value;

__syncwarp();

value = shared[lane ^ 1];

__syncwarp() 明确保证参与线程在 Warp 范围内同步。

如果只使用 Shuffle：

value = __shfl_sync(mask, value, source_lane);

该指令自身已经定义了参与线程和通信行为，但仍需要保证掩码和控制流使用正确。

---

二十一、存储顺序问题

考虑一个生产者线程：

data = 100;
flag = 1;

消费者线程：

if (flag == 1) {
    use(data);
}

程序员希望：

消费者看到 flag == 1
    ↓
也一定能看到 data == 100

但在并行系统中，以下因素可能使事情更复杂：

• 编译器重排；
• 硬件内存流水线；
• Cache；
• 写缓冲；
• 不同地址的访问完成顺序；
• 不同线程的观察顺序。

因此生产者—消费者通信通常需要：

• 原子变量；
• 合适的内存顺序；
• Fence；
• Barrier；
• 明确的同步协议。

不能只依赖普通的 flag = 1。

---

二十二、寄存器 Scoreboard 不能解决内存数据竞争

Scoreboard 能处理：

同一个 Warp 中
某条指令等待前面指令产生寄存器结果

但它不能自动推断：

Thread 0 写 shared[5]
Thread 17 读取 shared[5]

是否应该等待。

原因是：

• 地址可能运行时才能计算；
• 不同线程的语义由程序员定义；
• 两个访问也可能本来就允许并发；
• 自动同步所有潜在冲突会严重降低性能。

因此可以总结为：

寄存器依赖
→ 主要由硬件 Scoreboard 和编译器管理

跨线程内存依赖
→ 主要由程序员使用同步机制管理

---

二十三、GPU 为什么不大量使用 CPU 式乱序执行

高性能 CPU 会在单线程内部寻找可独立执行的指令，并动态乱序执行。

这需要复杂结构，例如：

• 寄存器重命名；
• 保留站；
• Reorder Buffer；
• 动态调度；
• 推测执行；
• 异常恢复。

GPU 的主要思路不同：

一个 Warp 等待
      ↓
执行另一个 Warp

也就是使用线程级并行性代替复杂的单线程乱序机制。

这样可以把更多芯片面积用于：

• 更多执行单元；
• 更大寄存器文件；
• 更多 Warp 状态；
• 更高显存带宽；
• Tensor Core。

GPU 并非完全没有动态调度，但通常不会像高性能 CPU 那样，为单个线程配备极其复杂的乱序执行窗口。

---

二十四、指令依赖和 Warp 调度的完整过程

假设 SM 中有四个 Warp：

Warp 0：
LOAD R1
ADD R2, R1, R3

Warp 1：
FP32 ADD

Warp 2：
等待同步

Warp 3：
Tensor Core MMA

执行过程：

1. Warp 0 发射 LOAD
2. Scoreboard 标记 R1 未就绪
3. Warp 0 的 ADD 暂时不能执行
4. 调度器选择 Warp 1
5. Warp 1 发射 FP32 ADD
6. 调度器选择 Warp 3
7. Warp 3 发射 Tensor Core 指令
8. Warp 0 的数据返回
9. R1 标记为 Ready
10. Warp 0 再次成为 Eligible Warp
11. 发射 ADD

这说明 GPU 并不一定消除某条指令的等待时间，而是用其他 Warp 的工作覆盖等待时间。

---

二十五、流水线停顿的几种依赖来源

1. 短延迟执行依赖

例如：

a = b + c;
d = a + e;

等待普通算术结果。

2. 长延迟内存依赖

例如：

a = global_memory[i];
d = a + e;

等待 Global Memory 数据。

3. Shared Memory 依赖

例如访问 Shared Memory 后立即使用结果，或者发生 Bank Conflict。

4. Tensor Core 累加依赖

连续矩阵累加使用同一累加寄存器时，也可能形成依赖链。

5. 原子操作依赖

多个线程对同一地址执行原子操作，可能产生序列化。

6. 同步依赖

Warp 等待：

__syncthreads();

直到其他线程到达。

---

二十六、依赖不一定意味着性能很差

几乎所有实际程序都有依赖。

例如：

x = A[i];
y = x * scale;
z = y + bias;

这是正常的计算链。

性能是否受严重影响，取决于：

• 依赖链有多长；
• 指令延迟有多大；
• 是否有其他独立指令；
• 是否有足够多可执行 Warp；
• 执行单元是否有其他工作；
• 访存是否命中缓存。

所以优化目标不是“消灭所有依赖”，而是：

提供足够的独立工作，覆盖无法避免的依赖延迟。

---

二十七、依赖、冒险和停顿的区别

这三个概念可以这样区分：

数据依赖

程序语义上的先后关系。

后一条指令需要前一条结果

数据冒险

因为指令重叠或乱序执行，可能违反依赖关系。

结果还没准备好，下一条就想读取

流水线停顿

硬件为了避免错误，暂时不发射或不推进相关指令。

等待数据准备完成

完整关系是：

图示：流水线停顿

---

二十八、GPU 是否使用数据转发

普通 CPU 流水线常使用 Forwarding 或 Bypassing：

执行单元结果
    ↓
直接送给下一条指令

而不是等结果先写回寄存器文件，再重新读取。

GPU 的具体转发网络取决于架构，通常也会有各种旁路和结果转发机制，以减少不必要的等待。

但从程序员角度，不需要显式控制这些通路。

入门时只需要理解：

• Scoreboard 判断结果是否可用；
• 可用结果可能通过写回或旁路提供；
• 如果结果尚不可用，Warp 会等待；
• 调度器转去执行其他 Warp。

---

二十九、从 RTL 角度看依赖处理

一个简化的依赖检查逻辑可以表示为：

图示：二十九、从 RTL 角度看依赖处理

结果写回时：

执行单元完成
    ↓
写回 dst
    ↓
Scoreboard 将 dst 标记为 Ready

对于多个 Warp，Scoreboard 状态必须带有 Warp 标识：

Warp ID + Register ID

因为不同 Warp 的同名寄存器属于不同上下文。

---

三十、本节核心结论

第一，数据依赖由程序本身决定，不是由流水线创造的。

第二，流水线和并行执行会使数据依赖表现为数据冒险。

第三，非流水线处理器中仍然存在数据依赖，但通常不会出现流水线数据冒险。

第四，RAW 是真依赖，后一条指令必须等待前一条结果。

第五，WAR 和 WAW 主要是寄存器或存储位置名称复用造成的依赖。

第六，GPU 使用 Scoreboard 跟踪 Warp 的寄存器结果是否准备好。

第七，一个 Warp 因依赖停顿时，Warp Scheduler 会尝试执行其他 Warp。

第八，同一线程的寄存器依赖通常由编译器和硬件自动处理。

第九，不同线程之间通过 Shared Memory 或 Global Memory 交换数据时，程序员必须正确同步。

第十，__syncthreads() 是 Block 级屏障，而 __syncwarp() 是 Warp 级同步。

第十一，Memory Fence 主要保证内存顺序，不等于线程屏障。

第十二，原子操作保证一次共享数据更新不可被其他线程破坏，但不等于全局同步。

第十三，不同 Block 通常不能在普通 Kernel 中使用 Block 屏障直接同步。

第十四，长依赖链会降低指令级并行性，但可以利用其他 Warp 隐藏延迟。

第十五，GPU 主要通过大量 Warp 调度，而不是复杂的 CPU 式单线程乱序执行来隐藏延迟。

下一课将进入 GPU 性能分析的核心模型：

第七课：GPU 到底是算力受限还是带宽受限

主要包括：

FLOPS 与显存带宽
Arithmetic Intensity
Compute Bound
Memory Bound
Roofline Model
为什么矩阵乘法适合 GPU
为什么数组加法通常跑不满算力
如何判断一个 Kernel 的性能瓶颈