第4章数据并行与任务并行

第4章：数据并行与任务并行

并行计算，说白了就是让多个计算单元同时干活。但怎么分活、怎么协调，这里头门道很深。我这些年调过的编译器，十有八九的优化空间都卡在「并行模型选型」上。

今天咱们就聊聊两种最核心的并行模型：数据并行和任务并行。嗯，先别急着翻书，我用自己的踩坑经历给你讲明白。

4.1 数据并行模型：SIMT 与 SIMD

数据并行，核心思想就一句话：同一指令，操作不同数据。你想想看，如果要对一万个数做加法，最笨的办法是一个一个加，聪明的办法是让一堆ALU同时加。

4.1.1 SIMD：单指令多数据

SIMD 是 CPU 上的老牌并行技术。从 MMX、SSE 到 AVX-512，Intel 折腾了二十多年。说白了，就是一条指令能同时处理多个数据元素。

核心特征：

数据宽度固定（128位、256位、512位）
所有处理单元同步执行
编译器自动向量化是关键

举个例子，用 AVX2 做向量加法：

// C 代码
float a[8], b[8], c[8];
for (int i = 0; i < 8; i++) {
    c[i] = a[i] + b[i];
}

// 编译器自动向量化后，实际生成：
// vmovups ymm0, [a]
// vaddps  ymm0, ymm0, [b]
// vmovups [c], ymm0

我个人习惯在写循环时，尽量让迭代次数是向量宽度的整数倍。为什么？因为非对齐的尾部循环，性能会断崖式下跌。我在项目中遇到过，一个图像滤波算法，就因为没注意对齐，性能差了3倍。

4.1.2 SIMT：单指令多线程

SIMT 是 GPU 的看家本领。和 SIMD 不同，SIMT 的「线程」是逻辑概念，硬件上实际跑的是 warp（NVIDIA）或 wavefront（AMD）。

我记得第一次调 GPU 内核时，死活想不通：为什么同一个 warp 里的线程，执行路径不一样时性能会崩？后来才明白，SIMT 本质上还是「锁步执行」——warp 里的线程必须走同一条指令流。

避坑指南：我曾经在一个流体力学模拟中，用了大量 if-else 分支。结果一个 warp 里32个线程，一半走 if，一半走 else，实际执行时间翻倍。因为硬件会把两个分支串行化！

SIMT 的典型编程模型是 CUDA：

__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

这里每个线程处理一个元素，但硬件上是以 warp（32线程）为单位调度。你想想看，如果 n 不是32的倍数，最后一个 warp 里有些线程就闲置了。嗯，这就是「线程发散」问题。

4.2 任务并行模型

任务并行和数据并行不同。数据并行是「人多活一样」，任务并行是「各干各的」。比如一个线程做矩阵乘法，另一个线程做数据压缩，两者互不干扰。

任务并行的核心挑战在于：任务划分和负载均衡。

4.2.1 静态任务分配

最简单的方式：编译时就确定每个线程/核要执行什么任务。比如 OpenMP 的 parallel for：

#pragma omp parallel for num_threads(4)
for (int i = 0; i < 1000; i++) {
    process(i);
}

这里把1000次迭代平均分给4个线程，每个线程处理250次。但问题来了——如果 process(i) 的计算量不均匀呢？

我的经验：静态分配适合计算量均匀的场景。如果任务粒度差异大，建议用动态调度。我曾经在粒子模拟中，因为粒子密度分布不均，静态分配导致3个线程闲死，1个线程累死。换成 dynamic 调度后，性能提升40%。

4.2.2 动态任务分配

动态分配让运行时系统来调度任务。OpenMP 的 dynamic 调度就是典型：

#pragma omp parallel for schedule(dynamic, 16)
for (int i = 0; i < 1000; i++) {
    process(i);
}

这里每个线程一次拿16个任务，做完再拿下一批。说白了就是「谁闲谁干活」，负载均衡自然就好了。

但动态调度有代价——任务窃取、同步开销。我建议 chunk size 不要太小，否则调度开销会吃掉并行收益。

4.3 同步与通信开销分析

并行计算最头疼的问题，不是怎么让计算跑起来，而是怎么让计算不乱。同步和通信，就是为此付出的代价。

4.3.1 同步开销

同步包括：屏障（barrier）、锁（lock）、原子操作（atomic）。每种都有开销。

同步方式	典型延迟	适用场景
__syncthreads()	~20 cycles	GPU 块内同步
pthread_mutex_lock	~50-100 ns	CPU 线程互斥
atomicAdd	~30-100 cycles	GPU 全局原子操作
MPI_Barrier	~1-10 μs	分布式节点同步

你想想看，如果一个 GPU kernel 里每个循环都加 __syncthreads()，那性能基本就废了。我见过一个新手写的代码，每个线程处理一个元素后都要同步一次，结果99%的时间花在等待上。

避坑指南：我曾经在优化一个归约算法时，用了全局原子操作来做累加。结果256个线程同时 atomicAdd，总线都快挤爆了。后来改成先块内归约，再全局原子操作，性能提升了8倍。

4.3.2 通信开销

通信开销包括：数据搬运（host↔device）、线程间数据共享、跨节点网络传输。

以 GPU 为例，数据从 CPU 传到 GPU 的 PCIe 带宽大约是 16 GB/s（PCIe 3.0 x16）。而 GPU 显存带宽可以到 900 GB/s。差了50多倍！

所以我的原则是：能少搬就少搬，能一次搬完就别分多次。

// 错误做法：每次迭代都搬数据
for (int i = 0; i < 100; i++) {
    cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
    kernel<<<grid, block>>>(d_data);
    cudaMemcpy(h_result, d_result, size, cudaMemcpyDeviceToHost);
}

// 正确做法：一次搬完，在 GPU 上循环
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
for (int i = 0; i < 100; i++) {
    kernel<<<grid, block>>>(d_data);
}
cudaMemcpy(h_result, d_result, size, cudaMemcpyDeviceToHost);

嗯，这里要注意：GPU kernel 启动本身也有开销（约 5-10 μs）。如果 kernel 执行时间太短，启动开销占比就很大。我建议每个 kernel 至少运行 100 μs 以上，否则不如在 CPU 上跑。

4.4 核心知识体系

下面这张图，是我梳理的本章知识结构。你可以把它当作调优时的检查清单：

这张图把并行模型的核心脉络理清了。你调优的时候，可以对照着看：当前瓶颈在数据并行还是任务并行？同步开销占比多少？通信是否可隐藏？

我的建议：刚开始做并行优化时，别急着上复杂模型。先用 profiler 看看热点在哪。我见过太多人一上来就搞动态任务窃取，结果发现瓶颈在内存带宽上——白费功夫。

好了，数据并行和任务并行就聊到这儿。记住一句话：没有银弹。SIMD 适合规则计算，SIMT 适合大规模数据并行，任务并行适合不规则负载。选对模型，性能翻倍；选错了，debug 到天亮。

公众号:蓝海资料掘金营，微信deep3321

📚 课程目录