光
矩阵乘法（GEMM）是理解 CUDA 优化最好的练习。本帖以 C = A × B 为例，梳理从朴素实现到分块优化的步骤，与 Shared Memory 基础帖互为补充。
一、朴素实现
每个线程计算 C 的一个元素，循环读取 A 的一行和 B 的一列。问题在于每个元素被重复从全局内存读取 N 次，算术强度很低，性能通常只有峰值的几个百分点。
二、Shared Memory 分块（tiling）
1. 每个 block 负责 C 的一个 TILE×TILE 子块（常用 16 或 32）；
2. 沿 K 维循环，每次把 A、B 各一个子块载入 shared memory；
3. __syncthreads() 后做子块乘加，再同步进入下一轮；
4. 全局内存访问量降为原来的 1/TILE。
三、进一步优化
1. 避免 bank conflict：A 子块按需转置存放或填充一列，让 warp 内访问落在不同 bank；
2. 寄存器分块：每线程计算 4×4 或 8×8 个输出，提升数据复用；
3. 向量化加载：使用 float4 一次读 16 字节；
4. 双缓冲：计算当前子块时预取下一个子块，隐藏访存延迟；
5. Tensor Core：使用 WMMA 或 MMA 指令，FP16/BF16 下吞吐提高数倍，需要按指定布局在 shared memory 中重排数据。
四、常见坑
1. 矩阵尺寸不是 TILE 整数倍时，边界判断遗漏导致越界；
2. 共享内存过大导致占用率下降，性能反而变差；
3. 自己写的 kernel 与 cuBLAS 差距大，生产环境建议优先用 cuBLAS/CUTLASS。
讨论：
1. 你们手写 GEMM 能做到 cuBLAS 的百分之多少？
2. Tensor Core 的数据重排开销你们怎么优化？
一、朴素实现
每个线程计算 C 的一个元素，循环读取 A 的一行和 B 的一列。问题在于每个元素被重复从全局内存读取 N 次，算术强度很低，性能通常只有峰值的几个百分点。
二、Shared Memory 分块（tiling）
1. 每个 block 负责 C 的一个 TILE×TILE 子块（常用 16 或 32）；
2. 沿 K 维循环，每次把 A、B 各一个子块载入 shared memory；
3. __syncthreads() 后做子块乘加，再同步进入下一轮；
4. 全局内存访问量降为原来的 1/TILE。
三、进一步优化
1. 避免 bank conflict：A 子块按需转置存放或填充一列，让 warp 内访问落在不同 bank；
2. 寄存器分块：每线程计算 4×4 或 8×8 个输出，提升数据复用；
3. 向量化加载：使用 float4 一次读 16 字节；
4. 双缓冲：计算当前子块时预取下一个子块，隐藏访存延迟；
5. Tensor Core：使用 WMMA 或 MMA 指令，FP16/BF16 下吞吐提高数倍，需要按指定布局在 shared memory 中重排数据。
四、常见坑
1. 矩阵尺寸不是 TILE 整数倍时，边界判断遗漏导致越界；
2. 共享内存过大导致占用率下降，性能反而变差；
3. 自己写的 kernel 与 cuBLAS 差距大，生产环境建议优先用 cuBLAS/CUTLASS。
讨论：
1. 你们手写 GEMM 能做到 cuBLAS 的百分之多少？
2. Tensor Core 的数据重排开销你们怎么优化？
