共享内存(Shared Memory)是 GPU 上由同一个线程块(Thread Block)内所有线程共同读写的一小块片上高速内存,它由程序员显式分配和管理,用来在计算过程中暂存并反复使用数据,从而避免反复访问显存。
把它想成一个小组的白板
想象一个项目组在开会:公司的档案库在楼下很远的地方(这就是全局显存,容量大,但取一次资料要等很久);每个人手里的小本子是私有的,只能自己看(这是寄存器);会议室墙上挂着一块白板,组里谁都能看、谁都能写,写上去的东西别人立刻就能用(这就是共享内存)。但白板只属于这个会议室,散会后就被擦掉,隔壁组也看不到。
GPU 的处境类似:一块显卡上有成千上万个线程,显存(Global Memory,一般是 HBM)容量大,但延迟高、带宽有限,是所有计算单元争抢的瓶颈。共享内存就在计算核心旁边,延迟低一到两个数量级。因此,做矩阵乘法、卷积、softmax、注意力这类“同一批数据要被反复使用”的计算时,常规做法是先把数据从显存搬进共享内存,让块内线程反复使用,最后再写回显存。这类手法常被称为分块(tiling)。
和相邻概念的区别
| 内存类型 | 谁能访问 | 由谁管理 | 特点 |
|---|---|---|---|
| 寄存器(Register) | 单个线程 | 编译器分配 | 最快、最稀缺 |
| 共享内存(Shared Memory) | 同一个线程块内的线程 | 程序员显式管理 | 快、容量小、可控 |
| 全局显存(Global Memory) | 所有线程 | 程序员 | 容量大、延迟高 |
| L1/L2 缓存(Cache) | 硬件按需缓存 | 硬件自动 | 不需管,但也控制不了 |
关键区别在于:缓存是硬件“猜”你要用什么,共享内存是你“指定”把什么放在手边。
用不好就成了瓶颈:bank 冲突
共享内存通常被划分成若干个 bank(常见实现是 32 个),硬件可以同时服务多个 bank 上的访问。一个 warp(通常是 32 个线程)同时访问共享内存时,如果各自落在不同 bank 上,一次就能完成;如果多个线程访问同一个 bank 里的不同地址,硬件只能排队,逐个处理,这就是 bank conflict(bank 冲突)。最坏情况下,访问会退化成几十次串行,共享内存带来的性能优势被抹平。如果多个线程读同一个 bank 的同一个地址,硬件会广播,这不算冲突。
一个经典场景是矩阵转置:把数据按行写进共享内存,再按列读出,而列方向的步长容易正好是 bank 数量的整数倍,结果所有线程都落在同一个 bank 上。常用办法是在数组里多补一列(padding)把地址错开。
对从业者的意义
共享内存是 GPU 性能优化中最“值钱”的部件之一:用得好,同样的显卡能跑出明显更高的吞吐;用得差,kernel 会卡在 bank 冲突或同步等待上。写 kernel 时要记住三件事:一是合并访问、按 bank 对齐;二是控制使用量,因为每个线程块能用的共享内存有限,占用太多会压低同时驻留的线程块数量,进而影响延迟隐藏;三是注意 __syncthreads() 这类块内同步的代价。此外,不同 GPU 架构对共享内存总量的规定不同,具体额度以厂商官方文档为准。
