想象你是餐厅后厨的掌勺人：后院有一座巨大的冷藏库，能装下一整季的食材，可每次取货都得穿过大半个院子，来回好几分钟。GPU 的全局内存就是这座冷藏库——容量大得惊人，所有计算单元都能取用，但单次访问要等上几百个时钟周期。聪明的厨师不会一趟趟往院里跑：他们把常用食材放在厨房门口的储物柜里（L2 缓存），把协作要用的料摆上中央拼盘台（共享内存），还会趁取货的空档先切手头的菜（延迟隐藏）。这篇教程带你逛一遍 GPU 的“后厨”：从 DRAM 冷藏库出发，途经 L2 储物柜，来到共享内存拼盘台，最后掌握合并取货、预替补货这些高效后厨的生存法则。学完之后，你会明白 CUDA 程序里那些“显存快慢”的门道，也能读懂 AI 框架为何如此在意数据搬运。

岗位关联

一句话价值：全局内存是 CUDA 程序的数据大动脉；掌握它的延迟特性与缓存、共享内存协作技巧，是 AI 训练与推理性能调优的必修课，也是 GPU 工程师面试的必考点。

目录
  1. 可视化演示
  2. 1. 全局内存概览：后院的大冷藏库
  3. 2. 缓存层次交互：门口的小储物柜
  4. 3. 共享内存缓冲：中央的拼盘台
  5. 4. 延迟隐藏策略：别让厨师闲下来
  6. 5. 访存模式优化：合并取货单
  7. 6. 编程实践要点：打造高效后厨
  8. AI 应用场景
  9. Python 演示
  10. 常见误区
  11. 面试题

可视化演示

三步展示GPU存储层次：金字塔结构（Reg→L1→Shm→L2→Glb延迟递增）、合并访问（coalesced 1次事务vs strided 8次）、延迟对比（1→30→200→600 cycles柱状图）。

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

GPU存储层次 Reg:寄存器（1周期，最快） Shm:共享内存（30周期，SRAM） Glb:全局内存（600周期，DRAM） latency: Reg(1) < L1/Shm(30) < L2(200) < Global(600) 从寄存器到全局内存延迟递增数百倍。

1. 全局内存概览：后院的大冷藏库

假设你是一家餐厅的总厨师长。后院有一座巨大的冷藏库，存着全店一个月的食材，从牛排到花椒面应有尽有——可它离灶台太远了，每次取料都要穿过院子、推开厚重的库门，来回耗上好几分钟。GPU 的全局内存（Global Memory）就是这座“后院冷藏库”：它容量最大、所有人都能用，但离计算单元最远、走得也最慢。

在 CUDA 编程模型里，全局内存是 GPU 的“主储物间”，物理上由 DRAM（动态随机存取存储器）实现。CPU 送进来的输入数据、GPU 算完要传走的结果，都先住在这里；每个线程都可以读写它。它和 CPU 内存最大的不同是位置：显存离 GPU 近，CPU 内存离 CPU 近，两者之间靠 PCIe 总线搬运。全局内存的优点是“装得下”——轻松几个 GB 到几十 GB；缺点是“走得慢”——单次访问延迟高达 400 到 800 个时钟周期。

虽然全局内存的带宽很高，现代显卡能做到数百 GB/s 甚至上 TB/s，但带宽像传送带的输送速度，延迟像“走过去再走回来”的固定脚程——传送带再快，脚程也省不掉。小数据量的随机访问，瓶颈几乎全在延迟上。这正是全局内存区别于寄存器、共享内存的根本特性：容量大、带宽高、延迟长。

理解了冷藏库的脾性，后厨的规矩就清楚了：尽量一次多取、按需取够，别一趟一趟瞎跑。CUDA 程序也一样——减少访问次数、加大单次访问的数据量、让数据在近处待着，是驾驭全局内存的第一课。

