《Programming Massively Parallel Processors》第四版 - 学习笔记与练习
本章系统梳理并行直方图计算及其优化技术：
- 原子操作（Atomic Operations）
- 私有化（Privatization）
- 共享内存优化
- 线程粗化（Thread Coarsening）
- 内存合并访问（Coalesced Memory Access）
相关博客笔记：第九章：并行直方图
实现5种直方图 kernel，对应书中图9.6、9.10、9.14。
代码位置：Exercise01/
实现列表：
| 实现 | 书中对应 | 特点 |
|---|---|---|
histogram_sequential |
- | CPU参考实现 |
histogram_parallel_basic |
图9.6 | 全局内存原子操作 |
histogram_parallel_private_shared |
图9.10 | 共享内存私有化 |
histogram_parallel_coarsening |
图9.14 | 线程粗化 |
histogram_parallel_coalesced |
- | 粗化 + 内存合并 |
核心代码：
// 私有化核心思想（图9.10）
__shared__ unsigned int histo_s[NUM_BINS]; // 私有直方图
// 初始化私有直方图
for (int bin = threadIdx.x; bin < NUM_BINS; bin += blockDim.x) {
histo_s[bin] = 0u;
}
__syncthreads();
// 统计到私有直方图（共享内存原子操作）
atomicAdd(&histo_s[alphabet_position / BIN_SIZE], 1);
__syncthreads();
// 合并到全局直方图
atomicAdd(&histo[bin], histo_s[bin]);cd Exercise01
make
make run================================================================
第九章：并行直方图
Histogram Computation - 5 Implementations
================================================================
配置:
数据长度: 10000000 (10 M)
BIN_SIZE: 4 (每桶字母数)
NUM_BINS: 7
=== 正确性验证 ===
1. CPU 顺序实现... 完成
2. 基础并行 (图9.6)... ✅ 结果正确！
3. 私有化共享内存 (图9.10)... ✅ 结果正确！
4. 线程粗化 (图9.14)... ✅ 结果正确！
5. 线程粗化 + 内存合并... ✅ 结果正确！
题目： 假设 DRAM 系统中每次原子操作的总延迟为 100 ns。对于同一全局内存变量的原子操作，我们能获得的最大吞吐量是多少？
解答：
每次操作需要 100 ns（即 100 × 10^-9 s/op）。一秒内可以执行：
1s / (100 × 10^-9 s/op) = 10^7 op/s
所以最大吞吐量是 1000 万次操作/秒。
题目： 对于支持 L2 缓存原子操作的处理器，假设每次原子操作在 L2 缓存中需要 4 ns，在 DRAM 中需要 100 ns。假设 90% 的原子操作命中 L2 缓存。对于同一全局内存变量的原子操作，大致吞吐量是多少？
解答：
平均每次操作耗时：
0.9 × 4 ns + 0.1 × 100 ns = 3.6 ns + 10 ns = 13.6 ns
吞吐量：
1s / (13.6 × 10^-9 s/op) ≈ 73.5 × 10^6 op/s
所以吞吐量约为 7350 万次操作/秒。
注：由于操作的是同一全局变量，这些原子操作必须串行执行。
题目： 在练习 1 中，假设 kernel 每次原子操作执行 5 次浮点运算。受原子操作吞吐量限制，kernel 执行的最大浮点吞吐量是多少？
解答：
根据练习 1，原子操作吞吐量为 1000 万次/秒。浮点吞吐量：
(5 FP ops/Atomic op) × (10^7 Atomic ops/s) = 5 × 10^7 FLOPS = 50 MFLOPS
这远低于现代 GPU 的浮点性能（通常是 TFLOPS 级别），相差多个数量级。
题目： 在练习 1 中，假设我们将全局内存变量私有化到共享内存，共享内存访问延迟为 1 ns。所有原来的全局内存原子操作都转换为共享内存原子操作。假设将私有化变量累加到全局变量的额外全局内存原子操作增加了 10% 的总执行时间。假设 kernel 每次原子操作执行 5 次浮点运算。受原子操作吞吐量限制，最大浮点吞吐量是多少？
解答：
共享内存原子操作吞吐量：
1s / (1 ns/op × 1.1) = 1s / (1.1 × 10^-9 s/op) ≈ 9.1 × 10^8 op/s
浮点吞吐量：
5 × 9.1 × 10^8 = 4.55 × 10^9 FLOPS = 4.55 GFLOPS
虽然比练习 3 提高了约 90 倍，但仍比 H100 等现代 GPU 的理论峰值（约数十 TFLOPS）低约 5 个数量级。
题目： 要执行原子加操作，将整数变量 Partial 的值加到全局内存整数变量 Total，应该使用以下哪个语句？
a. atomicAdd(Total, 1);
b. atomicAdd(&Total, &Partial);
c. atomicAdd(Total, &Partial);
d. atomicAdd(&Total, Partial);
解答：
atomicAdd 的函数签名是：
int atomicAdd(int* address, int val);第一个参数是要更新的变量的地址，第二个参数是要加的值。
答案是 D — atomicAdd(&Total, Partial)，即 Total 的地址和 Partial 的值。
题目： 考虑一个直方图 kernel，处理 524,288 个元素的输入，产生 128 个桶的直方图。kernel 配置为每个块 1024 个线程。
解答：
在图9.6 的 kernel 中，每个线程对全局内存执行一次 atomicAdd（第06行）。由于启动了 524,288 个线程，全局内存原子操作次数为 524,288 次。
解答：
在图9.10 中，只有在将共享内存私有直方图合并到全局内存时才执行全局原子操作。每个 Block 对每个桶执行一次 atomicAdd（如果 threadIdx.x < NUM_BINS），即每个 Block 最多 128 次。
Block 数量：
(524,288 + 1024 - 1) / 1024 = 512 个 Block
全局内存原子操作次数：
512 × 128 = 65,536 次
解答：
使用粗化因子 4 后，Block 数量减少：
512 / 4 = 128 个 Block
合并机制不变（每个 Block 对每个桶执行一次），全局内存原子操作次数：
128 × 128 = 16,384 次
- CUDA Toolkit: 11.0 或更高版本
- 编译器: GCC 7.5+ / Visual Studio 2019+ + NVCC
- GPU: 支持 CUDA 的 NVIDIA 显卡（计算能力 3.5+）
- 理解原子操作：读-改-写的原子性
- 私有化策略：减少争用是关键
- 共享内存 vs 全局内存：原子操作性能差 ~20 倍
- 线程粗化：减少 Block 数量和合并开销
- 内存合并：Grid-Stride Loop 是通用优化模式
完成本章学习后，继续学习：
- 第十章：归约
- 第十一章：前缀和
- 第十二章：合并
- PMPP 第四版 Chapter 9
- 第九章：并行直方图
学习愉快！ 🎓