第一章

TL; DR — 处理器的发展已经进入大规模并行和异构计算时代；GPU 通过大量线程提供高吞吐量，但要真正获得性能提升，程序员必须识别数据并行性，并同时处理任务划分、内存访问、数据传输、同步和可扩展性问题。

现代处理器性能提升正在从 提高单核频率 转向 增加并行计算单元。GPU 拥有大量计算核心，适合同时执行大量相似任务，因此能够为数据并行程序提供很高的吞吐量。

传统 CPU 性能提升受到功耗、散热和指令级并行度的限制，单纯提高时钟频率越来越困难。硬件厂商因此采用多核架构，但软件必须显式利用并行性，才能从新增核心中获得性能收益。这意味着程序员需要完成思维方式的转变：

  • 从顺序算法转向并行算法
  • 把计算拆分为大量可同时执行的任务
  • 关注数据访问、同步和任务之间的依赖
  • 考虑程序能否随处理器核心数量增长而扩展

CPU 主要针对低延迟和复杂控制流优化：

  • 核心数量较少，但单核能力强
  • 具有复杂的缓存和控制逻辑
  • 擅长分支、串行逻辑和不规则任务
  • 注重尽快完成单个任务

GPU 主要针对高吞吐量优化：

  • 拥有大量相对简单的计算核心
  • 将更多晶体管用于算术运算
  • 通过大量线程隐藏内存访问延迟
  • 擅长结构规则、计算密集、数据并行的任务

因此，GPU 并不是在所有场景下都比 CPU 快。只有当问题存在足够的并行性，并且计算与数据传输成本比例合理时，GPU 才能发挥优势。

现代系统通常采用 CPU 与 GPU 协同工作的异构架构：CPU 作为主机，负责程序控制、任务调度和串行部分；GPU 作为设备，负责大规模并行计算；数据需要在主机内存和设备内存之间进行管理或传输。

一个典型执行流程是：

  1. CPU 准备输入数据。
  2. 将数据传送到 GPU。
  3. CPU 启动 GPU 内核函数。
  4. GPU 使用大量线程并行处理数据。
  5. 将结果返回 CPU。

整体性能不能只看 GPU 内核的执行速度，还必须计算数据传输、内核启动和同步等开销。

GPU 最适合数据并行问题，即对大量数据元素执行相同或相似的操作。例如：

  • 向量和矩阵运算
  • 图像与视频处理
  • 科学计算和数值模拟
  • 机器学习
  • 大规模数据分析

优秀的并行程序应具有可扩展性：当硬件提供更多计算资源时，同一程序无需大幅修改就能利用这些资源。为此，程序通常需要创建数量远多于物理核心的线程，由硬件运行时负责调度。

CUDA 提供了一种面向 NVIDIA GPU 的并行编程平台。程序员可以编写内核函数，并指定大量线程执行该函数。CUDA 的关键价值包括：

  • 使用接近 C/C++ 的方式表达并行计算
  • 将线程组织为层次化结构
  • 由硬件和运行时完成线程调度
  • 允许程序员控制设备内存和数据移动
  • 使程序能够适配不同规模的 GPU

GPU 编程本质上是围绕大规模线程、数据并行和异构执行展开的。

GPU 拥有很高的理论运算能力，但实际性能还取决于：

  • 算法中可并行部分的比例
  • 内存访问效率
  • CPU 与 GPU 之间的数据传输
  • 线程间同步和通信
  • 分支导致的执行分歧
  • 是否有足够多的线程占满硬件
  • 算法的计算密度

根据 Amdahl 定律，程序中的串行部分会限制整体加速比。即使并行部分运行得非常快，只要串行部分或数据传输占比较高，总体提升仍然有限。

第二章

nvcc_process

第七章

TL; DR — 减少 DRAM 访问，提高计算强度，让程序从 内存带宽受限 转向 计算受限

  • 输出元素并行化；
  • 所有只读且所有线程共同访问的卷积核放入常量内存；
  • 会重复读取的数据分块放入共享内存，并讨论 halo 数据如何处理；

