前言

第一章虽然是引言,但信息量很大。系统梳理了并行计算的发展历程、CPU与GPU的架构差异、CUDA的诞生背景。理解这些历史和设计思想,对后续学习GPU编程很有帮助。

📦 配套资源:本系列文章配有完整的 GitHub 仓库,包含每章的练习题解答、CUDA 代码实现和详细注释。所有代码都经过测试,可以直接运行。

性能提升的转折点:从频率到并行

2004年之前:频率提升的黄金时代

在这个阶段,提高 CPU 频率是提升程序性能的主要方式。摩尔定律非常直观:晶体管数量每18个月翻倍,频率提升,性能自然增长。程序员无需特别关注性能优化,等待硬件升级就能获得免费的性能提升。

功耗墙的出现

2004年前后,CPU 发展遇到了物理极限——功耗墙(Power Wall):

  • 动态功耗正比于频率和电压的平方
  • 提高频率需要提高电压以保证稳定性
  • 功耗以超线性速度增长,散热成为瓶颈

当时 Pentium 4 的热设计功耗(TDP)已达130W以上,继续提升频率将导致功耗达到数百瓦,在消费级产品中难以实现。

转向多核

既然单核频率无法继续提升,芯片厂商转向多核:

  • 单核保持较低频率(控制功耗)
  • 通过并行提升整体性能
  • 利用不断增长的晶体管数量

但多核CPU受限于功耗和面积,通常只有几个到十几个核心,对于需要成千上万并发任务的应用(图形渲染、深度学习)还不够。

CPU vs GPU:设计哲学的根本差异

CPU:最小化延迟

目标:让单个线程尽可能快地完成任务

手段

  • 复杂控制逻辑:分支预测、乱序执行、寄存器重命名、推测执行
  • 庞大缓存:L1(32-64KB) + L2(256KB-1MB) + L3(8-64MB)
  • 强大计算单元:完备的ALU/FPU,SIMD指令集

代价:芯片上只有20-30%的晶体管用于计算,其余都是控制逻辑和缓存。单核成本高,核心数量有限。

适合:复杂控制流、不规则内存访问、串行依赖强的任务

GPU:最大化吞吐量

目标:通过海量线程提升整体吞吐量

手段

  • 大量简单核心:数千个核心,每个都很简单(无复杂分支预测/乱序执行)
  • SIMT模型:32个线程为一组(warp),共享指令单元
  • 高带宽内存:500-1000 GB/s(CPU只有50-100 GB/s)
  • 专注浮点运算:晶体管主要用于计算单元

特点:单核简单,但数量多。在数据并行任务上,浮点性能可达CPU的10倍以上。

适合:数据并行、计算密集、规则内存访问的任务

形象类比

  • CPU:经验丰富的专家,处理复杂问题效率高,但一次只能专注一件事
  • GPU:数千名新手组成的团队,单个能力有限,但通过大规模协作完成简单任务效率惊人

CUDA的诞生

CUDA之前的黑暗时代

早期GPGPU必须通过图形API(OpenGL/Direct3D)实现:

  1. 将数据编码成纹理(Texture)
  2. 编写顶点着色器和片段着色器
  3. 通过渲染到纹理获取结果
  4. 读回主内存

问题

  • 需要深入理解图形管线
  • 数据表示受限(只能用颜色通道)
  • 调试几乎不可能
  • 只有图形学专家能用

CUDA 的革新(2007)

硬件改进

  • 统一着色器架构
  • 支持任意内存读写(不局限于纹理)
  • 提供共享内存(Shared Memory)等可编程存储
  • 支持整数运算、位运算

软件生态

  • 基于 C/C++ 的扩展,学习成本低
  • 完整工具链(nvcc 编译器)
  • 丰富的数学库(cuBLAS、cuFFT)
  • 调试和性能分析工具

示例代码

1
2
3
4
5
6
__global__ void vecAdd(float *A, float *B, float *C, int n) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < n) {
C[i] = A[i] + B[i];
}
}

相比之前需要编写复杂着色器,CUDA代码直观得多。这是GPU通用计算普及的关键转折点。

异构计算:CPU 与 GPU 协作

GPU 虽强大,但不适合所有任务:

场景 GPU 优势 GPU 劣势
计算类型 大量浮点运算 复杂控制流
并行度 数据并行 线程数少的任务
内存访问 规则访问模式 不规则/随机访问
同步 独立计算 频繁同步

