《Programming Massively Parallel Processors》第四版 - 学习笔记与练习
本章系统梳理归并（Merge）操作及其 CUDA 优化技术：
- Co-rank 函数（二分搜索定位）
- 顺序归并算法
- 基础并行归并（全局内存）
- 分块并行归并（共享内存优化）
相关博客笔记：第十二章：归并
实现多种归并 kernel，对应书中图12.9、12.11-12.13。
代码位置：Exercise01/
实现列表：
| 实现 | 书中对应 | 特点 |
|---|---|---|
merge_sequential |
- | CPU参考实现 |
merge_basic_gpu |
图12.9 | 基础并行，每线程独立 co-rank |
merge_tiled_gpu |
图12.11-12.13 | 分块优化，共享内存 |
核心代码：
// Co-rank 函数核心思想
__host__ __device__ int co_rank(int k, float* A, int m, float* B, int n) {
int i = min(k, m);
int j = k - i;
int i_low = max(0, k - n);
int j_low = max(0, k - m);
while (active) {
if (i > 0 && j < n && A[i-1] > B[j]) {
// i 太大，需要减少
delta = cdiv(i - i_low, 2);
i -= delta; j += delta;
} else if (j > 0 && i < m && B[j-1] >= A[i]) {
// i 太小，需要增加
delta = cdiv(j - j_low, 2);
i += delta; j -= delta;
} else {
break; // 找到正确位置
}
}
return i;
}cd Exercise01
make
make run================================================================
第十二章：归并
Merge Operations - Multiple Implementations
================================================================
配置:
数组 A 长度: 10283
数组 B 长度: 131131
合并结果长度: 141414
=== 正确性验证 ===
1. CPU 顺序归并... ✅ 结果有序
2. 基础并行归并 (图12.9)... ✅ 结果正确！
3. 分块并行归并 (图12.11-12.13)... ✅ 结果正确！
题目： 假设需要合并两个列表 A=(1, 7, 8, 9, 10) 和 B=(7, 10, 10, 12)。C[8] 的协同排名（co-rank）值是多少？
解答：
首先得出合并结果：[1, 7, 7, 8, 9, 10, 10, 10, 12]
C[8] = 12，这是 B 中的最后一个元素。
执行 co_rank(8, A, 5, B, 4)：
初始:
i = min(8, 5) = 5
j = k - i = 8 - 5 = 3
i_low = max(0, 8-4) = 4
j_low = max(0, 8-5) = 3
检查条件:
i > 0 ✅ && j < n ✅ && A[4]=10 > B[3]=12 ❌ → 第一个 if 不满足
j > 0 ✅ && i < m ❌ (5 < 5) → 第二个 if 不满足
结束循环，返回 i = 5
答案：i = 5, j = 3
这意味着 C[8:] 从 A[5:] 和 B[3:] 开始，即不取 A 的元素（已用完），只取 B[3] = 12。
题目： 完成图12.6中线程2的协同排名函数计算。
解答：
线程2从 k=6 开始，计算 co_rank(6, A, 5, B, 4)：
初始:
i = min(6, 5) = 5
j = k - i = 6 - 5 = 1
i_low = max(0, 6-4) = 2
j_low = max(0, 6-5) = 1
检查条件:
i > 0 ✅ && j < n ✅ && A[4]=10 > B[1]=10 ❌ → 不满足
j > 0 ✅ && i < m ❌ → 不满足
结束循环，返回 i = 5
答案：i = 5
这意味着 C[6:] 从 A 取0个元素（已用完），从 B[1:] 开始取3个元素。
题目： 对于图12.12中加载 A 和 B 分块的 for 循环，添加 co_rank 函数调用，使得只加载当前 while 循环迭代中会被消耗的 A 和 B 元素。
解答：
while(counter < total_iteration){
// 确定需要多少 A 元素才能合并出 tile_size 个结果
int tileA = co_rank(tile_size,
A + A_curr + A_consumed, A_length - A_consumed,
B + B_curr + B_consumed, B_length - B_consumed);
int tileB = tile_size - tileA;
// 只加载需要的 A 元素
for(int i = 0; i < tileA; i += blockDim.x) {
if (i + threadIdx.x < tileA) {
A_S[i + threadIdx.x] = A[A_curr + A_consumed + i + threadIdx.x];
}
}
// 只加载需要的 B 元素
for(int i = 0; i < tileB; i += blockDim.x) {
if(i + threadIdx.x < tileB) {
B_S[i + threadIdx.x] = B[B_curr + B_consumed + i + threadIdx.x];
}
}
__syncthreads();优化效果：减少不必要的全局内存加载。
题目： 考虑对两个大小为 1,030,400 和 608,000 的数组进行并行归并。假设每个线程归并8个元素，块大小为1024。
结果数组长度：1,030,400 + 608,000 = 1,638,400 元素
每个线程归并8个元素，总共需要 1,638,400 / 8 = 204,800 个线程。
每个线程调用 co_rank 两次（起始和结束位置）。
答案：204,800 个线程执行全局内存二分搜索。
每个块处理 8 × 1024 = 8,192 个元素。
块数量：1,638,400 / 8,192 = 200 个块。
每个块只有线程0执行全局内存 co_rank（两次），其他使用共享内存。
答案：200 个线程执行全局内存二分搜索。
每个块有 1024 个线程，每个线程执行共享内存 co_rank 三次（行41、43、50）。
200 个块 × 1024 线程 = 204,800 个线程执行共享内存二分搜索。
- CUDA Toolkit: 11.0 或更高版本
- 编译器: GCC 7.5+ / Visual Studio 2019+ + NVCC
- GPU: 支持 CUDA 的 NVIDIA 显卡（计算能力 3.5+）
- 理解 co-rank：归并的关键是找到正确的分割点
- 二分搜索：O(log n) 复杂度，理解边界条件
- 共享内存优化：Block 级别只需2次全局 co-rank
- 负载均衡：通过 co-rank 实现均匀的工作分配
- 性能调优实践：使用 Nsight Compute 分析内存访问模式，对于不同大小的输入数组，比较基础归并和分块归并的性能差异，理解何时共享内存优化最有效
完成本章学习后，继续学习：
- 第十三章：排序
- 第十四章：稀疏矩阵计算
- 第十五章：图遍历
- PMPP 第四版 Chapter 12
- 第十二章：归并
学习愉快！ 🎓