Taccess=Tlatency+SdataBWpeak\displaystyle T_{access} = T_{latency} + \frac{S_{data}}{BW_{peak}}
一次全局内存访问的总耗时，等于固定的往返延迟，加上按峰值带宽把数据搬完所需的时间。数据量越大，传输时间越显著；但哪怕只取一个字节，延迟也躲不掉。
  1. 总访问时间 $T_{access}$从线程发出读请求到拿到数据的总时长，单位是时钟周期，是整个存储性能讨论的起点。
    TaccessT_{access}
  2. 固定延迟 $T_{latency}$无论数据多大都要付出的“脚程”，包括地址解析、DRAM 行激活等开销，约 400~800 个周期。
    TlatencyT_{latency}
  3. 数据传输量 $S_{data}$要搬运的字节数，比如一个 float 是 4 字节，一个 float4 是 16 字节。
    SdataS_{data}
  4. 峰值带宽 $BW_{peak}$显存每秒最多能搬多少字节，单位是 GB/s，反映传送带的输送速度。
    BWpeakBW_{peak}
  5. 除法 $\frac{S_{data}}{BW_{peak}}$用带宽把数据“抬上传输带”所需的时间，数据越大耗时越明显。
    SdataBWpeak\frac{S_{data}}{BW_{peak}}

2. 缓存层次交互：门口的小储物柜

老板发现厨师们翻来覆去只用那几十种常用食材，于是在厨房门口加了一排自动补货的小储物柜：常用料随用随取，缺了会有专人从后院补齐。这排储物柜就是 L2 缓存——全局内存的统一缓存层，也是 GPU 硬件自动管理的“贴心设计”。

所有对全局内存的访问，都要先敲 L2 的门。数据在柜子里，叫“命中”（hit），一伸手就拿走；不在柜子里，叫“缺失”（miss），L2 就得替你去后院取，同时把取回的东西在柜子里留一份，方便下次。L2 的容量从几 MB 到几十 MB，比全局内存小两三个数量级，但命中时的延迟只有全局内存的几分之一——约 200 个周期对 700 个周期。

对程序员来说，L2 是透明的，你无法直接控制它，但可以“迎合”它：让同一个线程块反复访问同一批数据，提高时间局部性；让数据排布紧凑，提高空间局部性。两者都能让命中率上升。平均访存延迟可以用命中率加权来估算：命中率高，平均延迟就贴近 L2 的快速通道；命中率低，就不得不频繁走 DRAM 的远路。

储物柜的补货策略由硬件决定，但厨师的使用习惯由你决定。想让 L2 命中率变高，就从数据布局和访问顺序下手——后厨的“顺手程度”是可以设计的。

Tavg=h×TL2+(1−h)×TDRAM\displaystyle T_{avg} = h \times T_{L2} + (1 - h) \times T_{DRAM}
平均访存延迟 = 命中率 × L2 命中延迟 + 缺失率 × DRAM 延迟。这个加权公式解释了为什么“数据局部性”是缓存性能的命门。
  1. 平均延迟 $T_{avg}$长期统计下，一次全局内存访问的期望耗时，是判断程序访存健康度的核心指标。
    TavgT_{avg}
  2. 命中率 $h$数据在 L2 缓存中被直接找到的概率，取值范围是 0 到 1，越高越好。
    hh
  3. L2 命中延迟 $T_{L2}$从 L2 中直接拿到数据的花费，约 200 个时钟周期，是“近路”。
    TL2T_{L2}
  4. 缺失率 $(1 - h)$数据不在 L2 中、必须去 DRAM 取的概率，与命中率互补。
    (1−h)(1 - h)
  5. DRAM 延迟 $T_{DRAM}$去全局内存（DRAM）取数的花费，约 700 个时钟周期，是“远路”。
    TDRAMT_{DRAM}
  6. 加权求和 $\times$两条路径按各自概率加权后相加，就是长期平均下来的访存成本。
    h×TL2+(1−h)×TDRAMh \times T_{L2} + (1 - h) \times T_{DRAM}

3. 共享内存缓冲：中央的拼盘台

光有储物柜还不够。几位厨师做同一道菜，经常要交换手头的配料：蒜末、葱花、调好的酱汁。与其每人跑一趟柜子，不如在案台中央摆一张拼盘台，大家把料往台上一放，一弯腰就能拿到。这张拼盘台就是共享内存（Shared Memory）。

