双十一深夜，快递分拣中心的灯还亮着。工人们分成 32 人一组，组长的任务是带着整组人从巨大的中央仓库把包裹搬到分拣台。有人聪明：安排 32 个人顺着同一条货架排队走，一次搬回 32 个相邻包裹——这叫“合并访问”；有人糊涂：让组员散到仓库各个角落，每人东奔西跑找一件，来回折腾。同样的 32 人、同样的货量，效率差了三五倍。这个故事里藏着 CUDA 性能优化最核心的学问：GPU 上的“中央仓库”就是全局内存，容量大但慢；“分拣台”是共享内存，快得多但容量小；工人口袋是寄存器，最快。搞懂这几层怎么配合、怎么避免“撞抽屉”（bank conflict）、怎么让数据提前到位（预取），你就掌握了 GPU 访存优化的基本功。这篇文章从一次双十一分拣任务说起，带你逐个击破访存优化的关键技法。

岗位关联

对 AI 从业者：大模型训练与推理的算子底层几乎都跑在 GPU 上，访存优化直接决定吞吐和成本；掌握这套方法论，你写出的 CUDA 算子性能可能翻倍，这是高性能 AI 工程岗位的核心竞争力。

目录
  1. 可视化演示
  2. 访存模型基础（仓库地形图）
  3. 合并访问模式（排队上货）
  4. 共享内存应用（中转货架）
  5. Bank冲突消除（撞抽屉）
  6. 寄存器重用与数据预取（口袋战术）
  7. 缓存优化策略（热门货架）
  8. 量化评估手段（绩效仪表盘）
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

步进式展示CUDA GPU访存优化四种核心技术：全局内存合并访问、共享内存减重复访存、Bank Conflict消除、寄存器重用与预取，点击步骤切换查看各优化策略的线程-内存访问模式。

本节配有交互动效（进入视口后自动加载，约 3KB）。

全局内存合并访问 同warp内32线程访问连续内存地址，GPU合并为单次事务 合并后仅需1次内存事务，带宽利用率接近峰值 Coalesced: 1 transaction / 128B 合并访问：32线程×4B=128B单次事务

访存模型基础（仓库地形图）

双十一分拣中心新来了 32 名临时工，组长第一件事是带他们认路。场地有三层：远处中央仓库堆着几百万件包裹，来回一趟要十分钟；工位旁的分拣台是全组共用，取件以秒计；每人腰间的工具袋，随手一摸就到。GPU 的世界完全对应：全局内存容量几十 GB，延迟数百个时钟周期；共享内存只有几十 KB，延迟几个周期；寄存器容量极小，但访问零延迟。

光看延迟不够，还要看带宽——传送带一分钟能运多少件，比单趟跑腿时间更能决定总吞吐。全局内存有每秒上千 GB 的带宽，但任何一次请求都要付出几百周期的固定等待；寄存器不存在“等待”，拿来即用。于是优化的第一直觉是：数据尽量放在离线程近的地方，减少长途跋涉。

访存以线程束（warp）为单位发起。一个 warp 就是 32 个线程组成的捆绑小队，全局内存的每次 load/store 都由整队 32 人共同触发。如果程序大部分时间在等数据返回，就属于“访存受限”（memory-bound），此时优化算术毫无用处，只能从访存模式下手。

记住这张地形图，先想清楚数据在哪一层被频繁使用，再设计搬运策略。后续所有优化，都是在仓库、分拣台、口袋之间找最省路线。

T=DB+L\displaystyle T = \frac{D}{B} + L
  1. D 数据量需要搬运的包裹总字节数，由任务规模决定。
    DD
  2. B 带宽单位时间内传输通道能搬运的字节数。
    BB
  3. \frac{D}{B} 传输耗时数据在通道上连续流动所消耗的时间。
    DB\frac{D}{B}
  4. L 访存延迟从发起请求到第一个数据落地的固定等待时间。
    LL
  5. T 总访问时间搬运耗时与往返等待之和，这就是最基础的访存耗时模型。
    T=DB+LT = \frac{D}{B} + L

合并访问模式（排队上货）

同样的 32 人小队，聪明的队长让大家站到同一条货架前，按编号顺序连续取下第 0 到第 31 个包裹，传送带一次吐出 32 件。GPU 全局内存的“传送带”按缓存行工作——一次事务最少取 32 字节，通常一整行 128 字节。让 warp 的一次请求恰好吃满一条缓存行，就是“合并访问”。

