我一度觉得自己对 GPU 性能分析挺有把握直到有一次在优化 GPT 推理的融合算子时被 Nsight Compute 里一组数字打脸DRAM Bandwidth 已经接近 80%SM 里的 FP32 计算单元利用率却不到 30%。算力明明有大量余量程序却跑不快唯一的解释就是——算力在等数据。后来我只是把循环里对一组中间结果的访问顺序从“逻辑顺序”改成“内存连续顺序”延迟立刻降了将近一半。从那时候起我就把 GPU 内存访问模式优化当成 AI 系统性能工程的核心课程而不只是 kernel 编程的选修项。这算是《AI 系统性能工程学习笔记》的第七篇。这篇我会把 GPU 访存优化从判定方法、底层机制、实战案例到 profiling 工具串起来讲一遍。目标读者是想深入调优 AI 推理/训练算子、写 CUDA kernel或者在给模型压延迟和提吞吐时总感觉“差了点什么”的工程师。1. 为什么很多 AI 算子的瓶颈在访存侧而不是算力侧1.1 用算术强度先判断方向我习惯在优化任何 kernel 之前先算一个非常粗糙的指标算术强度。这个概念不复杂就是“每从内存里读 1 字节数据对应多少浮点运算”单位是 FLOPs/Byte。GPU 芯片有自己的平衡点也就是峰值算力和显存带宽的比值。拿 A100 举例FP16 下峰值算力大约 312 TFLOPSHBM 带宽大约 1.6TB/s两者一除大概是 195 FLOPs/Byte。意思很清楚如果你的算子每字节访存只能支撑远低于 195 的浮点计算那它大概率是访存受限计算单元再怎么优化也很难把时间压下来。套到常见算子上看会更直观。向量加法要做 1 次浮点加却要读 2 个 float8 字节、写 1 个 float4 字节合计 12 字节访问量算术强度只有 0.08 FLOPs/Byte 左右。拿 0.08 和 195 比差了几千倍所以不管用什么花哨的向量化指令它都不可能把算力吃满优化方向只能是想办法减少跨层数据搬运、提高缓存命中。反过来看矩阵乘法在分块做得足够好之后算术强度能做到几百甚至上千属于典型的计算密集算子优化重点才应该放在把 Tensor Core 打满、减少线程空转、调整指令排布这些方向上。很多团队一上来就做大 kernel 融合、换更激进的编译器选项却不肯花十分钟算一下算术强度。方向选错之后后面投入的精力很容易全部打水漂。我现在的习惯是先算一个理论值再用 profiler 验证实际访存量和计算量两者对照基本能判断这个算子是“缺计算”还是“缺带宽”。1.2 计算单元空转是显性浪费第二个让我重视访存的原因是 GPU 硬件的工作方式。GPU 里的 SMStreaming Multiprocessor依靠大量线程并行来隐藏访存延迟一个线程发出访存请求后SM 不会干等而是切换执行其他可以运行的线程。这套机制本身没问题但它有一个前提——同一时刻必须存在足够多的独立访存请求可以并行发出去。如果线程访问的地址太离散内存控制器就得用很多次事务才能把数据凑齐每次事务之间的间隔就会暴露出来SM 再怎么能切换线程也填不满这些空档。以全局内存为例一个 warp 32 个线程同时发起访存时硬件按 32 字节的 sector 粒度和内存交互。如果 32 个线程访问的地址恰好落在一个连续的 128 字节区域里一次事务就全带回来了如果地址散落在几十个不同 sector那就要几十次事务处理。A100 的 L2 到 DRAM 带宽虽然高但它的上限就摆在那里事务数量越多每个事务能带回的有效数据比例就越低最终一定表现为 kernel 变慢。这个机制不是玄学是能在 profiler 里直接读指标验证的。我遇到太多“算力利用率一直上不去”的案例最后都归结到访存模式上。所以先记住一个结论访存优化的本质不是把访存量降到零——数据流动是必须的而是让每次内存事务搬运的数据尽量都被用上也就是提升流动效率。2. 访存优化的四个底层机制合并访问、sector、bank、对齐2.1 合并访问Memory Coalescing合并访问是 GPU 访存优化里最基础、也最容易被忽略的一条原则同一个 warp 的 32 个线程在访问全局内存时最好落在同一个连续内存区间里。线程 0 访问 A[base0]线程 1 访问 A[base1]一直到线程 31 访问 A[base31]这就是教科书式的完美合并。如果改成线程 i 访问 A[basei*64]地址一下子就散开了每次内存事务只能带回少量有效数据有效带宽会严重缩水。这种离散访问最容易出现在按“列优先”取数据的场景。比如一个按行存储的矩阵你想取某一列的所有元素对行做循环那访问地址天然跨步。解决思路通常是调整循环顺序或者先把矩阵转置再按连续方向访问。在 AI 领域也有一个非常典型的表现处理 [batch, channel, height, width] 张量时如果按 channel 维度循环取数而每个通道的数据在内存里是分离的那就等于逼着 warp 做离散访问。合并访问的真正价值不是“让每个线程用上缓存”而是“让一次内存事务覆盖尽可能多的有效地址”。理解了这一点很多优化的取舍就自然清楚了。比如判断要不要做算子融合、要不要改数据布局本质都在回答同一个问题经过这次改动之后每个内存请求能带回来多少字节有用数据2.2 sector 效率和缓存行的关系到 Ampere 以后更精确的访存粒度是 sector。一个 sector 是 32 字节一个 L2 cache line 通常是 128 字节也就是 4 个 sector。Nsight Compute 里的 sector 利用率指标会告诉你发出的内存事务里有效字节占多大比例。举个例子一个 warp 需要读 128 字节连续数据那么它最少只需要触发 1 个 128 字节的 wavefront4 个 sector 全部被利用。如果 32 个线程每个都去读一个 4 字节的 float并且地址间隔是 128 字节那每个 sector 里只有 4 字节被用到整体 sector 效率会非常难看。Memory Workload Analysis 页面里的 Memory Throughput 能直接显示你离硬件峰值有多远而 L1/TEX Hit Rate 和 L2 Hit Rate 可以帮助判断数据到底有没有在缓存里被复用。这一块的实际指导意义在于有时候你以为自己已经做了合并访问性能却依然不理想问题很可能出在 sector 粒度上。比如线程一次读 8 字节的 double或者一次读 16 字节的 float4同一个 warp 覆盖的连续范围会变大对事务数量的影响也很明显。所以优化访存时不能只看“连续不连续”还要看“每个线程一次性取多少字节”。向量化加载往往能同时改善合并访问和 sector 效率这也是为什么高手写 kernel 都喜欢用 float4、int4 这类类型。2.3 共享内存与 bank conflict共享内存是 GPU 片上的一块可编程缓存速度比全局内存快一个量级常用于保存会被重复使用的数据块。但共享内存本身也有一个经典陷阱叫 bank conflict。共享内存被划分成 32 个 bank每个 bank 每周期能提供一个 word通常 4 字节。如果同一个 warp 里有多个线程同时访问同一个 bank 的不同地址硬件就得把这一次访问拆成多次warp 的执行时间随之翻倍。最常见的情况是二维 tile 存在共享内存里线程按固定顺序读取列方向数据时地址恰好集中落在同一个 bank。共享内存数组 float tile[32][32] 是按行存储的线程 i 去读 tile[i][0] 时访问地址是第 i32 个 word。因为 bank 索引是对 32 取模i32 对 32 取模永远是 0于是所有线程都访问 bank 032 路冲突直接发生。性能掉一个数量级都不奇怪。解决办法非常经典在声明共享内存时给每一行加一个 padding。把 float tile[32][32] 改成 float tile[32][33]每行多出一个 word下一行起始地址就会偏移 1列方向的 bank 分布就被打散了。这个技巧在后面的矩阵乘法案例里我会实际展开。需要提醒的是bank conflict 并不容易在代码审查时发现因为问题往往藏在索引计算的细节里最好的方式是直接用 profiler 里的共享内存 conflict 指标去验证。2.4 数据对齐与连续分配还有一个经常被忽视的细节是内存对齐。GPU 访问全局内存时对 32 字节 sector 和 128 字节 cache line 的边界对齐很敏感。用 cudaMalloc 分配的内存通常是对齐的但片内偏移可能破坏对齐。比如一个 kernel 让线程去读结构体数组的某个字段如果结构体大小是 12 字节每个字段天然只按 12 字节对齐读取时就会跨 sector效率下降。在 AI 场景里这个问题的常见变体是 PyTorch 的 tensor 虽然底层是连续内存但经过 transpose、permute 之后逻辑上不再连续。直接把这种张量喂给自定义 CUDA kernel读写地址就变成跨步访问访存表现会很差。正确的做法是先调用 contiguous() 做一次真实的内存搬运或者让 kernel 内部意识到 stride 的存在并做针对性处理。不是所有 strided 访问都一定要消除但至少要知道它的成本在哪里。3. 一个能直接复现的案例从朴素 SGEMM 到分块 SGEMM为了避免整篇都是干巴巴的概念我用矩阵乘法SGEMM当例子把访存优化的完整思路走一遍。AI 场景里很多计算已经封装进 cuBLAS 或 Tensor Core 库了但 SGEMM 的结构高度通用理解了它再去看卷积、Attention 里的访存问题会顺畅很多。3.1 朴素版为什么慢最简单的 SGEMM kernel 是这样的每个线程负责计算输出矩阵 C 的一个元素通过一层循环累加 A 的一行和 B 的一列。__global__ void sgemm_naive(const float* A, const float* B, float* C, int N) { int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; float sum 0.0f; for (int k 0; k N; k) sum A[row * N k] * B[k * N col]; C[row * N col] sum; }如果把 block 定义成 16x16 线程每个线程计算一个 C 元素。对 C 中相邻线程来说B 的访问是连续的B[k * N col] 的 col 连续所以合并访问是成立的。但 A 的访问就完全不一样了row 不同的线程访问 A 的不同行地址间隔是 N 个 float完全离散。随着 k 不断变化A 的每一行都要重复从全局内存加载每个元素只被一个线程使用一次没有任何复用。这个版本的算术强度很低。矩阵乘法每输出一个结果需要读 2N 个浮点数、写 1 个浮点数如果 N 很大算术强度大概只有 0.5 FLOPs/Byte 左右离平衡点差了几百倍。所以朴素 SGEMM 跑不快不是计算指令的问题是访存模式的问题。很多人在这个阶段开始怀疑 GPU 算力其实冤枉硬件了。3.2 分块 共享内存怎么降低 HBM 压力矩阵乘法最经典的优化思路是分块tiling核心洞察是输出 C 的一个 tile只依赖 A 的若干行和 B 的若干列。如果让一个线程块负责计算 C 的一个小 tile那么它可以把 A 和 B 对应的数据块先搬到共享内存然后在这个块内反复复用避免每次都从全局内存读取。具体而言假设每个 block 负责 16x16 的 C tile。那么为了计算这个 tile我们需要在 k 方向上也分块例如每轮读入 A 的 16x16 块和 B 的 16x16 块到共享内存。这样一来全局内存里的每个数据元素只会被读取一次之后在共享内存里被多次复用。全局内存访问量从原来的 O(N^3) 级别下降到 O(N^3 / tile_size) 级别算术强度随之大幅上升。这个思路在 CPU 上看只是加了一层缓存但在 GPU 上它直接决定了 DRAM 带宽到底是不是瓶颈。真正高手写的 SGEMM 还会继续叠加优化用 float4 向量化加载、double buffering 隐藏拷贝延迟、用 Tensor Core 做矩阵乘累加。但这些高级技巧全都建立在一个基础之上先通过分块把全局内存的重复访问降下来。如果这一步没做好后面所有指令级优化都是在浪费精力。3.3 消除 bank conflict 的 pad 技巧把 A 和 B 的 tile 放进共享内存之后很快会遇到前面说的 bank conflict。假设共享内存数组声明成下面这样__shared__ float As[16][16]; __shared__ float Bs[16][16];如果线程按列方向读取 As比如需要 As[i][k]其中 i 是线程索引那么地址偏移是 i16 k。对 32 取模时16i 在 i 为偶数时得到 bank 0在 i 为奇数时得到 bank 16。也就是说大量线程会集中在少数几个 bank 上产生严重的 bank conflict。解决办法是给每一行都加一个 padding__shared__ float As[16][17]; __shared__ float Bs[16][17];现在地址偏移变成 i*17 k对 32 取模后i 在 0 到 15 之间变化时bank 分布基本被均匀打散冲突消失。这个操作只增加了一点共享内存占用却能把共享内存的有效带宽拉回接近峰值。我自己在写这类 kernel 时几乎已经形成肌肉记忆了凡是共享内存声明先考虑有没有 bank conflict 风险再决定要不要 pad。3.4 分块尺寸选择与实测数据分块尺寸不是越大越好。共享内存容量有限block 太大的话同一个 SM 上能同时驻留的 block 数会减少占用率下降调度灵活性也变差。以 A100 为例每个 SM 有 164KB 左右可配置共享内存。如果直接用 float 的 32x32 tileAs 和 Bs 加起来需要 32324*28KB一个 SM 可以放不少 block但如果换成 128x128 的 tile单 block 就要 128KB一个 SM 只能放一个 block隐藏延迟的能力反而下降。另一个关键参数是线程粒度。比较理想的做法是让一个线程计算多个输出元素比如 block 用 128 个线程每个线程负责 8x8 或 4x4 的结果块。这样做有双重好处一是提高寄存器里的数据复用减少共享内存访问次数二是能用向量化加载比如 float4进一步改善合并访问。我在 2080Ti 上做过一次简化版的分块 SGEMM 实验没有开 Tensor Core只用普通 FFMA 指令。朴素版大约 350 GFLOPS分块为 32x16、加上 padding 和 float4 加载之后大概能到 1800 GFLOPS 左右提升约 5 倍。再往后做 double buffering 和 unroll 也能继续涨但收益明显递减。不同显卡上的绝对数值会有差异但趋势一致访存优化带来的收益往往比指令级微调大得多而且更稳定。4. 把访存优化思维用到 AI 训练推理布局、Attention、KV Cache4.1 NCHW 和 NHWC布局变化会带来跨步访问AI 框架里最容易被忽略的访存问题之一是张量布局。PyTorch 在图像任务上默认使用 NCHWTensorFlow 在很多设备上默认 NHWC。这两种布局在卷积核里的访问模式差异非常大。对 NHWC 来说四个维度在内存中是连续排列的相邻像素在内存里紧挨着按像素维度做循环时很容易实现合并访问和 CPU 向量化。NCHW 则是通道维度的数据分开存放如果算子要按空间位置取多个通道的值跨步就会比较大。许多推理引擎在做卷积优化时会专门做布局转换把 NCHW 转成 NHWC 或者 NHWC1 这类更适合硬件合并访问的格式。另一个常见坑是 permute 之后的内存跳跃。PyTorch 里 tensor.permute() 并不会真正搬数据只是修改了 stride 元数据。逻辑上看着是整齐的多维张量实际内存布局已经乱了。直接把这个 tensor 传给自定义 CUDA kernel读写就成了跨步访问。正确的做法是先调用 contiguous()或者让 kernel 针对 stride 做专门优化。不要小看这一步我在实际项目里见过仅仅因为漏了一次 contiguous()kernel 耗时翻了 3 倍的案例。4.2 Attention 的访存密集特征与 FlashAttention 的思路Transformer 里最典型的访存密集场景是 Attention。标准实现会先算 QK^T 得到 S 矩阵然后 softmax最后乘 V。中间多个矩阵要写回全局内存再读回来这些中间结果的大小是 batch * num_heads * seq_len * seq_len长序列场景下非常惊人内存流量成了主要瓶颈。FlashAttention 的核心思路就是把中间阶段融合并重新组织访存不在全局内存里保存完整的 S 矩阵而是按 block 计算在 block 内部完成 softmax 的在线更新并且把结果直接用于和 V 的块乘累加。这么做最大的收益是降低全局内存流量而不是减少计算量。这个思想完全是第三章分块思想的延伸把一个大矩阵运算切成 tile让数据在片上存储里被最大化复用。如果你想深入理解 FlashAttention与其一上来就看那一堆数学推导不如先想清楚一件事——它到底把多少次全局内存读写省掉了。省掉了这些读写DRAM 压力降下来长序列推理的瓶颈自然就缓解了。这也是为什么 FlashAttention 在推理和训练里都能带来实打实的加速。4.3 KV Cache 分配与显存碎片化问题在大模型推理里KV Cache 是另一个和访存模式强相关的话题。每个请求都会生成一个不断增长的 KV Cache如果不预先分配连续的大块显存而是在每次需要扩展时都临时 cudaMalloc不仅会有分配延迟还会让显存碎片化。碎片化之后后续申请大块显存可能失败数据也可能物理分散L2 和 TLB 的命中率跟着下降。业界常见的做法是提前按最大长度预留连续空间或者用显存池统一管理。vLLM 的 PagedAttention 则更进一步把 KV Cache 切成固定大小的物理块用页表映射逻辑位置既缓解了碎片问题也避免了预留整块显存带来的浪费。如果你只在单机、短序列场景里调试小模型可能体会不深一旦上多路并发、长上下文KV Cache 的分配方式对吞吐的影响会非常明显。这些场景说明访存优化不只是 kernel 内部的事它也包含系统层面的内存布局和分配策略。做 AI 系统性能工程时眼光不能只停在 CUDA 代码里内存池、缓存策略和数据布局同样是第一线战场。5. 用 Nsight Compute 定位问题关键指标与一次排查链路5.1 关键指标怎么读遇到性能问题我最常用的工具是 Nsight Compute。打开一个 kernel 的 Memory Workload Analysis 页面重点看这几个指标指标含义我的判断习惯Memory Throughput内存子系统相对峰值的利用率接近 90% 以上说明内存侧饱和DRAM Throughput实际打到显存的带宽低且 Memory Throughput 高可能是共享内存或 L1/L2 交互瓶颈L1/TEX Hit Rate一级缓存命中率高不一定好要和实际耗时结合看L2 Hit Rate二级缓存命中率高说明复用不错低说明数据流太离散Achieved Occupancy实际占用率低于理论值要找寄存器或共享内存超限的原因Sector 效率每个内存事务有效 sector 比例低说明合并访问没有做好我不建议只盯着 Flops 利用率。很多 kernel 的 Flops 利用率不高根源是访存卡住了这时候去优化指令排布就是白费力气。正确的流程是先看内存侧指标再判断瓶颈在 DRAM、L2、L1 还是共享内存最后才回到指令和线程调度层面。5.2 一次真实排查链路我有一阵子在调一个 GELU LayerNorm 融合算子延迟总是比预期高 15%。Nsight Compute 显示 DRAM Throughput 不高但 Memory Throughput 很高L1 Hit Rate 也很高。这三个指标放在一起基本能排除“显存带宽打满”的可能我立刻怀疑问题出在片上缓存或共享内存交互上。后来把 kernel 源码翻了一遍发现有个地方把中间激活数组放进了共享内存而且数组没有做 paddingbank conflict 非常严重。改成按 float4 读取、并给共享内存加 padding 之后延迟立刻降了 18%。如果当时只靠读代码猜逻辑这个问题可能要找很久用 profiler 按“先看 Memory Workload再看 L1/L2 行为最后看指令统计”的顺序去筛其实很快就能锁定。profiling 不是可有可无的步骤它是访存优化的眼睛。特别是 bank conflict、sector 效率这些指标肉眼根本看不出来只有工具能告诉你实情。我踩过太多“以为改好了但实际上没变化”的坑后来都靠 profiling 拉回了正轨。6. 三个被误判的“伪优化”场景与我的踩坑教训6.1 L1 命中率高并不等于访存没问题我见过有人拿 L1 Hit Rate 95% 的结果证明缓存很健康。但 L1 命中率高可能只是因为数据量小、全塞进了缓存并不代表数据复用设计得好。更关键的是看 DRAM Throughput 和实际耗时。如果 kernel 依然慢问题可能在共享内存、寄存器溢出或者指令依赖上。指标要组合着看不能单点下结论。最稳妥的办法是每次改动只动一个变量记录 profiler 里的全套指标变化。6.2 显存带宽大不代表可以随便访问A100 的带宽再高也扛不住低效访存模式。每秒 1.6TB 听起来很多但大模型的中间张量动辄是 GB 级别几个来回就能把带宽吃干。不要觉得“反正带宽高多读几遍没关系”。凡是可以从全局内存搬到共享内存或寄存器里复用的数据都值得搬凡是能合并的读写都值得合并。访存模式差的 kernelILP 和 TLP 都救不回来。6.3 把“加同步”当成性能优化另一个常见误区是觉得访问共享内存前多写几个 __syncthreads() 总没有错。但多余的同步会让线程空转占用率越高代价越大。同步本身不产生任何计算价值它只是在等待数据依赖满足。正确的做法是仔细分析数据依赖只在必要的地方同步并且尽量用 double buffering 之类的技术把同步等待和计算重叠起来而不是用一堆同步去掩盖数据流没有想清楚的问题。最后说一点个人体会。访存优化的本质不是把代码改得花哨而是让数据在正确的层次、正确的时间以合并的方式流动。我在实际项目里养成了三个习惯优化前先算算术强度优化中每次只改一个变量并记录 profiling 结果优化后回到 Nsight Compute 看 Memory Workload 验证。这套流程听起来慢实际上比“想到哪改到哪”快得多。后面如果对 FlashAttention 的完整推导、或者 KV Cache 的显存池设计感兴趣我可以再单独展开写一篇。