共享内存是 GPU 芯片上的一块 SRAM，紧挨着计算单元，延迟只有几个时钟周期，比全局内存快一两个数量级。但它和 L2 有本质区别：L2 是硬件自动管理的缓存，共享内存是程序员管理的舞台——你用 __shared__ 关键字声明一块区域，自己负责把数据从全局内存搬进来、用完再搬回去。它是拼盘台，不是自动补货柜，台面上的料必须自己摆。

共享内存按线程块隔离：一个块只能访问自己那块“拼盘台”，块与块之间不能直接共享。因为是手动管理，使用前还要用 __syncthreads() 同步——就像大家约定“我放好了你再拿，别拿了空碟子”。每块共享内存容量只有几十 KB，但配合“一次搬运、全员复用”的用法，常常带来数量级的加速。

拼盘台的灵魂是协作：一个人去后院取一大块肉，全桌人分着用，而不是每人各跑一趟。共享内存要的就是这种“众筹搬运”的思维——把高频复用的数据放在离计算最近的地方。

Sspeedup≈TGMTSM≈10∼100\displaystyle S_{speedup} \approx \frac{T_{GM}}{T_{SM}} \approx 10 \sim 100
用共享内存替换等量的全局内存访问，理论加速比为两者延迟之比，量级在 10 到 100 倍之间。当然，实际收益还要扣除搬运与同步的开销。
  1. 加速比 $S_{speedup}$把某些全局内存访问改成共享内存访问后获得的提速倍数。
    SspeedupS_{speedup}
  2. 全局内存延迟 $T_{GM}$一次全局内存访问约 400~800 个时钟周期，是分母上的“慢基准”。
    TGMT_{GM}
  3. 共享内存延迟 $T_{SM}$一次共享内存访问仅数个时钟周期，是“快基准”。
    TSMT_{SM}
  4. 比值 $\frac{T_{GM}}{T_{SM}}$两者相除得到量级差距，约 10 到 100 倍。
    TGMTSM\frac{T_{GM}}{T_{SM}}
  5. 近似符号 $\approx$表示这是粗略量级估计，实际值因硬件代际、访问冲突和同步开销而异。
    ≈10∼100\approx 10 \sim 100

4. 延迟隐藏策略：别让厨师闲下来

有经验的厨师从不干等。他去后院取货的同时，手里继续切着案板上的菜；等货到了，再腾出手处理新菜。这叫“人不等货”。GPU 隐藏全局内存延迟，靠的是同一个道理：当一批线程在等数据时，立刻把计算单元让给另一批线程干活。

这个机制在 CUDA 里叫“延迟隐藏”。GPU 一次调度一个 warp（32 个线程的编队），当某个 warp 因为访存被卡住时，调度器立刻换上另一个能算的 warp。只要计算单元上有“闲不住”的线程可换，访存延迟就被掩盖住了。隐藏效果取决于占用率——同时驻留的线程越多，可供切换的“备胎”就越多。

除了靠硬件调度，程序员还可以主动隐藏：数据预取（提前把下一轮要用的数据搬进共享内存或寄存器）和计算访存重叠（让算法在等数据时先算别的）。两者本质都是“把取货的时间花在切菜上”。预取特别适合有规律的数据流，比如图像分块处理时提前加载下一块，让数据在路上跑的时候，计算已经在进行。

延迟隐藏的本质是“换人干活”：一个人去取货，另一个人顶上。线程越多、预取越勤，GPU 的等待时间就越短，后厨的流水线就越流畅。

Nwarps≈LmemoryLcompute\displaystyle N_{warps} \approx \frac{L_{memory}}{L_{compute}}
要隐藏访存延迟，大约需要多少个线程束？把访存延迟除以计算延迟即可。这个比值告诉你：等待有多长，就需要多少“备胎”来轮流顶班。
  1. 线程束数量 $N_{warps}$为了填满访存等待期，需要的可切换 warp 数量下限。
    NwarpsN_{warps}
  2. 访存延迟 $L_{memory}$一次全局内存访问的延迟，约 400~800 个周期，是“等待的时长”。
    LmemoryL_{memory}
  3. 计算延迟 $L_{compute}$一次算术或逻辑运算的延迟，通常几个周期，是“单位干活的时长”。
    LcomputeL_{compute}
  4. 除法 $\frac{L_{memory}}{L_{compute}}$等待时长是干活时长的多少倍，就需要多少个 warp 轮流顶上。
    LmemoryLcompute\frac{L_{memory}}{L_{compute}}
  5. 近似 $\approx$实际还受吞吐、资源限制影响，这里只给数量级的参考。
    ≈\approx