合并访问的技术表述是：warp 内第 t 个线程访问的地址与 t 顺序对应，且连成连续区域。写成代码就是 data[base + tid]。这个模式对 float 数组最自然：32 个线程的 32 个 float（128 字节）正好填满一条缓存行，一次事务全部解决。相邻线程访问相邻地址，就像排队上货一样高效。

反例是跨步访问：data[base + tid * 8] 会让请求散落在 8 条不同的缓存行里，一次访问被迫拆成 8 次事务。想象 32 个人分别跑去 8 个不同货架取件——跑腿总次数没变，但传送带每趟只装一点点，利用率暴跌。判断方法是用 Profiler 观察 Global Load Transactions 与请求数的比值。

排队上货比散落翻货快三四倍，这就是合并访问的威力。凡是看到跨步、转置、按列访问，先回到这张“排队图”重新设计线程与地址的映射。

E=BusefulBtotal×100%\displaystyle E = \frac{B_{\rm useful}}{B_{\rm total}} \times 100\%
  1. B_useful 有用字节数程序真正需要的数据量，即 32 个线程实际消费的字节。
    BusefulB_{\rm useful}
  2. B_total 总线搬运字节数硬件实际从显存搬回的字节数，包含被浪费的相邻缓存行数据。
    BtotalB_{\rm total}
  3. E 合并访问效率有用字节占总搬运字节的比例，越接近 100% 代表合并越完美。
    E=BusefulBtotalE = \frac{B_{\rm useful}}{B_{\rm total}}
  4. \times 100% 换算百分数把效率从小数换算成直观的百分比。
    ×100%\times 100\%

共享内存应用（中转货架）

如果全组反复要用同一批包裹，聪明的组长会先把它们批量搬到分拣台，接下来全组都在台上取用——跑仓库的次数从“每取一件跑一次”变成“每批跑一次”。共享内存就是 GPU 里的分拣台：位于片上、速度接近寄存器，且能被线程块内所有线程并行读写。

经典套路叫“分块”（tiling）：把大数组切成一个个 tile，批量拷入共享内存，循环体内只从共享内存取值。全局访存次数从“每算一次就访存一次”降为“每处理一个 tile 访存一次”，代价只是一次拷贝的开销。数据被复用得越多，摊薄到每次使用的搬运成本就越低。

矩阵转置是教科书级案例。朴素转置读行写列，写回时 32 个线程访问同一列的不同行，地址跨度大，非合并访问。改造思路：按行合并读入共享内存 tile，再从共享内存按列合并写出，两侧全局访问都变成连续模式。团队协作必须 __syncthreads() 齐步走——先写完 tile 的线程要等全组都写完才能开始读，否则会读到半成品数据。

共享内存分块没有必须公式，但“复用倍数”是核心衡量指标：当 tile 尺寸为 $T\times T$、每个线程处理 $T$ 个元素时，全局访存次数降为原来的 $\frac{1}{T}$。tile 越大复用越多，但共享内存容量有限，需要在容量与复用之间取平衡。

1. T tile 边长：分块矩阵的边长，决定一次搬运多少数据。

TT

2. \frac{1}{T} 访存缩减比：tile 越大，全局访存被均摊的次数越少。

1T\frac{1}{T}

中转载台的策略总结：能复用才搬，搬一次用好几次，用前面讲的地形图和合并访问规则，共同决定这个搬运计划是否划算。

Bank冲突消除（撞抽屉）

分拣台看似一块大桌面，底下其实是 32 个独立抽屉（Bank），每个抽屉同一时刻只能被一只手取走一件。硬件共享内存同样划分为 32 个 bank，相邻的 4 字节轮流落入相邻 bank。只要每个工人伸手取的是不同抽屉，32 人就能同时拿到货。

当 warp 里多个线程同时访问同一个 bank 的不同地址时，硬件被迫排队串行，这就是 bank conflict。比如 32 个线程全撞向 bank 0，一次访问要拆成 32 拍完成，分拣台从并行变成龟速。常见触发场景：共享内存二维数组按列访问、行宽恰好是 32 的整数倍。

