📦 Shared Memory Tiling — 矩阵乘法数据流

模块 2 — GPU 硬件架构与 CUDA 编程入门  |  PPT 第 35-36 页  |  Demo: M=N=K=8, TILE=4

💡 Shared Memory Tiling

问题: 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 实现、优化及性能分析