5. 访存模式优化：合并取货单

假设一道菜要用十种食材，散落在冷藏库的不同角落。笨办法是十位厨师各跑一趟，每人取一样——十个人各花几分钟在走路上。聪明的办法是：把取货单按货架顺序排好，让一个采购员推着车沿货架一趟取完。CUDA 的内存合并（Coalescing）就是这个道理。

当一个 warp 的 32 个线程访问连续的内存地址时，硬件把这 32 次访问合并成少数几个内存事务，一次搬回整段数据。相反，如果线程各访问各的散落地址，硬件就得拆成多次事务，带宽被白白浪费在“跑空车”上。合并访存是全局内存最立竿见影的优化手段——通常只是把索引从 i * stride 改成 i，带宽利用就能翻好几倍。

合并访存要求地址连续且对齐：理想情况下，32 个线程访问 32×4=128 字节的连续区域，一次事务完成。如果访问跨度大，效率就按跨度成比例下降。用公式衡量：有用字节数除以实际搬运字节数，越接近 1 说明每一趟车都装满了有用的货。

合并访存是“性价比之王”：往往只是调整数据布局或索引写法，带宽利用就能翻好几倍——相当于给冷藏库的取货单重新排了序，一趟车装满了再走。

Ecoalesce=SusefulStransferred\displaystyle E_{coalesce} = \frac{S_{useful}}{S_{transferred}}
合并访存效率 = 真正用到的字节数 ÷ 硬件实际搬运的字节数。这个比值衡量带宽有没有被浪费在“空车”上。
  1. 合并效率 $E_{coalesce}$0 到 1 之间的比值，越接近 1 说明带宽利用越充分。
    EcoalesceE_{coalesce}
  2. 有用字节 $S_{useful}$真正被线程消费的数据量，比如 32 个线程各读 4 字节，共 128 字节。
    SusefulS_{useful}
  3. 实际搬运字节 $S_{transferred}$GPU 从显存实际取回的字节数，包含被浪费的填充与重复数据。
    StransferredS_{transferred}
  4. 比值 $\frac{S_{useful}}{S_{transferred}}$有用货载占整车货载的比例，是带宽效率的直观度量。
    SusefulStransferred\frac{S_{useful}}{S_{transferred}}
  5. 优化目标让一个 warp 访问 128 字节连续且对齐的区域，效率趋近于 1。
    Suseful→StransferredS_{useful} \rightarrow S_{transferred}

6. 编程实践要点：打造高效后厨

高效的后厨，光有冷藏库、储物柜和拼盘台还不够，还得有章法：什么料放柜子、什么料上桌、采购频率多高，都要心里有数。CUDA 程序的存储策略也一样，需要“分层设计”。

分层设计的口诀是：最常用、每个线程独有的数据放寄存器；块内协作、需要反复读的数据放共享内存；大块、低频率访问的数据放全局内存；只读数据可以走常量内存或纹理内存等专门路径。再配合性能分析工具（NVIDIA Nsight Compute 等）查看访存吞吐、L2 命中率和 stall 原因，找到真正的瓶颈再动手——先度量，再优化。

还要注意资源平衡：共享内存和寄存器都是稀缺资源，用得太多会拉低占用率，削弱延迟隐藏，甚至导致寄存器溢出（spill）到慢速的局部内存。优化收益还受 Amdahl 定律约束：如果访存占了 90% 的时间，即使把访存速度提升 10 倍，整体加速也只有约 5 倍。所以要优先优化真正占时间的环节。

后厨优化没有银弹：先度量（哪个环节最慢），再动手（改进最慢的环节），最后复度量。CUDA 存储优化同样是“度量—优化—再度量”的循环，而全局内存永远是值得第一个审视的地方。