解法是“错位填充”（padding）：把 s[TILE][TILE] 改成 s[TILE][TILE + 1]，让行与行之间错开一个元素，原本落在同一 bank 的地址被弹到不同抽屉。另一个特殊机制是广播：当多个线程读的是同一地址，硬件只发一次广播，人人有份，不算冲突。

撞抽屉问题用 Nsight 的 conflict 计数器一眼可见。设计共享内存布局时，把“几号线程访问几号 bank”当成排座位表来规划，必要时用填充列把邻居们错开。

Tbank=k×Tbase\displaystyle T_{\rm bank} = k \times T_{\rm base}
  1. k 冲突路数同一 bank 内同时被多少线程访问，就称几路冲突。
    kk
  2. T_base 无冲突访问时间没有冲突时，一次共享内存访问只需一个时钟周期。
    TbaseT_{\rm base}
  3. T_bank 实际访问时间冲突路数乘以单次时间，2 路冲突就慢一倍，32 路冲突慢 32 倍。
    Tbank=k×TbaseT_{\rm bank} = k \times T_{\rm base}

寄存器重用与数据预取（口袋战术）

工人最常用的螺丝刀绝不放进仓库，而是插在口袋里。GPU 也有这样的“口袋”：寄存器，每个线程私有、零延迟、无冲突。编译器通常会自动把局部变量放寄存器，但程序员的写法决定了它能否成功——把循环内反复使用的中间量提取到局部变量，就能避免每次迭代都重新读全局内存。

两个进阶手法是循环展开与数据预取。循环展开（unroll）让多个迭代共用一份已经加载进寄存器的数据，减少重复访存；数据预取（prefetch）则是在当前迭代计算的同时，提前把下一轮要用的数据从全局内存搬到寄存器或共享内存，让访存延迟与计算时间重叠，工人算账的时候传送带已经在路上。

但寄存器不是无限的。一个 SM 的寄存器文件总量固定，每个线程用太多寄存器会降低可驻留线程数，活跃 warp 变少，延迟掩盖能力变弱；更糟的是寄存器溢出（spill），被挤出的数据落回慢速的“局部内存”（物理上在全局内存里），得不偿失。

口袋战术的核心是平衡：既别让工人空手跑仓库，也别把口袋塞成炸药包。用 --maxrregcount 或 __launch_bounds__ 约束寄存器用量，在小步实验中找到吞吐峰值。

O=WactiveWmax×100%\displaystyle O = \frac{W_{\rm active}}{W_{\rm max}} \times 100\%
  1. W_active 活跃线程束数一个 SM 上当前能并行驻留的 warp 数量。
    WactiveW_{\rm active}
  2. W_max 硬件上限线程束数SM 硬件设计的最大 warp 容量。
    WmaxW_{\rm max}
  3. O 占用率活跃 warp 占硬件上限的百分比，反映并发掩盖延迟的能力。
    O=WactiveWmaxO = \frac{W_{\rm active}}{W_{\rm max}}
  4. \times 100% 换算百分数把占用率显示为直观的百分比形式。
    ×100%\times 100\%

缓存优化策略（热门货架）

仓库经理发现某些货被反复取用，于是在门口设了“热卖区”。GPU 的 L2 缓存就是所有流处理器共享的热卖区：数据第一次访问落到显存，后续访问若命中 L2，就直接从片上取回，省掉一趟长途运输。对于被多线程反复读的数据，L2 的命中率直接影响整体速度。

对只读数据，可以用 __ldg() 引导它走只读数据缓存路径，绕过常规的 L2/全局路径，对纹理等只读访问更友好。数据布局同样决定命中率：结构体数组（AoS）让同一线程的字段散落各处，而数组结构体（SoA）让所有线程按序访问同一字段，缓存行友好得多，GPU 上优先 SoA。

对齐常常被忽略。将大数组起始地址按 128 字节对齐（cudaMalloc 默认对齐 256 字节），保证 warp 单次事务恰好吃满整条缓存行，不会跨行拆成两次事务。此外，CPU 与 GPU 之间的数据走 PCIe，带宽仅为显存的几十分之一；传输应该大块、异步、可与计算重叠，而不是频繁做小碎粒同步拷贝。

热卖区解决的是“重复访问”的痛点；首次访问的痛点仍要靠合并访问解决。两级配合，访存系统才算完整。