考查一个简化的二维卷积的代码

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
__global__ void convolution2D_basic(
    const float* N,
    const float* F,
    float* P,
    int radius,
    int width,
    int height)
{
    int outCol = blockIdx.x * blockDim.x + threadIdx.x;
    int outRow = blockIdx.y * blockDim.y + threadIdx.y;

    if (outCol >= width || outRow >= height) {
        return;
    }

    float value = 0.0f;

    for (int fRow = 0; fRow < 2 * radius + 1; ++fRow) {
        for (int fCol = 0; fCol < 2 * radius + 1; ++fCol) {
            int inRow = outRow - radius + fRow;
            int inCol = outCol - radius + fCol;

            if (inRow >= 0 && inRow < height &&
                inCol >= 0 && inCol < width) {
                value += F[fRow * (2 * radius + 1) + fCol]
                       * N[inRow * width + inCol];
            }
        }
    }

    P[outRow * width + outCol] = value;
}
  • 边界处存在分支发散：同一个 warp 中，有些线程位于图像内部，不会越界；有些线程位于边缘，部分滤波器位置越界；有些线程位于角落，大量滤波器位置越界。因此同一个 warp 内不同线程的 if 结果不一致，产生控制流分歧。（对大图像和小滤波器来说，边界像素只占总像素的一小部分，所以这一问题通常不是最严重的瓶颈）
  • 全局内存访问过多：假设滤波器大小为 ，则每个输出需要读取 个输入元素， 个滤波器元素，最后写一个输出元素。（每个 float 按 4 个字节计算），则浮点数操作数为 ；需要进行 次乘法和 次加法，浮点数操作为 ；得到算存比为

这对 GPU 来说非常低，严重受到 DRAM 带宽限制，大量计算单元在等待数据而不是计算。

常量内存和常量缓存

卷积滤波器通常比较小，运算期间不会被修改，并且所有线程按照相同顺序读取同一个滤波器元素，这种性质适合 CUDA 的常量内存与常量缓存。

假设滤波器半径是编译期常量：

1
2
3
4
#define FILTER_RADIUS 2
#define FILTER_WIDTH (2 * FILTER_RADIUS + 1)

__constant__ float F[FILTER_WIDTH][FILTER_WIDTH];  // 在函数外部声明

__constant__ 表示这个变量位于设备常量地址空间中。然后从主机端复制：

1
cudaMemcpyToSymbol(F, F_h, FILTER_WIDTH * FILTER_WIDTH * sizeof(float));

其优势来自专门的常量缓存。当一个 warp 中所有线程访问同一个常量地址时，硬件将这个值广播给整个 warp，原本可能是 32 个线程分别提出相同数据请求；使用常量缓存后，读取一次就可以向多个线程广播。

1
2
F[fRow][fCol]   // 所有线程的 fRow 和 fCol 相同
F[threadIdx.x]  // 同一个 warp 中不同线程访问不同地址时，常量访问可能被串行化

参考资料

附录

Amdahl 定律

TL; DR — 在问题规模固定、串行与并行比例固定、主要关注缩短执行时间、不考虑并行化开销的前提下，Amdahl 定律 用于估算当程序只有一部分能够被并行化时，增加处理器数量或加速并行部分，整个程序最多能获得多大的加速。

  • 整体优化效果取决于被优化部分在原运行时间中的占比；
  • 少量串行代码也可能限制最大加速比；
  • 处理器数量增加时，性能收益通常递减；
  • 优化前应先测量瓶颈，优先优化耗时占比最大的部分；
  • 并行化时必须把通信、同步和数据传输纳入总成本；

设 表示程序中可以并行或被优化的比例， 无法并行的串行比例； 表示处理器数量，或并行部分的加速倍数。则程序整体加速比（优化前运行时间/优化后运行时间）表示为