Soverall=1(1−p)+pk\displaystyle S_{overall} = \frac{1}{(1 - p) + \frac{p}{k}}
整体加速比 = 1 ÷（不可优化的占比 + 可优化部分提升后的占比）。它提醒你：把力气花在最花时间的环节上，而不是最容易改的环节上。
  1. 整体加速比 $S_{overall}$优化后程序相对于优化前的总提速倍数。
    SoverallS_{overall}
  2. 被优化部分占比 $p$访存等被改进环节在总耗时中的比例，比如 0.9 表示占了 90%。
    pp
  3. 不可优化占比 $(1 - p)$剩下的、无法受益于本次优化的时间比例。
    (1−p)(1 - p)
  4. 加速倍数 $k$被优化环节自身提升了多少倍，比如访存快了 10 倍。
    kk
  5. 分母 $\frac{p}{k}$优化后，被优化部分的新耗时是原来的 $1/k$，与不变部分相加就是新总耗时，取倒数即加速比。
    1(1−p)+pk\frac{1}{(1 - p) + \frac{p}{k}}

AI 应用场景

场景 1：深度学习训练中的批量数据搬运

阶段：数据加载与训练性能优化

解决的问题：训练时 CPU 预处理好的数据要先拷入 GPU 全局内存，再由 kernel 读取。如果 kernel 频繁小批量随机访问全局内存，数百周期的延迟会严重拖慢训练；传统写法往往在数据搬运上浪费大量时间。利用共享内存暂存数据块 + 预取策略，可以让计算单元持续忙碌，而不是等数据。

具体案例：在 BERT 等大模型训练中，矩阵乘法（GEMM）和 embedding 查询是访存大户。采用共享内存分块 GEMM（tiled GEMM），配合数据预取与 CUDA Graphs 合并内核启动，端到端训练吞吐可提升 30% 以上。

场景 2：图像与视频处理中的滤波加速

阶段：预处理管线（图像缩放、滤波、归一化）

解决的问题：图像滤波的每个输出像素都要读周围邻居，若每个线程直接访问全局内存，相邻线程会重复读取大量像素，浪费带宽。传统实现受限于访存瓶颈，吞吐上不去。用共享内存保存图像块（tile），配合边界 halo 处理，把重复读取降到最低。

具体案例：在批量图像归一化与高斯模糊管线中，将 32×32 像素块载入共享内存，相邻线程复用边缘像素，预处理吞吐提升 2~3 倍，卷积类 AI 模型的预处理瓶颈显著缓解。

场景 3：向量检索中的带宽优化

阶段：RAG 系统的向量相似度检索

解决的问题：百万级向量库做余弦相似度时，需要流式读入大量向量数据，访存带宽成为吞吐瓶颈。传统实现若按线程随机取向量，非合并访问浪费带宽。通过合并访存 + 共享内存暂存查询向量，带宽利用率大幅提高。

具体案例：在 500 万条 768 维向量的检索服务中，把查询向量放入共享内存、底库向量按连续地址合并读取，单卡 QPS 提升约 2 倍，检索延迟明显下降。

Python 演示

