很多人以为GPU性能上不去是“算力不够”，但调过几次AI模型和CUDA内核之后就会发现，大部分瓶颈其实卡在“数据喂不进去”这一步——GPU的算力再强，内存访问模式不合理，运算单元也只能饿着肚子等数据。这个系列到第七篇，我把GPU内存访问模式这块硬骨头彻底啃了一遍，这篇笔记就聊聊怎么“打通访存任督二脉”，把显存带宽吃干榨净。
这篇内容主要围绕GPU内存访问模式的原理、实际代码层面的优化手法、以及如何用工具定位访存瓶颈来展开。适合正在做AI推理加速、CUDA内核优化、或者踩过“明明算力不低但程序跑不快”这种坑的同学。文章不会讲太多玄乎的理论，更多是实际调优时能直接上手用的东西。
1. GPU访存：算力的隐形天花板
1.1 算力和带宽的失衡是原罪
先摆一组数据。以NVIDIA A100为例，它的FP16稠密算力大约312 TFLOPS，显存带宽大约1.6TB/s。按照“每读一个数至少做一次乘加”的理想模型，要喂饱312 TFLOPS的算力，每秒钟需要读取至少156万亿次FP16数据。但1.6TB/s带宽只够读大约800万亿字节，折合成FP16只有400万亿个数，和156万亿差距不小。更麻烦的是，这只是理论值，实际代码里又有索引计算、分支、依赖等待这些开销，访存带宽被浪费的比例比想象中高得多。
这个失衡的后果是：你做任何矩阵运算、卷积、Transformer推理，只要数据复用率不够，计算单元大部分时间都在空转。判断一个内核是否“访存受限”，业界有个很实用的指标叫算术强度——总计算量除以总访存量，单位是FLOP/Byte。如果这个值低于设备的平衡点（A100大约在200 FLOP/Byte上下），瓶颈就在访存；高于平衡点，核心就是算力受限。这个判断直接决定了你优化的大方向：是减少访存，还是优化计算。
我自己调过一个小型卷积核，一开始算术强度只有40多，怎么改线程排布都提升不大。后来把计算重排、做了数据复用，算术强度拉到150以上，速度一下快了将近3倍。所以内存优化最核心的一句话是：降低总访存量，或者提高每次访存带来的计算量，两者至少占一个。
1.2 延迟隐藏并不意味着内存不重要
GPU解决内存延迟的方式是多线程并发隐藏延迟，而不是像CPU那样依赖大缓存和分支预测。当一个线程在等数据时，硬件立刻切换到另一组线程执行计算。听起来似乎访存延迟不重要了，但有一个硬限制：每秒钟所有线程等数据的总等待时间必须小于总执行时间。如果你有100个线程，每个线程花70%的时间在等待数据，那么最多只能隐藏70%的延迟，剩下的30%就是实实在在的空洞。
这也是为什么占用率（occupancy）如此重要。占用率过低，线程数量不够淹没延迟，访存慢的问题就原形毕露。占用率过高也可能带来副作用，比如寄存器溢出和缓存抖动。实战中我一般通过调整block大小和寄存器用量，把占用率控制在50%~75%之间，再配合访存优化，往往能达到不错的平衡。
2. 内存访问模式的底层逻辑
2.1 合并访问：GPU访存的命根子
GPU从全局内存读数据，最小的传输单位不是单个字节，而是内存事务。以现代NVIDIA架构为例，一次事务的粒度通常是32字节或64字节，而且硬件按128字节的cache line来管理缓存。如果一个warp（32个线程）同时访问的地址恰好落在一个连续的128字节段内，硬件就能通过一次事务把数据全部取回，这就是合并访问（coalesced access）。反之，如果32个线程各跳各的，访问地址散布在不同cache line上，硬件就要拆成多次事务，访存效率呈级数下降。
生活化的类比是快递分拣。32个人到仓库取各自订购的商品，如果32个人的货刚好码在一个货板上，分拣员一趟就能全搬出来；如果每个人订的东西分散在不同货架，分拣员就得跑好几趟，表面上看每个人的取货时间没变，但整个仓库的出货效率被拖垮了。GPU访存也是同一个道理——带宽资源是共享的，一次能搬多少有用的数据，决定了吞吐量。
细心的同学应该能发现，合并访问关心的是“同一warp内32个线程的地址分布”，而不是“整个block的访问模式”。所以写着CUDA代码时，永远要问自己：一个warp的线程，它们此刻访问的内存地址连续吗？
2.2 对齐：从cache line层面理解访存
合并访问之外，对齐也常被人忽略。硬件读取128字节cache line时，如果数据没有按128字节边界对齐，那么一次请求可能要跨越两个cache line，白白多读一段无用数据。CUDA里可以用__align__在结构体上指定对齐，也可以在分配内存时用cudaMalloc（它天然对齐到至少256字节）。对于自管理的内存池，我习惯手动对齐到512字节，虽然多花几个字节，但后续做向量化读取时确实省心。
还要注意，合并访问是“连续”，对齐是“边界”，两者互相配合。一个最简单的例子：一个float数组，第0~31个元素给warp 0读，第32~63给warp 1读，只要保证warp的起始元素是32的倍数（同时保证起始地址是128字节对齐），每个warp就能吃满一次cache line事务。写kernel时用blockIdx.x * blockDim.x + threadIdx.x这种经典索引，天然满足对齐和合并。
2.3 访存模式的量化评估方法
判断一段代码的访存模式好不好，不能靠肉眼。我常用的方式是看三个核心指标：全局内存吞吐率（global memory throughput）、cache命中率（L1/L2 hit rate）、以及内存事务数（memory transactions）。NVIDIA Nsight Compute里可以直接看到这些数据，尤其是Memory Throughput和L1/TEX Cache Hit Rate两项，基本能还原内核真实的访存行为。
对比时注意，事务数才是关键。有时候吞吐率看着很高，但因为大量数据被重复读取，实际有用的字节数可能只有一小部分。Nsight Compute里有一个Memory Workload Analysis的视图，会分读写分别列出sectors和transactions，一眼就能看出是不是在做无用功。后面我会专门讲怎么用这个分析问题。
3. 矩阵转置实战：从非合并到合并访问
3.1 一个经典的访存反面教材
矩阵转置是访存优化的经典教学案例。朴素写法一般长这样：
__global__ void transpose_naive(const float* in, float* out, int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (x < width && y < height) { out[x * height + y] = in[y * width + x]; } }这个kernel的索引是线程坐标（x, y）直接映射到输入矩阵的(y, x)。当同一个warp的线程沿着x方向展开时，x连续，但in的下标是按y * width + x算的，注意这里y固定、x连续，in的访问是连续的；问题出在写回out时，out的下标是x * height + y，同一个warp里x连续，导致y方向的步长是height，所以写回是完全非合并的。反过来如果交换x和y的映射，读又变成非合并。也就是说，朴素转置始终有一半的访存是低效的。
在我的测试里（GTX 1080，矩阵4096x4096，float），这个朴素版本的全局内存吞吐只有大约80GB/s，远低于理论带宽。原因是写回时每个warp要触发几十个事务，把大量cache line都污染了，整体表现惨不忍睹。
3.2 用共享内存做分块转置
想同时保证读和写都合并，就得引入共享内存，把矩阵切成小块（tile），先将数据按合并方式读入共享内存，再从共享内存按合并方式写出去。共享内存没有cache line的概念，按任意模式访问代价都低很多。
经典的分块转置代码如下：
#define TILE_SIZE 32 __global__ void transpose_tiled(const float* in, float* out, int width, int height) { __shared__ float tile[TILE_SIZE][TILE_SIZE + 1]; int x = blockIdx.x * TILE_SIZE + threadIdx.x; int y = blockIdx.y * TILE_SIZE + threadIdx.y; if (x < width && y < height) { tile[threadIdx.y][threadIdx.x] = in[y * width + x]; } __syncthreads(); int xT = blockIdx.y * TILE_SIZE + threadIdx.x; int yT = blockIdx.x * TILE_SIZE + threadIdx.y; if (xT < height && yT < width) { out[yT * height + xT] = tile[threadIdx.x][threadIdx.y]; } }注意我故意把共享内存定义成[TILE_SIZE][TILE_SIZE + 1]，而不是[TILE_SIZE][TILE_SIZE]。这是为了避免bank conflict，后面还会细说。这里先理解整体流程：读取in的时候，每个warp访问连续地址，合并；写回out之前，先从共享内存中按转置位置取数据，此时共享内存的访问模式虽然有些交错，但由于加了一列padding，实际冲突被压到很低；最终写回时又恢复合并访问。
这一版在我的测试里能跑到约350GB/s，比朴素版提升了4倍多。还没上更激进的向量化，只是纯靠访存模式优化，效果已经很显著。
3.3 向量化读取再进一步
如果还想再压榨一点，可以用float4向量化读取。一次读4个float，相当于把内存事务的利用率推到极限。这时候要注意线程与数据的映射：假设宽度为width，向量化后逻辑上宽度变成width/4，下标计算要小心对齐。
__global__ void transpose_vec4(const float4* in, float4* out, int width, int height) { __shared__ float4 tile[TILE_SIZE][TILE_SIZE + 1]; int x = blockIdx.x * TILE_SIZE + threadIdx.x; int y = blockIdx.y * TILE_SIZE + threadIdx.y; if (x < width && y < height) { tile[threadIdx.y][threadIdx.x] = in[y * width + x]; } __syncthreads(); int xT = blockIdx.y * TILE_SIZE + threadIdx.x; int yT = blockIdx.x * TILE_SIZE + threadIdx.y; if (xT < height && yT < width) { out[yT * height + xT] = tile[threadIdx.x][threadIdx.y]; } }这里把原始数据当作float4数组处理，宽度也相应除以4。实测在4096x4096 float矩阵上，这版能到约480GB/s，已经比较接近硬件上限。注意向量化要求起始地址16字节对齐，用cudaMalloc拿到的内存都满足这个条件，自己写内存池时必须小心。
4. 更细致的访存优化维度
4.1 共享内存bank conflict到底怎么回事
共享内存虽然比全局内存快得多，但它不是无限带宽的。现代GPU的共享内存被划分成32个bank，每个bank在一个时钟周期内只能响应一次访问。如果同一个warp的多个线程访问同一个bank中不同地址，硬件就会把这次访问串行化，称为bank conflict。最坏情况下32路冲突，共享内存的带宽直接降到1/32。
回到刚才的矩阵转置，如果不加那列padding，转置后的数据写入tile[threadIdx.x][threadIdx.y]时，同一warp内正好有许多线程访问同一bank的不同行，冲突严重。加了1列padding之后，地址映射被打散，冲突就不再频繁发生。这个小技巧在不少并行算法里都适用，遇到共享内存访问瓶颈，先试试加padding。
初学时会觉得共享内存才多大，哪用得着这么精打细算。但真实场景里，卷积、FlashAttention这类算子都重度依赖共享内存，bank conflict带来的损耗可以轻易让二三十个百分点的性能消失。我在优化一个attention kernel时，仅仅改padding就获得了约18%的收益，这笔账相当划算。
4.2 只读路径：纹理内存与__ldg
现代GPU的L1缓存和纹理缓存是合并的，但通过__ldg或者只读缓存路径访问数据，有时能获得更好的缓存行为，尤其是当数据访问模式有较强空间局部性、又不想被常规读写混淆缓存时。const __restrict__修饰的指针在CUDA里通常会自动走只读路径，所以养成写const __restrict__的习惯，等于免费给编译器一个优化信号。
纹理内存还有一套独立的硬件插值单元，适合做图像类操作，但它真正的优势是二维局部性——纹理缓存在水平、垂直方向上都有较好的缓存策略，对于二维数组按行列混合访问的场景，比全局内存缓存更友好。普通的数值计算里用得不多，但做图像处理、体渲染时值得一试。
4.3 页锁定内存与零拷贝
CPU和GPU之间的PCIe传输是另一个经常被忽略的访存瓶颈。常规的malloc分配的内存是分页的，GPU访问前需要先锁定页面，拷贝过程会多一次复制。cudaHostAlloc分配的页锁定内存（pinned memory）则绕开了这个步骤，传输带宽可以提升不少。
页锁定内存的使用有一个小陷阱：分配过多会占用系统可用内存，反而拖慢CPU端性能。我一般只在数据需要反复传输、且单次拷贝量较大的场景用pinned memory，并控制总量不超过物理内存的10%~20%。接着还可以配合cudaMemcpyAsync和多个CUDA stream，把传输和计算重叠起来，这属于流水线优化，和访存模式是两条独立的优化维度，但加在一起效果异常明显。
零拷贝内存（cudaHostAllocMapped）则允许GPU直接访问CPU内存地址，省去显式拷贝。听起来很美，但PCIe延迟和带宽远不如显存，所以Zero-Copy只适合小数据量、访问频率低的场景。线程数一多、数据一密集，零拷贝会变成新的瓶颈。
4.4 数据布局：AoS与SoA的选择
访存模式的好坏不仅取决于kernel怎么写，更取决于数据在内存里怎么排列。最常见的选择是AoS（Array of Structures）和SoA（Structure of Arrays）。举个例子：如果有100万个粒子的坐标和速度，AoS是一个结构体数组，每个结构体包含x、y、z、vx、vy、vz；SoA则是6个独立数组，分别存x坐标、y坐标……GPU并行处理这类问题，SoA往往远优于AoS，因为一个warp内32个线程访问连续元素的x分量时，SoA保证完全合并，而AoS则会在不同分量之间跳动。
判断该用哪种布局时，核心看“一个warp同时处理的数据，在内存中是否连续”。如果每个线程处理一个粒子且需要所有分量，AoS可能更好，因为局部性更优；如果每个线程只管一个分量，SoA就是更优解。AI框架里许多算子因为数据是NCHW还是NHWC排列，性能差距达到两位数百分比，本质上也是这个道理。日常写代码时，一定要在数据结构设计阶段就考虑访存模式，而不是等kernel写完了再来修修补补。
5. 完整排查流程与工具实战
5.1 用Nsight Compute定位访存瓶颈
先声明，我踩过最大的坑就是“感觉不行就怀疑访存，然后瞎改”。正确的打开方式是先用工具量化，再对症下药。NVIDIA Nsight Compute应该是目前最好用的GPU kernel分析器，它能给出一个内核的**SOL（Speed of Light）**模型，直接显示当前内核有多少时间浪费在访存上，有多少浪费在计算等待上。
运行的方式很简单：ncu --set full ./your_application，然后打开报告找到Memory Workload Analysis。重点看三项：
- Mem Busy%：内存管线实际忙碌比例，如果这个值已经很高，说明瓶颈确实在访存。
- Max Bandwidth：达到的全局内存带宽占理论峰值的百分比。
- L1/TEX Hit Rate：命中率太低说明存在大量重复取数或缓存策略不当。
如果Mem Busy%不高，但内核依旧跑得慢，问题可能出在占用率不够、block调度不合理或者kernel启动参数上，这时候就不该继续折腾访存模式，而是先从占用率和指令混合下手。
5.2 常见访存问题速查表
为了方便平时排查，我把常见的访存问题整理成一个对照表，遇到性能异常先拿来对一下：
| 现象 | 可能原因 | 解决方向 |
|---|---|---|
| 全局内存吞吐远低于峰值 | warp内地址不连续 | 改写线程到数据的映射，确保合并访问 |
| L1命中率极低 | 数据复用差或每次读取粒度太碎 | 引入共享内存分块，或改向量化读取 |
| 共享内存成为瓶颈 | bank conflict严重 | 加padding，或重排共享内存布局 |
| 内核延迟高但占用率高 | 依赖链过长、缓存抖动 | 减少寄存器使用，或拆分kernel |
| 小数据量拷贝耗时占比高 | 固定开销超过了数据本身 | 使用页锁定内存、stream重叠或减少拷贝次数 |
| 多维度数组遍历慢 | 布局和访问顺序不匹配 | 调换存储顺序为NHWC或SoA布局 |
这个表不算完备，但覆盖了平时90%以上的访存性能问题。遇到新问题，我会按“先量化、再对照、再修改”的顺序处理，很少再靠瞎试。
5.3 几件实践中总结的事
第一，别迷信理论带宽。硬件标称的显存带宽，实际能跑到70%~80%就算很好了。能到90%以上的内核，通常用了向量化、合并访问、大块连续读写，并且尽量减少了写放大。不同GPU的L2策略不同，同一套代码在A100和RTX 4090上的访存表现可能差很远，优化一定要绑定具体硬件来测。
第二，block size对访存效率影响巨大。经验上，block size取128或256较稳妥，太小的block导致warp数量不足，无法隐藏延迟；太大的block又可能导致L1局部性变差。每个kernel的最优值不太一样，用Nsight Compute做个简单的block size扫描，比凭感觉设定更靠谱。
第三，共享内存和寄存器是可以互换的。有时候寄存器溢出，性能骤降；有时候共享内存占用过高，导致占用率打折。调整这两者的平衡，常常比直接改访存模式见效更快。我一般先把寄存器占用压到每个线程不超过32个，再看共享内存占用，尽量让每个SM同时驻留足够多的block。
6. 一次真实的内核优化全程记录
拿一个实际例子复盘整个流程。之前我在优化一个AI模型里的GELU激活融合kernel，核心任务是把激活函数、残差加法和LayerNorm之前的均值方差计算合并到一次显存读取中。初始版本很朴素，每个线程处理一个元素，读取一次输入、计算、写回一次输出，做了两次完整内存遍历。
用Nsight Compute一看，内存吞吐只有约40%的峰值，奇怪的是L1命中率也不低，查下去才发现问题出在写回时的部分cache line覆盖上——由于每个线程只写一个float，warp内刚好覆盖32个float的128字节范围，按理说应该没问题。但配合上LayerNorm需要二次读取同一批数据，导致每个元素实际上被读了两到三次。
优化方案是改成两阶段：第一阶段每个block读取一大块数据到共享内存并完成激活和残差加，第二阶段直接从共享内存读取做方差均值统计，最后再写回结果。这样全局内存只读一次、写一次，总访存量直接减少一半。改完后内存吞吐上升到75%左右，整体kernel时间缩短了接近50%。
这个过程最值得说的不是某个技巧，而是**“先量化，再动手”**。如果没有Nsight Compute的SOL分析，我大概率会去改线程映射、调block大小，折腾半天也找不到真正的问题在重复读取上。
7. 一些值得养成的开发习惯
写GPU代码和写CPU代码的思维模式很不一样。CPU上，编译器帮你做了大量缓存优化；GPU上，硬件的调度逻辑更直白，线程和数据的映射关系几乎决定一切。我总结了几条值得长期坚持的习惯。
第一，写kernel之前，先在纸上画出线程到数据的映射图。不需要多精细，用一个小warp举例，从左到右列出线程0到31的访问地址，看是否连续、是否对齐。这一步如果能形成肌肉记忆，很多访存问题在写代码阶段就能避免。
第二，把const __restrict__当成默认写法。它不仅帮助编译器走只读路径，还能给阅读代码的人一个明确的信号：这块内存只读，且没有别名。别小看这个，有时候它能带来意想不到的编译优化。
第三，用profile驱动优化，而不是直觉。每改一处，都跑一遍Nsight Compute对比前后指标。哪怕性能没有提升，也可以确认“改这一处没有让其他环节变差”。
第四，先做简单版本，再做复杂版本。很多优化技巧叠加在一起反而互相干扰。比如先保证合并访问，加向量化，再加共享内存分块，每步都做性能回归，这样能准确判断是哪一环带来的收益。
第五，养成看SASS的习惯。CUDA C代码和最终硬件执行的指令之间，隔了一层编译器优化。有时候你以为的高效写法，生成的SASS根本不是那么回事。Nsight Compute里能看到每个内核的SASS，重点关注是否存在本地内存溢出（local memory spill），如果有，第一件事是减寄存器、调整循环结构，而不是继续优化访存模式。
)
)
