GPGPU 编程入门

GEA 把 GPGPU 讲成「SIMD 向量化推到极致」：GPU 结合 SIMD（向量化 ALU）与 MIMD（一种抢占式多线程），即 NVIDIA 所说的 SIMT。适合 GPU 的任务必须是数据并行的——每个元素的计算不依赖其他元素的结果。写法与 CPU 向量化相反：不把循环硬编码成每次 4/8/16 个，而是保留单通道的循环体，抽成计算核函数 (compute kernel)，由 GPU 替我们「跑循环」。一个「GPU 线程」其实就是一条 SIMD 通道。

数据并行计算（卷 I 4.11.1）

GPU 是专为高度数据并行的计算设计的协处理器。各厂商、各产品线设计差异很大，但都遵循 SIMT 的一般原则。
判断标准：对任一元素的计算独立于其他元素，可按任意顺序处理。SIMD 向量化中的 DotArrays_ref 就是例子：正因各轮迭代独立，才能用 SSE/AVX 一次算 4/8/16 个。
设想 1024 通道的 VPU：1024 个以内的元素一次迭代就算完。GPU 大致如此，但并不用 1024 宽的 SIMD：其 SIMD 单元通常 8 或 16 通道宽，按 32 或 64 个元素一批处理，而且有很多个这样的单元，所以能并行处理成千上万个元素。
像素着色器（片元着色器）每帧处理数百万像素、顶点着色器处理数十万到数百万顶点，正是数据并行；现代 GPU 也允许写通用的计算着色器，只要计算基本相互独立，通常比 CPU 更高效（图形 API 侧见计算着色器）。

计算核函数（卷 I 4.11.2）

CPU 向量化要改写代码：换成 __m128 等向量类型、用内建函数、循环每次固定处理 N 个。GPGPU 则保留标量的「单通道」写法，把循环体抽成核函数 DotKernel(i, r, a, b)，处理一个输入元素、产出一个输出元素，类似像素着色器一进一出、顶点着色器一进一出。GPU 等于替我们执行 for 循环，对每个元素调用一次核函数。
着色语言大多语法接近 C：DirectX 的 HLSL、OpenGL 的 GLSL、NVIDIA 的 Cg 和 CUDA C、OpenCL。有的要求核函数放在独立源文件；OpenCL 和 CUDA C 是 C++ 的扩展，核函数可写成普通 C/C++ 函数，只需少量语法调整。
CUDA C 版本用 __global__ 标记，循环下标取自编译器提供的「魔法」变量 threadIdx.x（就像成员函数中的 this），并检查 if (i < count)。

执行核函数（卷 I 4.11.3）

书中 CUDA 示例的流程如下：

flowchart LR
  A["cudaMallocManaged 分配 CPU/GPU 都可见的缓冲"] --> B["memcpy 输入数据"]
  B --> C["DotKernel_CUDA<<<G, N>>> 提交给驱动"]
  C --> D["cudaDeviceSynchronize 等待 GPU 完成"]
  D --> E["拷回结果，cudaFree 释放"]

原书问题（4.11.2–4.11.3 示例代码）：(1) 缓冲声明成 int *cr, *ca, *cb，应为 float*；(2) 核函数声明为 DotKernel_CUDA(int count, float* r, const float* a, const float* b)，调用时却写成 (cr, ca, cb, count)，参数顺序不符；(3) threadIdx.x 只是线程在本线程块内的下标，而每个线程块最多 1024 个线程，所以 <<<1, count>>> 在 count 超过 1024 时会启动失败；通用写法是 i = blockIdx.x * blockDim.x + threadIdx.x，并用 <<<(count + 255) / 256, 256>>> 之类的配置（依据：NVIDIA CUDA C++ Programming Guide 的 Thread Hierarchy 一节）。另外，cudaMallocManaged 分配的托管内存本身 CPU 可直接读写，可以直接在其中生成数据，memcpy 只是为了沿用原有数组。

书中称「OpenCL 和 CUDA C 是 C++ 的扩展」也不准确：CUDA C++ 确实与宿主代码同源编译；但 OpenCL 的核函数是用 OpenCL C（或 C++ for OpenCL）写的独立源码，通常在运行时由驱动编译，宿主端只通过 C API（clEnqueueNDRangeKernel 等）调用，不能像 <<<>>> 那样直接在 C++ 中调用核函数。另外 NVIDIA Cg 已于 2012 年停止开发。

上图中 cudaDeviceSynchronize() 让 CPU 等 GPU 完成，作用类似 pthread_join()（见进程、线程与纤程）。

<<<G, N>>> 中 N 是输入数据的维度，即要 GPU 执行的循环次数，可以是一、二或三维；和 C/C++ 一样，多维数组本质上是按特殊方式索引的一维数组（[row][column] 等价于 [row*numColumns + column]）。
G 是线程组数（NVIDIA 称 thread block）。只有一个线程组的计算任务只能在一个计算单元 (CU) 上运行；CU 本质上是 GPU 里能执行一条指令流的核。G 越大，驱动越能把工作分散到多个 CU。