python45 行
# global_memory_sim.py# 模拟朴素版与共享内存分块版矩阵乘的访存延迟差 DRAM_LAT = 700    # 全局内存(DRAM)延迟, 时钟周期L2_HIT_LAT = 200  # L2 缓存命中延迟, 时钟周期SM_LAT = 10       # 共享内存访问延迟, 时钟周期  def naive_cycles(n, tile):    """朴素版: 每次乘加都直接读全局内存"""    total = 0    for bi in range(0, n, tile):        for bj in range(0, n, tile):            for bk in range(0, n, tile):                for i in range(bi, bi + tile):                    for j in range(bj, bj + tile):                        for k in range(bk, bk + tile):                            total += 2 * DRAM_LAT   # 读 A 和 B 两个数    return total  def tiled_cycles(n, tile):    """分块版: 每个 tile 只搬一次, 计算时读共享内存"""    total = 0    for bi in range(0, n, tile):        for bj in range(0, n, tile):            for bk in range(0, n, tile):                # 把 A 的子块和 B 的子块搬进共享内存                total += 2 * tile * tile * L2_HIT_LAT                # 块内计算: tile 乘 tile 个输出, 每个做 tile 次乘加                total += tile * tile * tile * SM_LAT    return total  def main():    n, tile = 64, 16    naive = naive_cycles(n, tile)    tiled = tiled_cycles(n, tile)    print(f"朴素版全局内存直访: {naive:,} 周期")    print(f"共享内存分块版:     {tiled:,} 周期")    print(f"理论加速比:         {naive / tiled:.1f} 倍")  if __name__ == "__main__":    main()
目的：用延迟数字模拟 CUDA 矩阵乘法的两种访存策略，直观展示共享内存分块为什么能带来数量级加速。
  1. total += 2 * DRAM_LAT # 读 A 和 B 两个数朴素版每次乘加都要从全局内存取两个数，按 700 周期计入，代价高昂。
  2. total += 2 * tile * tile * L2_HIT_LAT分块版每块只需把两个 16×16 的 tile 搬进共享内存，按 L2 命中延迟计。
  3. total += tile * tile * tile * SM_LAT块内计算改为读共享内存，每次只花 10 个周期，是加速的核心来源。
  4. for bi in range(0, n, tile)把 64×64 矩阵切成 4×4 个 tile，这是分块算法（tiling）的骨架。
  5. naive / tiled两种策略总延迟之比，就是共享内存暂存带来的理论加速比。

常见误区

误区“我的显卡显存 24GB，比朋友的 8GB 大三倍，训练一定快三倍吧？”

正解容量只决定“装不装得下”，不决定“跑多快”。AI 训练往往卡在带宽和延迟上：显存再大，如果线程随机访问全局内存、没有合并访存，带宽利用率会低得可怜。就像冷藏库修得再大，厨师一趟趟空跑也快不起来。选卡看容量，更要看带宽、算力和缓存大小。

为何常见：厂商宣传单上“大显存”最显眼，容量也最好理解，性能反而被下意识等同于容量。

误区“我声明了 `__shared__` 数组，硬件应该会自动帮我缓存全局内存的数据吧？”

正解不是。共享内存和 L2 缓存完全是两回事：L2 是硬件透明管理的缓存，共享内存则是程序员显式搬运的 SRAM。你需要亲手写“从全局内存读入 → 写进共享内存 → __syncthreads() → 再读共享内存”。它是拼盘台，不是自动补货柜——台面上的料必须自己摆。

为何常见：CPU 的缓存完全透明，新手大脑里“硬件自动优化”的预期太强，容易把这个设定搬到 GPU 上。

误区“全局内存延迟这么高，有没有办法把所有数据都塞进共享内存？”

正解物理上不可能。共享内存每线程块只有几十 KB，寄存器更少，而全局内存有几十 GB。关键不是“不用”，而是“少用、用得好”：大块数据放全局内存，用合并访问提高带宽效率，用共享内存暂存高频复用的数据，用预取把延迟藏起来。

为何常见：新手刚被 800 个周期的延迟数字吓到，容易忘掉容量层级的物理现实——数据总得有地方安放。

误区“延迟要靠线程来藏，那把 GPU 的线程全部拉满，一定最快！”

正解线程数超出某个平衡点后，性能不升反降。原因在于资源守恒：共享内存、寄存器总量固定，线程太多会导致每线程可用资源减少、占用率排不进去，甚至寄存器溢出到慢速的局部内存。合理做法是用占用率分析，找到“刚好藏住延迟”的平衡点。

为何常见：“大力出奇迹”的直觉在并行世界里很诱人，但资源守恒定律不会因为冲动而改变。

面试题

Q1. 全局内存（Global Memory）在物理上主要由什么器件实现？

单选题 · easy

  • A.SRAM 高速缓存颗粒
  • B.DRAM 动态存储颗粒← 正确答案
  • C.Flash 闪存芯片颗粒
  • D.EEPROM 寄存器颗粒

答案：B。全局内存由 DRAM 实现，容量大但延迟高达数百周期；SRAM 用于共享内存和片上缓存，Flash 用于持久化存储，EEPROM 与显存无关。这正是“为什么全局内存慢”的物理根源。

Q2. 访问一次全局内存通常需要多少个时钟周期？