缓存命中本身没有固定公式，但“有效利用”可以用命中率表达：$H = \frac{N_{\rm hit}}{N_{\rm access}}$，命中次数占总访问次数的比例。命中率越高，走 L2 的快速路径越多。

1. N_hit 命中次数：请求的数据已在缓存中的次数。

NhitN_{\rm hit}

2. N_access 总访问次数：程序发起的全部访存请求数。

NaccessN_{\rm access}

3. H 命中率：可以粗略理解为“热卖区直接拿到的比例”。

H=NhitNaccessH = \frac{N_{\rm hit}}{N_{\rm access}}

量化评估手段（绩效仪表盘）

分拣中心大堂挂着一块绩效屏，实时显示每个小组的步数、取件数、空跑时间。CUDA 开发者也有这块屏：Nvidia Nsight Compute 能统计全局内存吞吐、L2 命中率、bank conflict 次数、活跃 warp 数，一眼定位瓶颈在哪一层。

第二步是用 Roofline 模型判断程序属哪类受限。把机器的算力（FLOP/s）和带宽（Byte/s）画成二维坐标，二者比值就是“机器拐点”。再把程序的计算强度与拐点比较：低于拐点，程序受带宽限制，优化访存才有意义；高于拐点，受计算限制，该优化算术。

优化后统一用加速比量化收益：$S = T_{\rm old} / T_{\rm new}$，或直接观察 DRAM 吞吐是否逼近硬件峰值。建议每次只改一个变量、跑同一基准、记录数据，形成自己的优化清单：合并访问 → 共享内存复用 → 消除 bank conflict → 寄存器与预取 → 缓存布局。

没有数据支撑的优化只是玄学。像绩效屏一样把每次改动晒在阳光下，访存优化才能成为可复制的工程方法论。

---

I=FB\displaystyle I = \frac{F}{B}
  1. F 浮点运算量程序执行的总浮点运算次数，单位 FLOPs。
    FF
  2. B 访存字节数程序访问 DRAM 的字节总量。
    BB
  3. I 计算强度每搬运 1 字节数据能完成多少次浮点运算，是 Roofline 模型的横轴坐标。
    I=FBI = \frac{F}{B}

AI 应用场景

LLM 推理中的算子融合

阶段：推理引擎 Kernel 实现与调优

解决的问题：大模型生成 token 时，RMS Norm、Softmax、残差加等算子频繁读写中间张量。朴素实现会把激活值反复写回全局内存再读出来，访存占比极高、时延感人。利用访存优化：算子融合让中间量留在共享内存和寄存器，每次访问保证相邻线程合并。

具体案例：vLLM 的 fused RMSNorm 将归一化、加残差、乘缩放一次完成，中间张量不出 SM；KV cache 写入采用合并访问与对齐布局，显著降低了显存带宽压力。

图像预处理中的 HWC→CHW 转置

阶段：数据加载与布局转换

解决的问题：训练管线中图像常以 HWC 读入，模型需要 CHW。朴素转置是典型非合并访问，带宽浪费严重。用共享内存分块转置，tile 写入时用填充错开 bank，两侧都变成合并访问。

具体案例：DALI 预处理管线中，用 32×32 共享内存 tile 完成 HWC→CHW，实测吞吐可提升数倍，训练数据读取不再成为瓶颈。

推荐系统 Embedding 查询

阶段：访存密集型 Kernel 编写

解决的问题：推荐系统的大 Embedding 表按 ID 随机取向量，天然非合并，且高频 ID 反复访问。结合缓存优化：热 Embedding 放在 L2/共享内存，布局改 SoA，查询批次按 ID 排序后批次内合并，大幅提升吞吐。

具体案例：某广告推荐系统将 Embedding 查询 Kernel 的 DRAM 吞吐从峰值的 40% 提升到 85%，单查询时延下降 60%。

Python 演示

