38. 并行计算与 GPU 编程模型

1. 并行体系结构分类

1.1 Flynn 分类法

按指令流与数据流的多少分为四类：

1.2 SIMD 与 SIMT

• SIMD：一条指令同时作用在一组数据上，程序员显式使用向量寄存器（如 AVX-512 的 512 位）。适合规整的数据并行。
• SIMT（单指令多线程）：GPU 的模型，程序员写的是标量线程代码，硬件把 32 个线程（一个 warp）锁步执行。同一 warp 内分支不一致会产生分支发散，两条路径串行执行。

关键区别：SIMD 是显式向量化（程序员管），SIMT 是隐式向量化（硬件管，程序员写标量）。

1.3 共享内存与消息传递

MIMD 又分两种编程模型：

• 共享内存：所有线程看到同一地址空间，用锁/原子同步。OpenMP、pthread、TBB 属此类，适合单机多核。
• 消息传递：各进程有独立地址空间，靠收发消息通信。MPI 属此类，适合集群。

2. 并行性能定律

2.1 Amdahl 定律

设程序中不可并行部分占比为 s，用 N 个处理器加速：

S(N) = 1 / (s + (1 - s) / N)

当 N→∞ 时，S→1/s。若 10% 串行，最多加速 10 倍。Amdahl 定律说明：优化串行部分是提升上限的唯一途径，这是固定问题规模下的悲观结论。

2.2 Gustafson 定律

若问题规模随处理器数一起增长：

S(N) = N - s(N - 1)

规模可扩展时，加速比近似线性。Gustafson 说明大问题下并行效率不会因串行部分而封顶——这解释了为什么 HPC 偏爱「更大规模」而非「更快单核」。

2.3 并行效率与可扩展性

• 效率 E = S / N，衡量处理器利用程度。
• 强扩展：固定规模加核，效率会下降（受 Amdahl 限制）。
• 弱扩展：规模随核数增长，效率常能保持（Gustafson 视角）。

2.4 开销来源

实际加速远低于理论值，因为还有：通信开销、同步开销、负载不均、访存带宽瓶颈、缓存一致性流量。

3. 共享内存并行与 OpenMP

3.1 OpenMP 基本用法

OpenMP 靠编译指令把串行循环并行化：

#include <omp.h>
#include <stdio.h>

int main(void) {
    long sum = 0;
    #pragma omp parallel for reduction(+:sum)
    for (int i = 0; i < 1000000; i++) {
        sum += i * i;
    }
    printf("threads=%d sum=%ld\n", omp_get_max_threads(), sum);
    return 0;
}

reduction(+:sum) 让每个线程有私有副本，最后安全合并，避免手工加锁。

3.2 常见子句

• schedule(static|dynamic|guided)：静态按块分，动态按需领，guided 先大后小。
• num_threads(N)：指定线程数。
• critical / atomic：临界区与原子操作。
• collapse(2)：把两层循环合并成一个并行维度。

gcc -fopenmp -O2 par.c -o par
OMP_NUM_THREADS=8 ./par

3.3 任务并行

除循环并行外，#pragma omp task 可表达不规则任务图，配合 taskwait 等待，适合递归分治（如快排、树遍历）。

4. 数据竞争与 false sharing

4.1 数据竞争

多个线程无同步地访问同一内存且至少一个写，结果依赖执行顺序，即数据竞争。它是未定义行为，不是「偶尔出错」。

// 错误：counter++ 是读-改-写三步，非原子
#pragma omp parallel for
for (int i = 0; i < N; i++) counter++;   // 结果必然小于 N

// 正确
#pragma omp atomic
counter++;

4.2 伪共享 false sharing

缓存以 cache line（通常 64 字节） 为单位。两个线程写同一 line 内不同变量时，line 在两个核心间来回弹跳，造成严重性能损失：

struct { long a; long b; } s;      // a 与 b 同处一个 cache line，坏