如果程序只有几百个线程(Thread),在 GPU 上可能不如 CPU。GPU 需要数千甚至数百万线程才能充分利用硬件。

典型 CUDA 程序流程

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
1. 主机(Host)初始化
- 分配主机内存
- 初始化数据

2. 数据传输到设备(Device)
cudaMemcpy(..., cudaMemcpyHostToDevice);

3. 启动 GPU 核函数(Kernel)
kernel<<<gridDim, blockDim>>>(args);

4. GPU 并行执行

5. 结果传回主机
cudaMemcpy(..., cudaMemcpyDeviceToHost);

6. 主机端后续处理

现代深度学习框架(PyTorch、TensorFlow)底层都是这个模式:

  • 模型构建在 CPU
  • 前向/反向传播在 GPU
  • 优化器更新回到 CPU

并行编程的挑战

1. 算法并行化

不是所有算法都能直接并行化:

  • 递归算法(快速排序、树遍历)
  • 动态规划(许多DP问题有串行依赖)
  • 累积计算(前缀和、累积乘积)

有时需要完全重新设计算法(如并行前缀和),而不是简单并行化串行版本。

2. 内存带宽瓶颈

假设GPU峰值10 TFLOPS,内存带宽500 GB/s。如果计算需要读2个float(8B)做1次加法:

  • 计算需求:10 TFLOPS
  • 内存支撑:500 GB/s ÷ 8B = 62.5 GFLOPS
  • 实际性能:只能达到峰值的0.6%

这就是内存受限(Memory-Bound)问题。优化方向:

  • 使用Shared Memory减少全局内存访问
  • 优化访问模式实现合并访问(Coalesced Access)
  • 提高算术强度(每字节数据做更多计算)

3. 性能不确定性

并行程序性能高度依赖输入数据。例如数组过滤:

1
2
3
if (data[i] > threshold) {
output[count++] = data[i];
}
  • 全部满足条件 → 负载均衡,性能好
  • 少数满足 → 大部分线程空闲
  • 分布不均 → warp分支发散

同样的代码,不同输入可能有数倍性能差异。

4. 调试复杂性

并发错误在串行程序中根本不存在:

  • 竞态条件(Race Condition):多线程同时访问同一内存
  • 死锁(Deadlock):线程间相互等待
  • 非确定性错误:随机出现,难以重现

性能优化也需要深入理解硬件细节:

  • Warp(32线程一组)
  • Occupancy(活跃线程比例)
  • Bank冲突、寄存器溢出等

本书的学习目标

1. 并行编程技术

以CUDA为平台,学习:

  • 基本概念(线程、同步、内存模型)
  • CUDA语法和API
  • 常见并行模式(Map、Reduce、Scan、Stencil)
  • GPU硬件架构

2. 并行计算思维

更深层次的能力:

  • 识别问题中的并行性
  • 设计高效并行算法
  • 权衡并行度、负载均衡、通信开销
  • 针对硬件优化算法

3. 正确性与性能

实践技能:

  • 避免并发错误,测试并行程序
  • 使用profiling工具
  • 系统化性能调优
  • 编写可移植的代码

这些方法论比单纯学API更有价值,能迁移到其他并行平台(OpenCL、SYCL、HIP)。

小结

第一章建立了并行计算的大框架:

历史脉络:功耗墙导致从频率提升转向并行计算,这是必然趋势。

架构差异:CPU追求低延迟(复杂核心、大缓存),GPU追求高吞吐(简单核心、高带宽)。理解这点,就明白为什么某些任务适合GPU,某些不适合。

CUDA价值:不仅是技术,更是生态。虽然是专有平台,但完整的工具链和库让它成为GPU编程的事实标准。

挑战认识:并行编程确实比串行复杂,但掌握的思维方式和优化方法具有通用性。

下一章将进入 CUDA 编程实战,通过向量加法这个经典例子,学习 CUDA 程序的基本结构和核心概念。


参考资料:

  • Hwu, W., Kirk, D., & El Hajj, I. (2022). Programming Massively Parallel Processors: A Hands-on Approach (4th Edition). Morgan Kaufmann.
  • NVIDIA CUDA官方文档

本文 GitHub 仓库: https://github.com/psmarter/PMPP-Learning