python37 行
import time  def transactions_needed(threads, stride, line_size=32):    """计算一个线程束访问触发的缓存行事务数"""    lines = {t * stride // line_size for t in range(threads)}    return len(lines)  def measure_access(stride, threads=32, iterations=200000):    """模拟按地址顺序读取数据并返回耗时"""    addrs = [t * stride for t in range(threads)]    data = tuple(range(max(addrs) + 1))    start = time.perf_counter()    total = 0    for _ in range(iterations):        for a in addrs:            total += data[a]    elapsed = time.perf_counter() - start    return elapsed, total  def main():    print("=== CUDA 访存模式模拟(32 线程/线程束) ===")    for stride in (1, 2, 4, 8, 16, 32):        tx = transactions_needed(32, stride)        print(f"stride={stride:>2}: 事务数={tx}, 事务放大={tx}x")     t1, s1 = measure_access(1)    t2, s2 = measure_access(32)    print(f"合并访问耗时: {t1:.4f}s")    print(f"跨步访问耗时: {t2:.4f}s, 慢 {t2 / t1:.1f} 倍")    print(f"校验值: {s1 + s2}")  if __name__ == "__main__":    main()
目的：用 Python 模拟 32 线程的访存事务合并效果，直观展示 stride 对缓存行事务数的影响。
  1. lines = {t * stride // line_size for t in range(threads)}用集合去重，统计需要触碰多少条缓存行
  2. t * stride生成线程 t 要访问的地址，stride 越大地址越分散
  3. data = tuple(range(max(addrs) + 1))构造足够长的“内存”序列，模拟地址空间
  4. for a in addrs: total += data[a]模拟真实访存：按地址逐个取数
  5. transactions_needed(32, stride)返回事务放大倍数，stride=1 为 1 条，stride=32 为 32 条
  6. t2 / t1计算跨步访问相对合并访问慢了多少倍

常见误区

误区共享内存比全局内存快，我把能搬的数据全部塞上分拣台，访问一定更快！

正解分拣台容量有限，一个线程块占用的共享内存越多，SM 能同时容纳的线程块就越少，活跃 warp 变少，访存延迟无法被并发掩盖。结果往往是“全局访存少了、等待变多了”，总吞吐反而下降。共享内存的正确用法是“分块复用”，而不是无脑囤货。

为何常见：只比较了单次访问速度，忽略了存储容量与占用率之间的连锁反应。

误区遇到非合并访问，套一层共享内存转置就万事大吉了。

正解数据拷入共享内存那一刻同样要遵循 32 bank 的寻址规则；如果 tile 写回时行间地址恰好撞同一 bank，冲突带来的串行化可能抵消全部收益。正确姿势是先用本章方法分析 bank 分布，必要时加 padding，再决定是否值得动用共享内存。

为何常见：把共享内存当成缓存黑盒，没意识到它内部还有 bank 结构约束。

误区我在 kernel 里声明了局部数组，它应该放在寄存器里，访问超快。

正解寄存器文件按线程配额有限（如最多 255 个）。局部变量一旦超出配额，编译器会把多余部分溢出（spill）到“局部内存”——它物理上在全局内存里，只不过带私有缓存。大数组、动态索引、取地址操作都可能落回慢速内存，速度可能比原来更差。

为何常见：混淆了 C 语言“局部变量”与 CUDA“寄存器”，不知道编译器存在 spill 机制。

误区我的地址明明是连续的，为什么 Profiler 还是报告多次事务？

正解合并访问还要求线程编号与地址顺序一一对应，且起始地址对齐缓存行。若第 0 号线程访问高地址、第 31 号线程访问低地址，或 base 未按 128 字节对齐，连续地址也会被拆成多次事务。检查 base + tid 的映射方向与 base 的对齐情况。

为何常见：只盯“连续”两个字，忽略了 warp 内线程与地址的对应关系和对齐边界。

误区GPU 核多线程多，把 block 开到最大，访存问题自然被淹没。

正解访存受限程序的瓶颈是 DRAM 带宽，不是线程数。线程再多，每秒能搬的字节总量由硬件带宽决定；只有让每个线程的每次访问都合并、对齐、少重复，才能逼近带宽上限。盲目加线程只会增加调度开销。

为何常见：把 CPU 多核思维搬上 GPU，混淆了并行度与字节吞吐两个概念。

面试题

Q1. CUDA 中访问延迟最低的存储单元是？

单选题 · easy

  • A.全局内存
  • B.共享内存
  • C.寄存器← 正确答案
  • D.常量内存

答案：C。寄存器是每个线程私有的片上存储，零延迟、无冲突；共享内存是片上但可能撞 bank；全局内存要付出数百周期延迟。工程师应优先把热点数据放寄存器，就像工人把螺丝刀放口袋。

Q2. 一个线程束（warp）包含多少个线程？

单选题 · easy

  • A.16 个线程
  • B.32 个线程← 正确答案
  • C.64 个线程
  • D.128 个线程

答案：B。warp 是 GPU 调度的基本单位，固定 32 个线程。全局内存的合并访问、共享内存的 bank conflict，都站在这 32 人小队的视角分析。

Q3. 关于共享内存的位置，下列说法正确的是？

单选题 · easy

  • A.与显存位于同一片区域
  • B.位于 SM 片上← 正确答案
  • C.位于主机内存中
  • D.位于 PCIe 总线中

答案：B。共享内存物理上在流处理器（SM）内部，访问速度接近寄存器，容量远小于全局内存。它对应故事里的“分拣台”，适合做块内数据交换。

Q4. 满足合并访问最关键的条件是？

单选题 · mid

  • A.相邻线程访问相邻地址← 正确答案
  • B.相邻线程访问同一地址
  • C.同一线程访问相邻地址
  • D.随机线程访问同一区域

答案：A。warp 内第 t 个线程访问地址 base+t，顺序对应且连续，一次事务吃满缓存行。只有地址连续而线程映射错乱，或者访问同地址，都无法获得高效的合并访问。

Q5. 共享内存通常划分为多少个 bank？

单选题 · mid

  • A.16 个独立存储 bank
  • B.32 个独立存储 bank← 正确答案
  • C.64 个独立存储 bank
  • D.128 个独立存储 bank

答案：B。共享内存被硬件分为 32 个 bank，与 warp 大小一致。相邻 4 字节轮流落入相邻 bank，设计布局时让 32 个线程命中不同 bank 即可并行无冲突。

Q6. 缓解共享内存 bank conflict 的有效手段有？

多选题 · mid

  • A.给共享内存数组填充偏移列← 正确答案
  • B.让多线程访问同一地址广播← 正确答案
  • C.调整数据布局错开线程映射← 正确答案
  • D.提高全局内存缓存行容量

答案：ABC。填充偏移列让行与行错开 bank；广播让同地址多线程一次取到；调整映射改变线程到 bank 的分布。全局缓存行与共享内存 bank 无关，D 说法错误。

Q7. 以下哪些手段能减少对全局内存的重复访问？

多选题 · mid

  • A.使用共享内存缓存热数据← 正确答案
  • B.用寄存器复用中间结果← 正确答案
  • C.循环展开并重用加载值← 正确答案
  • D.增加循环迭代的总次数

答案：ABC。共享内存、寄存器、循环展开都能让一份数据被反复使用而不重访全局内存。D 只是增加计算量，与减少访存背道而驰。

Q8. 一个 warp 以 stride=4 读取 float 数组（缓存行含 32 个 float），一次访问触发多少条缓存行事务？

单选题 · hard

  • A.1 条缓存行事务
  • B.4 条缓存行事务← 正确答案
  • C.8 条缓存行事务
  • D.32 条缓存行事务

答案：B。32 个线程的地址为 0,4,8,……,124，跨越 0~127 共 128 个 float，正好 128/32=4 条缓存行。这解释了为什么跨步访问会把一次请求放大成 4 次事务。

Q9. Roofline 模型中“计算强度”的定义是？

单选题 · hard

  • A.计算量除以访存量← 正确答案
  • B.访存量除以计算量
  • C.线程数乘以每线程访存量
  • D.带宽除以前端频率

答案：A。计算强度 I = FLOPs / Byte，表示每搬运 1 字节能做多少次计算。与机器拐点比较即可判定是访存受限还是计算受限，指导优化方向。

Q10. 关于共享内存 bank conflict，下列说法正确的有？

多选题 · hard

  • A.不同 bank 之间的访问互不干扰← 正确答案
  • B.同一 bank 多线程访问会被串行化← 正确答案
  • C.多线程读同一地址可走广播不冲突← 正确答案
  • D.填充 padding 不能缓解任何冲突

答案：ABC。不同 bank 并行无冲突；同一 bank 不同地址的访问会排队；同地址走广播机制不算冲突；padding 恰好是消除冲突的经典手段。D 颠倒了事实。