// 修复：填充到 64 字节对齐，各占独立 line
struct __attribute__((aligned(64))) Padded { long v; char pad[56]; };
Padded p[NUM_THREADS];

perf c2c record ./prog      # 检测 cache line 争用
perf c2c report

4.3 内存序与原子

弱内存模型下需显式屏障。C11 的 atomic_int 与 memory_order_acquire/release 提供可移植的顺序保证，避免用 volatile 当同步（volatile 不保证原子与顺序）。

5. MPI 消息传递模型

5.1 基本模型

MPI 是进程级并行：每个 rank 是独立进程，靠 MPI_Send / MPI_Recv 通信。SPMD 风格，同一份程序按 rank 走不同分支。

#include <mpi.h>
int main(int argc, char **argv) {
    MPI_Init(&argc, &argv);
    int rank, size;
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
    MPI_Comm_size(MPI_COMM_WORLD, &size);
    if (rank == 0) {
        int x = 42;
        MPI_Send(&x, 1, MPI_INT, 1, 0, MPI_COMM_WORLD);
    } else if (rank == 1) {
        int y;
        MPI_Recv(&y, 1, MPI_INT, 0, 0, MPI_COMM_WORLD, MPI_STATUS_IGNORE);
    }
    MPI_Finalize();
}

5.2 点对点通信

• 阻塞式：MPI_Send 可能等到对方接收才返回（取决于缓冲）。
• 非阻塞式：MPI_Isend / MPI_Irecv 立即返回，用 MPI_Wait 收尾，可重叠通信与计算。
• 死锁：两个 rank 都先 Send 再 Recv 且消息超缓冲，会互等。解法是 MPI_Sendrecv 或非阻塞配对。

5.3 集合通信

集合通信由实现优化（树形、环形、递归倍增），手写点对点几乎总是不如内置集合操作快。

mpicc mpi_demo.c -O2 -o demo
mpirun -np 4 ./demo

6. GPU 编程模型

6.1 线程层次

CUDA 的三层结构：

Grid（一次 kernel 启动）
  └── Block（可含 1024 线程，块内可共享内存同步）
        └── Thread（最小执行单元）

• warp：32 个线程一组，硬件锁步调度。warp 是调度与访存的基本单位。
• 一个 block 内可有多个 warp；block 被分配到某个 SM（流多处理器）上。

6.2 索引计算

__global__ void add(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}
// 启动：<<<grid, block>>>
add<<<(n + 255) / 256, 256>>>(a, b, c, n);

6.3 内存层次

共享内存是 GPU 性能的核心：把数据从全局内存搬到共享内存复用，可把访存次数降低一个量级。

6.4 合并访存

一个 warp 的 32 个线程访问连续且对齐的全局内存时，硬件合并成一次事务；若地址分散，则拆成多次，带宽利用率骤降。

// 好：连续访问，合并
c[i] = a[i] + b[i];
// 坏：跨步访问，非合并
c[i] = a[i * stride];

7. CUDA 与 OpenCL 核心概念对照

OpenCL 的优势是跨厂商（NVIDIA/AMD/Intel/FPGA），CUDA 的优势是生态与性能调优工具（Nsight、cuBLAS、cuDNN）。

7.1 统一内存与异步

• 统一内存（UM）：cudaMallocManaged 让 CPU 与 GPU 共享一个地址空间，缺页时驱动自动迁移，简化编程但可能引入隐式迁移开销。
• 流（stream）：同一流内串行，不同流可并发。异步拷贝 + 多流是隐藏传输延迟的标准手段。

cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream);
kernel<<<grid, block, 0, stream>>>(d_a);
cudaStreamSynchronize(stream);

8. 异构编程性能瓶颈

8.1 三大瓶颈

1. PCIe 传输：Host↔Device 带宽有限（PCIe 4.0 x16 约 32GB/s），远超计算时间就成了瓶颈。对策是增大计算密度、减少往返、用 pinned 内存。
2. 占用率（occupancy）：SM 上活跃 warp 数与上限之比。太低无法掩盖访存延迟，太高可能寄存器溢出。
3. 分支发散：同一 warp 走不同分支会串行化两条路径。

