很多人以为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来管理缓存。如果一个warp32个线程同时访问的地址恰好落在一个连续的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矩阵4096x4096float这个朴素版本的全局内存吞吐只有大约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怎么写更取决于数据在内存里怎么排列。最常见的选择是AoSArray of Structures和SoAStructure of Arrays。举个例子如果有100万个粒子的坐标和速度AoS是一个结构体数组每个结构体包含x、y、z、vx、vy、vzSoA则是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分析器它能给出一个内核的**SOLSpeed 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覆盖上——由于每个线程只写一个floatwarp内刚好覆盖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如果有第一件事是减寄存器、调整循环结构而不是继续优化访存模式。