💡 Shared Memory Tiling
问题: naive 矩阵乘法中,Global Memory 每个元素被读取
方案: 切成 TILE 小块,每次加载一个 tile 到 Shared Memory 后反复读取
A/B 和 As/Bs 的关系:
As 是从 A 复制的 TILE×TILE 子块
Bs 是从 B 复制的 TILE×TILE 子块
Block 内线程协作加载 → 在 Shared Memory 中反复读取 → 计算 →
为什么 C 始终高亮同一块?
C 的一个子块需要 K 维上的所有 tile 累加才能完成:
两个 tile 沿 K 维度接力,部分和累加到同一个 C 子块。
C 全矩阵共 4 个子块 (2×2),每个子块都需要遍历全部 TILES 个 tile。
Shared Memory ~20ns vs Global Memory ~600ns — 快 30 倍
参考: 矩阵乘法的 CUDA 实现、优化及性能分析
问题: naive 矩阵乘法中,Global Memory 每个元素被读取
K 次方案: 切成 TILE 小块,每次加载一个 tile 到 Shared Memory 后反复读取
A/B 和 As/Bs 的关系:
As 是从 A 复制的 TILE×TILE 子块
Bs 是从 B 复制的 TILE×TILE 子块
Block 内线程协作加载 → 在 Shared Memory 中反复读取 → 计算 →
__syncthreads() → 加载下一 tile为什么 C 始终高亮同一块?
C 的一个子块需要 K 维上的所有 tile 累加才能完成:
C[0:4,0:4] = A[0:4,0:4]×B[0:4,0:4] (Tile 0) + A[0:4,4:8]×B[4:8,0:4] (Tile 1)两个 tile 沿 K 维度接力,部分和累加到同一个 C 子块。
C 全矩阵共 4 个子块 (2×2),每个子块都需要遍历全部 TILES 个 tile。
Shared Memory ~20ns vs Global Memory ~600ns — 快 30 倍
参考: 矩阵乘法的 CUDA 实现、优化及性能分析