8.2 占用率权衡

占用率不是越高越好。寄存器用量决定每 SM 能驻留多少 warp：

nvcc --ptxas-options=-v kernel.cu   # 查看寄存器/共享内存用量

调优顺序通常是：先保证合并访存，再用共享内存复用，最后才调占用率。

8.3 性能分析工具

ncu --set full ./app            # Nsight Compute，逐 kernel 分析
nsys profile ./app              # Nsight Systems，看时间线与传输

关注指标：sm__throughput、dram__throughput、l1tex__t_sectors_per_request（访存合并度）。

9. 并行算法案例

9.1 归约 Reduce

把 N 个数求和。朴素做法是「每个线程加一个数，再串行合并」，快但不充分。树形归约让每步活跃线程减半：

// 共享内存树形归约（块内）
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
    if (tid < s) sdata[tid] += sdata[tid + s];
    __syncthreads();
}

关键点：__syncthreads() 保证每步同步；避免 tid % 2 的取模写法（会发散），用 tid < s。

9.2 前缀和 Scan

计算 out[i] = sum(in[0..i])。并行做法是 Hillis-Steele 算法：

for (d = 1; d < n; d *= 2)
    for each i in parallel:
        if (i >= d) a[i] += a[i - d];

复杂度 O(n log n)，但并行深度只有 log n。Blelloch 工作高效版把复杂度降到 O(n)，代价是两次遍历（上扫 + 下扫）。

9.3 矩阵乘

朴素三重循环是 O(n³) 且访存极差。GPU 优化的核心是 tiling（分块）：把 A、B 的子块搬进共享内存复用：

__shared__ float As[TILE][TILE], Bs[TILE][TILE];
for (int t = 0; t < n / TILE; t++) {
    As[ty][tx] = A[(ty + t*TILE)*n + tx];
    Bs[ty][tx] = B[(t*TILE + ty)*n + tx];
    __syncthreads();
    for (int k = 0; k < TILE; k++)
        sum += As[ty][k] * Bs[k][tx];
    __syncthreads();
}

共享内存复用把全局访存从 O(n³) 降到 O(n³/TILE)，这是 GPU 矩阵乘提速的关键。

10. 常见陷阱

• 用 Amdahl 定律吓自己：它假设问题规模固定，实际扩大规模常能线性加速，要结合 Gustafson 判断。
• 忽略伪共享：多线程写相邻变量导致 cache line 弹跳，性能比单线程还差，需按 64 字节对齐填充。
• 把 volatile 当同步：它不保证原子性和内存序，跨线程同步必须用原子或锁。
• MPI 阻塞调用顺序不当：双方都先 Send 会死锁，用非阻塞或 MPI_Sendrecv。
• 手写点对点替代集合通信：内置 Allreduce 有拓扑优化，手写环形往往更慢。
• GPU 非合并访存：跨步访问让带宽利用率跌到 1/32，先改数据布局再谈其他优化。
• 忽略 warp 发散：kernel 里的 if/else 会让同一 warp 串行执行两条路径，尽量让分支按 warp 对齐。
• 盲目追求高占用率：寄存器压力与共享内存用量会限制驻留 warp 数，占用率高不等于快。
• Host 与 Device 频繁往返：每次 cudaMemcpy 都是瓶颈，应批量传输并用流重叠。
• 共享内存缺 __syncthreads()：漏一次同步就会读到脏数据，且错误随机出现极难排查。

参考文章

• 计算机组成原理 — 流水线、SIMD 指令与多核结构
• 并发与同步 — 锁、原子操作与内存序
• 缓存一致性 — 多核缓存协议与伪共享的硬件根源
• 算法复杂度 — 并行算法的工作量与深度分析
• 进程与线程 — 线程调度与并行执行的基本单位