GPU 线程与线程组（卷 I 4.11.4）

GPU 核函数和 C++ 函数一样被编译成机器指令流，所以 GPU「线程」也是一条指令流；但 GPU 执行线程的方式不同，因此这个词含义略有不同。
计算单元像精简的 CPU 核：取指/译码单元、可能有缓存（L1、本地数据存储 LDS）、一个标量 ALU（可执行单通道指令）、一个大寄存器文件，通常约 4 个 SIMD 单元（作用相当于 CPU 的 VPU）。CU 不做推测执行或乱序执行，只是逐条执行指令。
以 AMD Radeon GCN 架构为例：SIMD 宽 16 通道。驱动把输入缓冲分成 64 个元素一块，每块在一个 CU 上调用一次核函数，称为一个 wavefront（NVIDIA 称 warp）。每条指令由 SIMD 单元锁步作用于 16 个元素；SIMD 内部是 4 级流水线、需要 4 个周期，为了不让流水线 4 个周期里闲 3 个，CU 对另外三组 16 个元素重复同一指令——所以 wavefront 是 64 个元素，而 SIMD 只有 16 宽。
因此 「GPU 线程」指的是一条 SIMD 通道，可理解为原来未向量化循环的一次迭代，或核函数的一次调用（一个输入、一个输出）。GPU 实际按 wavefront 一次处理 64 个元素只是实现细节；正因程序员不必关心具体 GPU 如何向量化，计算核函数和图形着色器才能写得可移植。

从 SIMD 到 SIMT（卷 I 4.11.4.1）

SIMT 一词强调 GPU 不只用 SIMD，还用一种抢占式多线程在 wavefront 之间分时。原因：指令一旦访存，SIMD 单元就要长时间等待内存控制器。为填补这些大延迟槽，SIMD 单元在多个 wavefront（来自同一个或许多不相关的着色器程序）之间切换：一个停顿就切到另一个，只要还有可运行的 wavefront，单元就保持忙碌。
GPU 切换频率极高，像 CPU 那样把寄存器存回内存太慢。所以每个 SIMD 单元有非常大的寄存器文件，物理寄存器数通常约为单个 wavefront 可用逻辑寄存器数的十倍，最多约十个 wavefront 的寄存器同时常驻，切换时完全不需要保存或恢复寄存器。

sequenceDiagram
  participant S as SIMD 单元
  participant A as wavefront A
  participant B as wavefront B
  S->>A: 执行指令
  A-->>S: 访存，停顿
  S->>B: 切换（寄存器常驻，无需保存）
  B-->>S: 访存，停顿
  S->>A: 数据已返回，继续

上图示意：一个 wavefront 停顿时 SIMD 单元立即转去执行另一个。

机器学习与张量核心（卷 I 4.11.5）

NVIDIA 新 GPU 的张量核心 (tensor core)（2017 年 Volta 架构首次引入）工作方式类似 CUDA 核心，但为机器学习（推理、预测、神经网络训练）优化。数学上张量是矩阵的推广；在 ML 中指算法要处理的任意 n 维数值表。
高速来自对少于 32 位格式的融合乘加 (MADD) 指令。书中列出的格式：FP64、FP32、FP16/bfloat16、Int8、Int4、Int1，以及 TF32：占 32 位内存，但只有 10 位尾数、8 位指数和 1 位符号，共用 19 位；比 FP32 最多快 8 倍（这是 NVIDIA 对 A100 的宣传口径：TF32 张量核心与同卡 FP32 CUDA 核心的峰值之比）。书中说张量核心靠「少于 32 位的格式」提速，但同时列出 FP64/FP32，两者并不矛盾：低位宽格式吞吐更高，FP64 张量核心主要用于 HPC。各格式按架构逐代加入（如 Int1/Int4 见于 Turing/Ampere，后续架构又有 FP8 等），并非每代都全支持。张量核心是精度换速度的取舍，可行是因为 ML 算法比图形应用更能容忍低精度（浮点格式见数值表示与浮点数）。

与其他来源的关系

本库已有 RTR4 的 GPU 并行架构与计算着色器。RTR4 说 NVIDIA warp 为 32 个线程、由 8 到 64 个核心以 SIMD 方式锁步执行；GEA 以 AMD GCN 为例，给出 SIMD 16 通道 × 4 周期流水线 = 64 元素 wavefront 的具体来由。两书描述一致，只是 GEA 补充了「为什么是 64」和「寄存器文件约十倍于单个 wavefront 所需」这两个细节。
延伸阅读（书中所列）：CUDA 入门、OpenCL 学习资源、HLSL 编程指南、OpenGL 着色语言入门、AMD RDNA 架构白皮书。