单选题 · easy

  • A.大约 2 到 5 个周期
  • B.大约 40 到 80 个周期
  • C.大约 400 到 800 个周期← 正确答案
  • D.高达 4000 到 8000 周期

答案：C。全局内存（DRAM）的访存延迟约为数百个时钟周期，是共享内存的几十到上百倍。这正是延迟隐藏和共享内存暂存策略存在的根本原因。

Q3. 在 GPU 存储层次中，哪个部件是全局内存的统一缓存层？

单选题 · mid

  • A.每个线程私有的 L1 缓存
  • B.全局统一的 L2 缓存← 正确答案
  • C.每线程块独立的共享内存
  • D.只读专用的纹理路径缓存

答案：B。L2 缓存是所有全局内存访问的统一汇聚点，命中时延迟从约 700 周期降到约 200 周期；L1 和共享内存按线程块隔离，纹理缓存只服务只读路径。

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

单选题 · mid

  • A.硬件自动管理数据的存取
  • B.程序员显式管理且延迟极低← 正确答案
  • C.容量与全局内存相当
  • D.所有线程块都可直接访问

答案：B。共享内存是程序员手动管理的 SRAM，延迟仅几个周期，但每块容量只有几十 KB，且只对本线程块可见，使用前需要同步。

Q5. 什么是“内存合并（Coalescing）”？

单选题 · mid

  • A.同一 warp 访问连续地址被合并为少量事务← 正确答案
  • B.不同 warp 的数据被合并进同一缓存行
  • C.多个内核函数被合并为一次启动
  • D.显存与主机内存合并为一个地址空间

答案：A。合并访存指硬件把同一 warp 对连续地址的访问合并成少数内存事务，最大化带宽利用；其余选项描述的是缓存行共享、kernel 合并等无关概念。

Q6. 下列哪些策略可以有效隐藏全局内存延迟？

多选题 · mid

  • A.提高占用率增加可切换的线程束← 正确答案
  • B.数据预取到共享内存或寄存器← 正确答案
  • C.减少内核函数的启动次数
  • D.让计算与访存操作重叠执行← 正确答案

答案：ABD。延迟隐藏靠“等待时换人干活”：占用率高则切换备胎充足，预取和计算访存重叠让等待期有事可做；减少启动次数主要省的是调度开销，而非访存延迟。

Q7. 平均访存延迟公式 T = h × T1 + (1 − h) × T2 中，h 的含义是？

单选题 · hard

  • A.共享内存的数据命中率
  • B.L2 缓存的数据命中率← 正确答案
  • C.全局内存的带宽利用率
  • D.当前线程块的占用率

答案：B。h 是 L2 命中率，T1、T2 分别是 L2 命中与缺失（去 DRAM）的延迟；两者按概率加权得到平均访存延迟，这解释了为何数据局部性设计如此重要。

Q8. 线程块内数据就绪前，通常需要执行哪类操作来同步？

单选题 · easy

  • A.块内同步屏障指令← 正确答案
  • B.全局内存栅栏指令
  • C.设备级同步函数
  • D.原子加锁保护操作

答案：A。CUDA 中线程块内用 __syncthreads() 建立屏障，确保所有线程写完共享内存后再读取；栅栏、设备级同步和原子操作针对的是跨块或全局顺序问题。

Q9. 以下哪些做法能提升全局内存访问效率？

多选题 · hard

  • A.让线程按连续地址访问← 正确答案
  • B.保持访问地址对齐到 128 字节← 正确答案
  • C.用共享内存暂存可复用的数据← 正确答案
  • D.每次访问减少请求的数据量

答案：ABC。连续且对齐的访问触发合并事务，共享内存暂存减少重复访问全局内存；减少单次请求数据量反而降低带宽利用效率，与优化目标相反。

Q10. 共享内存与全局内存的访问延迟差距大约是？

单选题 · mid

  • A.两者完全相同没有差距
  • B.共享内存快 10 到 100 倍← 正确答案
  • C.全局内存快 10 到 100 倍
  • D.两者都由 DRAM 技术实现

答案：B。共享内存（SRAM）延迟约几个周期，全局内存（DRAM）约数百周期，量级上相差 10~100 倍。这正是把高频复用数据搬进共享内存的核心动机。