《Programming Massively Parallel Processors》第四版 - 学习笔记与练习
本章系统梳理 GPU 性能优化技术：
- 内存合并（Memory Coalescing）
- 分区露营（Partition Camping）
- 指令混合与吞吐量
- 线程粗化（Thread Coarsening）
- 资源平衡与占用率权衡
- Warp 执行效率
- 数据预取与双缓冲
相关博客笔记：第六章：性能方面的考虑
本练习实现对应图 6.4 设计的矩阵乘法 kernel，演示角转换（Corner Turning）技术。
代码位置：Exercise01/
功能：
- 行主序 Tiled 版本：标准的 Tiled 矩阵乘法实现
- 列主序 Tiled 版本：使用角转换技术，按列访问 N 矩阵以保持合并访问
核心优化：
// 列主序 Tiled 矩阵乘法 Kernel（Corner Turning）
__global__ void TiledMatrixMulKernelColMajor(float* M, float* N, float* P, int m, int n, int o) {
__shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
__shared__ float Nds[TILE_WIDTH][TILE_WIDTH];
// ... 省略变量定义 ...
for (int ph = 0; ph < numTiles; ++ph) {
// 加载 M（行主序，正常访问）
Mds[ty][tx] = M[row * n + ph * TILE_WIDTH + tx];
// 加载 N（列主序：按列访问以保持合并）
// N 存储为列主序：N[col * n + row]
Nds[ty][tx] = N[col * n + (ph * TILE_WIDTH + ty)];
__syncthreads();
for (int k = 0; k < TILE_WIDTH; ++k) {
Pvalue += Mds[ty][k] * Nds[k][tx];
}
__syncthreads();
}
}cd Exercise01
make
make run========================================
第六章：性能方面的考虑
Column-Major Matrix Multiplication
========================================
矩阵大小: 4 × 3 × 2
行主序矩阵乘法结果:
列主序矩阵乘法结果:
✅ 两种方法结果一致！
本代码对比三种矩阵乘法实现的性能差异，演示 Thread Coarsening 优化技术。
代码位置：Exercise02/
三种实现对比：
| 版本 | 特点 | 算术强度 |
|---|---|---|
| 朴素版本 | 直接全局内存访问 | 0.25 OP/B |
| Tiled (32×32) | 共享内存优化 | 8 OP/B |
| Tiled + Coarsening×4 | 每线程计算 4 个元素 | 12.8 OP/B |
核心优化：
// Thread Coarsening：每个线程计算 COARSE_FACTOR 个输出元素
float Pvalue[COARSE_FACTOR] = {0.0f}; // 4 个累加器
for (int ph = 0; ph < numTiles; ++ph) {
// M 的 Tile 只加载一次
Mds[ty][tx] = M[row * n + ph * TILE_WIDTH + tx];
// 但 N 的 Tile 加载 COARSE_FACTOR 次，计算 4 个输出
for (int c = 0; c < COARSE_FACTOR; ++c) {
int col = colStart + c * TILE_WIDTH;
Nds[ty][tx] = N[(ph * TILE_WIDTH + ty) * o + col];
__syncthreads();
for (int k = 0; k < TILE_WIDTH; ++k) {
Pvalue[c] += Mds[ty][k] * Nds[k][tx];
}
__syncthreads();
}
}cd Exercise02
make
make run========================================
第六章：性能方面的考虑
Thread Coarsening Matrix Multiplication
========================================
矩阵大小: 1024 × 1024 × 1024
测试迭代次数: 10
Thread Coarsening Factor: 4
=== 正确性验证 ===
✅ 所有方法结果一致！
=== 性能测试 ===
朴素矩阵乘法: 57.273 ms
Tiled 矩阵乘法: 53.636 ms (1.07x vs naive)
Tiled + Coarsening: 48.123 ms (1.19x vs naive, 1.11x vs tiled)
题目： 编写一个对应图 6.4 设计的矩阵乘法 kernel 函数。
解答：
图 6.4 展示的是处理列主序（Column-Major）矩阵的 Tiled 矩阵乘法。关键点是使用**角转换（Corner Turning）**技术：不按行访问 N 矩阵（会导致非合并访问），而是按列访问以保持合并访问，然后在共享内存中重新排列数据。
__global__ void TiledMatrixMulKernelColMajorOrder(float *M, float *N, float *P,
int m, int n, int o) {
// 角转换：按列访问 N 矩阵（合并访问）
// 共享内存用于重排数据
__shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
__shared__ float Nds[TILE_WIDTH][TILE_WIDTH];
int by = blockIdx.y;
int bx = blockIdx.x;
int ty = threadIdx.y;
int tx = threadIdx.x;
int row = by * TILE_WIDTH + ty;
int col = bx * TILE_WIDTH + tx;
float PValue = 0;
for (int ph = 0; ph < (n + TILE_WIDTH - 1) / TILE_WIDTH; ph++) {
// 加载 M（行主序，正常访问）
if (row < m && (ph * TILE_WIDTH + tx) < n)
Mds[ty][tx] = M[row * n + ph * TILE_WIDTH + tx];
else
Mds[ty][tx] = 0.0f;
// 加载 N（列主序：按列访问以保持合并）
// N 在内存中是列主序存储：N[col][row] = N[col * n + row]
if ((ph * TILE_WIDTH + ty) < n && (col < o))
Nds[ty][tx] = N[col * n + (ph * TILE_WIDTH + ty)];
else
Nds[ty][tx] = 0.0f;
__syncthreads();
for (int k = 0; k < TILE_WIDTH; k++) {
PValue += Mds[ty][k] * Nds[k][tx];
}
__syncthreads();
}
if (row < m && col < o)
P[row * o + col] = PValue;
}关键点：
N[col * n + (ph * TILE_WIDTH + ty)]：按列访问 N，相邻线程访问相邻内存地址- 共享内存
Nds存储后，计算时按正常方式访问
题目： 对于 Tiled 矩阵乘法，在 BLOCK_SIZE 的可能取值范围内，哪些值能完全避免对全局内存的非合并访问？（只考虑方形块）
解答： BLOCK_SIZE 必须是 32 的倍数（32、64 等）
分析：
合并访问要求 Warp 中的 32 个连续线程访问连续的内存地址。
- 如果 BLOCK_SIZE < 32，一个 Warp 会跨越多行
- 例如 BLOCK_SIZE = 16：Warp 的线程 0-15 访问第 0 行，线程 16-31 访问第 1 行
- 这导致部分访问是非合并的
验证：
| BLOCK_SIZE | Warp 划分 | 合并情况 |
|---|---|---|
| 16 | 1 个 Warp 跨 2 行 | ❌ 非合并 |
| 32 | 1 个 Warp = 1 行 | ✅ 完全合并 |
| 64 | 2 个 Warp = 2 行 | ✅ 完全合并 |
实践限制：
- 共享内存大小限制通常使 BLOCK_SIZE > 64 不可行
- 32 是最常用的选择
题目： 考虑以下 CUDA kernel，判断每个内存访问是合并的、非合并的、还是不适用合并概念：
__global__ void foo_kernel(float* a, float* b, float* c, float* d, float* e) {
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
__shared__ float a_s[256];
__shared__ float bc_s[4*256];
a_s[threadIdx.x] = a[i]; // 第 5 行
for(unsigned int j = 0; j < 4; ++j) {
bc_s[j*256 + threadIdx.x] =
b[j*blockDim.x*gridDim.x + i] + c[i*4 + j]; // 第 7 行
}
__syncthreads();
d[i + 8] = a_s[threadIdx.x]; // 第 10 行
e[i*8] = bc_s[threadIdx.x*4]; // 第 11 行
}解答：
解答： ✅ 合并访问
a[blockIdx.x*blockDim.x + threadIdx.x]：相邻线程的 threadIdx.x 连续，访问连续内存地址。
解答： ⚪ 不适用
a_s 是共享内存，合并访问概念只适用于全局内存。
解答： ✅ 合并访问
b[j*blockDim.x*gridDim.x + blockIdx.x*blockDim.x + threadIdx.x]：对于固定的 j，相邻线程访问连续地址。
解答： ❌ 非合并访问
c[(blockIdx.x*blockDim.x + threadIdx.x)*4 + j]：相邻线程之间跳跃 4 个元素，步长为 4。
解答： ⚪ 不适用
bc_s 是共享内存。
解答： ⚪ 不适用
a_s 是共享内存。
解答： ✅ 合并访问
d[blockIdx.x*blockDim.x + threadIdx.x + 8]：相邻线程访问连续地址（偏移 8 不影响合并）。
解答： ⚪ 不适用
bc_s 是共享内存。
解答： ❌ 非合并访问
e[(blockIdx.x*blockDim.x + threadIdx.x)*8]：相邻线程之间跳跃 8 个元素，每个线程访问不同的内存段。
题目： 计算以下矩阵乘法 kernel 的浮点运算与全局内存访问比（OP/B）：
假设矩阵 M 大小为 (m, n)，矩阵 N 大小为 (n, o)，使用 float32（4 字节）。
解答： 0.25 OP/B
分析：
对于结果矩阵的每个元素：
- 加载 M 的一行：n 次内存读取
- 加载 N 的一列：n 次内存读取
- 执行乘法：n 次
- 执行加法：n-1 ≈ n 次
内存访问: 2n × 4 bytes = 8n bytes
浮点运算: 2n (n 次乘法 + n 次加法)
比值: 2n / 8n = 0.25 OP/B
解答： 8 OP/B
分析：
使用 Tiling 后，每个线程只需从全局内存加载 n/32 个元素（M 和 N 各 n/32）：
- 同一 Tile 被 32 个线程共享使用
- 每个元素从全局内存只读 1 次，在共享内存被使用 32 次
内存访问: 2 × (n/32) × 4 bytes = n/4 bytes
浮点运算: 2n
比值: 2n / (n/4) = 8 OP/B
相比朴素版本提升 32 倍。
解答： 12.8 OP/B
分析：
Thread Coarsening 让每个线程计算 4 个输出元素：
- M 的 Tile 被复用 4 次：n/32/4 = n/128 次加载
- N 的 Tile 正常加载：n/32 次加载
- 计算量不变：2n 操作
内存访问: (n/128 + n/32) × 4 bytes = (n/128 + 4n/128) × 4 = 5n/32 bytes
浮点运算: 2n
比值: 2n / (5n/32) = 2n × 32 / 5n = 12.8 OP/B
相比 Tiled 版本提升 1.6 倍。
- CUDA Toolkit: 11.0 或更高版本
- 编译器: GCC 7.5+ / Visual Studio 2019+ + NVCC
- GPU: 支持 CUDA 的 NVIDIA 显卡（计算能力 3.5+）
- 优先保证合并访问：检查全局内存访问模式
- 使用 Profiler：
ncu比猜测更有效 - 权衡占用率：计算密集型可接受低占用率
- Grid-stride loop：灵活处理任意大小数据
- 实测验证：不同 GPU 性能特性不同
完成本章学习后，继续学习：
- 第七章：卷积
- 第八章：模板
- 第九章：并行直方图
- PMPP 第四版 Chapter 6
- 第六章：性能方面的考虑
学习愉快！ 🎓