让处理器数量趋近于无穷：

称为最大理论加速比。

在 CPU 和 GPU 协同计算中，串行或不可加速的部分可能包括：

  • CPU 上的串行逻辑
  • CPU 与 GPU 之间的数据传输
  • GPU 内核启动开销
  • 线程同步
  • 无法有效并行的算法阶段
  • 文件读取和网络访问
  • 负载不均衡造成的等待

基本公式假设并行部分可以被完美地平均分配，但实际系统还存在：

  • 创建和调度线程的开销
  • 锁竞争与同步开销
  • 缓存一致性开销
  • 内存带宽限制
  • 处理器之间的通信
  • 工作分配不均
  • GPU 分支分歧
  • CPU 与 GPU 数据复制

因此，更接近实际情况的时间模型可以写成：

其中 表示并行化带来的额外开销，并且这个开销不一定随 线性增长，只是这里用来概括各种附加成本。因此，增加处理器有时不仅不再提速，甚至可能使程序变慢。

Gustafson 定律

TL; DR — 在总运行时间固定、问题规模能够随处理器数量增加而扩大、串行部分耗时基本不变、并行工作能够均匀分配且暂不考虑并行化开销的前提下，Gustafson 定律 用于估算增加处理器后，系统能够处理多大规模的问题，以及由此获得的可扩展加速比。

  • 增加处理器的主要价值不只是缩短固定任务的运行时间，也可以是在相同时间内处理更大的问题；
  • 当问题规模扩大时，并行部分可以随处理器数量增长，而串行部分通常保持相对稳定；
  • 只要串行比例足够小，整体加速比就可以接近处理器数量；
  • Gustafson 定律没有 Amdahl 定律中固定的最大加速上限，但串行部分仍会降低并行效率；

设 表示程序在使用 个处理器运行时，串行部分占总运行时间的比例， 表示并行部分占总运行时间的比例； 表示处理器数量。将并行程序在 个处理器上的总运行时间归一化为 。如果改为单处理器执行，串行部分耗时仍为 ，而并行部分需要依次执行，耗时变为 。因此，程序的可扩展加速比表示为

整理后也可以写成：

其中：

  • 是使用 个处理器时测得的串行时间比例；
  • 是并行时间比例；
  • 表示相同时间尺度下，单处理器完成等量工作的时间与多处理器运行时间之比。

当处理器数量增加时：

如果 保持不变，则加速比可以随 近似线性增长，不存在类似 Amdahl 定律 这样的固定理论上限。但其并行效率为：

代入 Gustafson 定律可得：

让处理器数量趋近于无穷：

因此，Gustafson 定律虽然允许加速比随处理器数量增长，但串行比例仍然决定了系统能够达到的长期并行效率。

在 CPU 和 GPU 协同计算中，适合通过扩大问题规模获得收益的场景包括：

  • 处理更高分辨率的图像或视频
  • 扩大矩阵和张量的规模
  • 增加科学模拟的网格数量
  • 增加机器学习的批次大小
  • 同时处理更多数据样本
  • 提高数值计算的精度或迭代次数
  • 扩大数据库查询或数据分析的范围

Amdahl 定律与 Gustafson 定律的核心区别是：

对比项 Amdahl 定律 Gustafson 定律
问题规模 固定 随处理器数量扩大
关注目标 缩短固定任务的运行时间 在固定时间内处理更多工作
串行比例 通常基于单处理器原始程序 通常基于多处理器运行时间
加速上限 受串行部分限制，存在固定上限 理想情况下可随处理器数量增长
典型问题 固定任务最快能运行多快 更多处理器能处理多大的任务
适用场景 固定数据集、固定工作量 科学计算、大数据、仿真和模型训练

简而言之，Amdahl 定律强调固定工作量下串行部分对加速比的限制；Gustafson 定律强调通过扩大并行工作量，可以让更多处理器持续发挥作用。