想象你是一家高档餐厅的主厨，面前站着几百位帮厨。你喊一声“切洋葱！”，所有帮厨同时拿起菜刀、同时落刀、同时切出均匀的洋葱丁——没有人单独行动，也没有人需要思考“我该切什么”，因为指令来自你一个人。而隔壁的CPU，则像一位凡事亲力亲为的大厨：他头脑聪明、手法灵活，但他只有一双手，一次只能切一颗洋葱。GPU的秘密就藏在这场“一人指挥、千人动手”的后厨协奏曲里：它用成千上万个简单的小核心，听着同一条指令，批量处理海量数据。而CUDA，就是给这位“后厨总管”写的菜谱——告诉它如何把一场万人宴的配菜工作，拆成一层层的分工与协作。

岗位关联

掌握GPU并行计算与CUDA，意味着你能在手握千核万线程的“超级后厨”里自如调度——这是深度学习训练、科学计算、实时图像处理等高性能AI应用落地的核心技能。

目录
  1. 可视化演示
  2. SIMT并行架构
  3. Warp线程束调度
  4. CUDA线程层次
  5. GPU内存层次
  6. 性能优化技巧
  7. 核函数与SM
  8. 异构协同计算
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

GPU通过SIMT架构实现大规模并行计算，CUDA是NVIDIA GPU的编程模型。本文讲解GPU的SIMT架... 切换模式查看不同维度。

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

GPU SI：核心原理 动态展示关键流程与数据流动 交互控件调节参数实时渲染 f(x) = transform(x) 核心概念的可视化呈现

SIMT并行架构

你再次站上主厨位，面前是几百位帮厨。你喊“切洋葱！”，他们齐刷刷落刀；你喊“削土豆！”，他们又齐刷刷开始削皮。每一位帮厨面前都有自己的食材，但所有人的动作都由你的口令统一指挥——这就叫“单指令、多线程”（Single Instruction Multiple Threads，SIMT）。GPU的核心思想正是如此：与其让一位聪明的大厨（CPU核心）飞快地连续处理一百件事，不如让一百个笨一点但听话的帮厨（GPU核心）同时各做一件简单的事。

SIMT架构中，“指令”是唯一的，执行它的是海量线程。每个线程有自己的数据、自己的寄存器，但同一时刻执行的是同一条机器指令。这一点与SIMD（单指令多数据）不同：SIMD像一位帮厨用一把多格刀具一次切出四片土豆片（数据并行、向量化），而SIMT像一百位帮厨各拿一把刀同时切各自的土豆（线程并行、可独立寻址）。GPU内部由多个流式多处理器（Streaming Multiprocessor，SM）构成，每个SM里又有几十到上百个简单的CUDA核心。评价一个后厨的产出能力，关键指标不是单个帮厨多快，而是整体吞吐量——每秒能切出多少盘菜。

理解了吞吐量，你就明白了为什么GPU的单个核心频率往往低于CPU，却能获得几十倍的算力——正如让一百个普通帮厨同时切菜，远比一位特级厨师单打独斗更快。GPU走的就是“人多力量大”的路线。

吞吐量=NSM×Ncore/SM×f×IPC\displaystyle \text{吞吐量} = N_{\text{SM}} \times N_{\text{core/SM}} \times f \times \text{IPC}
吞吐量等于SM数量乘以每个SM的CUDA核心数，再乘以时钟频率和每周期指令数——就像后厨的出菜速度取决于操作台数量、每台帮厨人数、手速和配合节奏。
  1. 统计核心总数先算出整个GPU里有多少个CUDA核心，这是后厨的总人数。
    NSM×Ncore/SMN_{\text{SM}} \times N_{\text{core/SM}}
  2. 乘上时钟频率频率代表每秒能发出多少条指令，即帮厨的“手速”。
    ff
  3. 乘上每周期指令数IPC（Instructions Per Cycle）代表每个时钟周期真正完成了几条指令，即配合的默契度。
    IPC\text{IPC}
  4. 组合成吞吐量公式把三者相乘，就得到GPU每秒能执行的指令总量。
    NSM×Ncore/SM×f×IPC\displaystyle N_{\text{SM}} \times N_{\text{core/SM}} \times f \times \text{IPC}

Warp线程束调度

帮厨们站成整齐的队列，每32人一队，队长喊“齐步走”，全队必须同时迈出同一只脚。GPU中这32个线程的编队就叫Warp（线程束），它是硬件调度和执行的最小单位。无论你的程序里写了多少个线程，GPU实际运行指令时，总是32个线程一组、齐步执行。这带来一个重要的纪律：Warp内任何一个线程都不能独自“先走一步”。

这种“锁步执行”机制意味着，Warp内的线程在执行同一条指令时，如果某个线程的数据不符合条件需要走不同的路径，整队人都会被拖住——这就是后面要讲的分支发散。而聪明的调度器（Warp Scheduler）有另一个绝招来掩盖等待：当一队帮厨因为去仓库取料（访问内存）而停下时，调度器立刻切换到另一队已经准备好的帮厨继续切菜。这种“哪队能上哪队上”的轮换，叫延迟隐藏（Latency Hiding）。

有了足够的Warp轮换，GPU的每一个CUDA核心都像上了发条一样永不停歇。这正是GPU能“以量取胜”的底层秘密：不是算得快，而是等得少。

Nthread≥D周期×R发射\displaystyle N_{\text{thread}} \ge D_{\text{周期}} \times R_{\text{发射}}
要完全隐藏内存延迟，需要的线程总数至少要达到延迟周期数乘以每周期可发射的指令数——就像后厨里必须备足多队帮厨，一队去取料时，另一队立刻顶上。
  1. 记录内存延迟周期一次内存访问从发出请求到数据返回，CPU周期数记为延迟。
    D周期D_{\text{周期}}
  2. 记录每周期发射率每个周期硬件能发射几条指令，记为发射率。
    R发射R_{\text{发射}}
  3. 两者相乘得出隐藏延迟所需线程数延迟越长，需要的备用队就越多。
    D周期×R发射\displaystyle D_{\text{周期}} \times R_{\text{发射}}
  4. 举例验证若延迟为400周期、每周期发射1条指令，则至少需要400个线程（约13个Warp）才能让调度器总有活干。

CUDA线程层次

假设你要承办一场万人宴。作为后厨总管，你不会让一万个帮厨挤成一团，而是会划分成多个操作台（Block），每个操作台负责一道菜，操作台之间各有分工又要互相配合。整个宴会的后厨体系叫Grid（网格），它包含了这一次任务启动的所有线程；每个操作台是一个Block；操作台上每位帮厨是一个Thread。CUDA编程中，你通过<<<grid, block>>>这样的语法来布置整个后厨。

这个层次结构是逻辑上的，你可以在启动核函数时自由指定，比如<<<128, 256>>>表示128个Block、每个Block有256个线程，总共32768个线程。硬件会自动把这些线程映射到各个SM上。每个线程如何找到自己负责的数据？它需要三个“坐标”：blockIdx（我在第几个操作台）、blockDim（一个操作台有多少人）、threadIdx（我在该操作台的几号位）。三者一组合，就能算出全局唯一的线程ID。

这种“宴会—操作台—帮厨”的三层结构，让GPU程序既能大规模并行，又能灵活协作。Block内线程可以通过共享内存通信、通过同步指令对齐进度；而Block之间相对独立，便于硬件自由调度到任意SM上。

tid=bx×Bx+tx\displaystyle tid = b_x \times B_x + t_x
全局线程ID等于当前Block的编号乘以Block内线程数，再加上线程在Block内的编号——就像你报出“3号操作台第70号帮厨”，全场就能立刻定位到是哪一位。
  1. 确定Block编号当前线程所在Block在Grid中的序号。
    bx=blockIdx.xb_x = \text{blockIdx.x}
  2. 确定Block大小每个Block在x方向上有多少个线程。
    Bx=blockDim.xB_x = \text{blockDim.x}
  3. 确定线程块内编号线程在Block内部的编号。
    tx=threadIdx.xt_x = \text{threadIdx.x}
  4. 组合计算全局ID块编号乘以块大小，再加上块内编号，得到全局唯一线程ID。
    tid=bx×Bx+tx\displaystyle tid = b_x \times B_x + t_x
cuda1 行
kernel<<<Grid数量, Block线程数>>>(参数);

GPU内存层次

后厨里，食材分门别类存放在不同地方：整扇的猪肉在冷库（全局内存）里，离操作台远但容量巨大；每个操作台边上有一个调料架（共享内存），同一操作台的帮厨们伸手就能取用，速度快得多；而每位帮厨手里的菜刀和料理勺（寄存器）则是专属私人物品，速度最快、没有之一。GPU的内存世界，就是这样一座秩序井然的后厨仓库。

全局内存（Global Memory）是GPU中容量最大、最慢的内存，所有线程都可访问，如同冷库向所有操作台开放；共享内存（Shared Memory）是每个Block内部共享的高速存储，容量很小（几十到一百多KB），但速度快了约一个数量级，需要程序员显式管理；寄存器（Register File）分布在SM内部，分配给每个线程私有，速度最快但数量有限。此外还有常量内存和纹理内存，分别像“贴在墙上的菜谱”和“摆盘参考图”，适合广播式的只读访问。一个典型的优化手法是“分块（tile）”：把冷库里的大块食材一次性搬到操作台的调料架旁边，之后所有帮厨都在架上取料，避免反复跑冷库。

正确使用共享内存，是CUDA性能优化的第一课。记住口诀：冷库要少跑，调料架要摆满，手里的刀要随身带——离得越近，拿得越快。

Mtile=⌈NTtile⌉\displaystyle M_{\text{tile}} = \left\lceil \frac{N}{T_{\text{tile}}} \right\rceil
分块访存的次数等于数据总量除以每块能装下的数据量向上取整——就像把一万颗土豆按每筐五百颗分装，一共要搬二十筐，搬完就不用再跑冷库。
  1. 统计数据总量需要处理的数据条数记为数据总量。
    NN
  2. 确定分块容量每个Block的共享内存一次能装下的数据条数记为分块容量。
    TtileT_{\text{tile}}
  3. 相除并向上取整得到需要搬运的块数，每个线程块处理一块。
    ⌈NTtile⌉\displaystyle \left\lceil \frac{N}{T_{\text{tile}}} \right\rceil

性能优化技巧

午市高峰，你喊“切洋葱！”却发现洋葱不够了。30号帮厨被派去剥蒜，可整队32人都在等他——因为Warp是锁步前进的，一个人系鞋带，全队都得停下来。这就是分支发散（Branch Divergence）：当Warp内的线程因为if/else等分支条件走向不同路径时，硬件只能先执行一半线程的路径，再倒回来执行另一半，相当于把一条队伍拆成两支依次走过，总时间直接翻倍。

另一个关键优化是访存合并（Memory Coalescing）。想象帮厨们按取料单顺序排成一列，依次从冷库窗口取货：第0号拿第0格、第1号拿第1格……窗口小哥把相邻的货一口气全递出来，一趟搞定；但如果大家东一个西一个地点货，窗口小哥就得满仓库跑好几趟。GPU访问全局内存时，如果同一个Warp的32个线程访问的是连续的内存地址，硬件会合并成一个或少数几个内存事务，带宽利用率大幅提升。

优化GPU程序，本质上就是管好这两件事：让Warp里的兄弟们走同一条路（避免发散），去同一个窗口按顺序取料（访存合并）。后厨的秩序，就是性能。

Beff=SdataTtrans\displaystyle B_{\text{eff}} = \frac{S_{\text{data}}}{T_{\text{trans}}}
有效带宽等于传输的数据量除以实际发生的内存事务次数——合并访存时一次事务搬回32份数据，有效带宽就是分散访存的好几倍。
  1. 统计传输数据量Warp实际需要的数据总字节数。
    SdataS_{\text{data}}
  2. 统计内存事务数硬件实际发起的内存事务次数。相邻地址能合并为更少事务。
    TtransT_{\text{trans}}
  3. 两者相除得有效带宽事务次数越少，单位时间的有效吞吐越高。
    SdataTtrans\displaystyle \frac{S_{\text{data}}}{T_{\text{trans}}}

核函数与SM

一道“红烧肉”菜谱写好了，把它复印几百份，发到每一个操作台，所有帮厨同时照着做——这就是核函数（Kernel）。在CUDA中，用__global__修饰的函数就是核函数，它在GPU上执行，由<<<grid, block>>>指定启动的线程规模。每个SM就像一个操作台，收到菜谱后，把Block分配给内部的CUDA核心去执行。一块GPU上有几十个SM，每个SM能同时容纳多个Block，共同构成庞大的并行执行体系。

SM（Streaming Multiprocessor，流式多处理器）是GPU物理执行的核心单元。它内部包含：几十到上百个CUDA核心（负责整数和浮点运算）、共享内存、寄存器堆、Warp调度器，以及专门的特殊功能单元。Block被派发到SM后，会被进一步拆分成Warp来调度执行。SM能同时容纳多少线程不是无限的，它受寄存器数量和共享内存容量的约束。衡量SM忙碌程度的指标叫占用率（Occupancy），即活跃的Warp数占最大可容纳Warp数的比例。

写核函数时，不要追求把每一个SM塞满，而要像经营后厨一样讲究平衡：人坐满但调料架不够用，一样做不出菜。合理的Block大小、适度的占用率，才能让性能最大化。

占用率=活动Warp数最大Warp数\displaystyle \text{占用率} = \frac{\text{活动Warp数}}{\text{最大Warp数}}
占用率就像操作台的“上座率”——上座率越高，SM里坐满的帮厨越多，调度器越容易找到替补队来隐藏延迟，但上座率也不是越高越好，太高会挤占共享内存空间。
  1. 统计活动Warp数当前SM上正在轮换执行的Warp数量。
    活动Warp数\text{活动Warp数}
  2. 统计最大Warp数SM硬件支持的最大并发Warp数，由寄存器堆和共享内存容量决定。
    最大Warp数\text{最大Warp数}
  3. 两者相除得到一个0到1之间的比例，即占用率。
    活动Warp数最大Warp数\displaystyle \frac{\text{活动Warp数}}{\text{最大Warp数}}

异构协同计算

我们把视角拉回整家餐厅：老板（CPU）负责接单、排菜、协调各桌客人的节奏，他头脑灵活但手慢；后厨帮工们（GPU）负责批量备菜，手快但只能照着菜谱干活。现代计算机系统里，CPU和GPU各司其职，通过PCIe总线这个“传菜升降机”交换数据，这套协同工作模式叫异构计算（Heterogeneous Computing）。程序员通常在CPU上写主逻辑，把需要海量并行计算的部分（比如矩阵乘法、图像滤波）交给GPU。

异构协同的第一个铁律是“CPU主控，GPU加速”：CPU负责流程控制、数据准备、启动核函数；GPU负责执行大规模数据并行计算。第二个铁律是“先搬运，再计算，再搬回”：数据必须先通过PCIe从主存拷贝到显存，GPU算完后再拷回主存。PCIe的带宽（如PCIe 4.0 x16约为32GB/s）远低于显存带宽（几百GB/s到数TB/s），因此数据搬运的时间开销可能盖过计算收益。

实际工程中，一个常见的经验准则是“算得久不如搬得多”：如果计算量很小、数据量很大，用GPU可能比CPU更慢——炒一盘菜不必启动整个后厨。聪明的工程师会尽量让数据留在显存中反复复用，减少CPU与GPU之间的来回搬运。异构协同的本质，是让对的角色干对的活，并算清传递的代价。

Tcopy=SdataBPCIe\displaystyle T_{\text{copy}} = \frac{S_{\text{data}}}{B_{\text{PCIe}}}
搬运时间等于数据量除以PCIe带宽——就像传菜升降机一次只能装那么多菜，数据量大了，光“上菜下单”就要耗半天。
  1. 统计待搬运数据量需要在CPU和GPU之间传输的数据总字节数。
    SdataS_{\text{data}}
  2. 确定PCIe带宽PCIe总线的理论峰值带宽，实际还会打折扣。
    BPCIeB_{\text{PCIe}}
  3. 相除得到搬运时间用数据量除以带宽，估算拷贝开销。
    SdataBPCIe\displaystyle \frac{S_{\text{data}}}{B_{\text{PCIe}}}

AI 应用场景

深度学习训练加速

阶段：模型训练与推理加速

解决的问题：神经网络训练涉及海量矩阵乘法与卷积运算，CPU串行计算耗时以周为单位；GPU的SIMT架构将矩阵乘法拆解为成千上万个线程并行执行，训练时间从数周压缩到数小时。

具体案例：训练一个GPT类大语言模型时，注意力机制中的矩阵乘（如Q×Kᵀ）动辄是上亿次浮点运算。工程师将矩阵分块（tile）放入共享内存，让Warp内线程按合并访存模式读取数据，配合CUDA的cuBLAS库，在A100 GPU上实现数十TFLOPS的算力利用。这个“后厨备菜”的思路，正是GPU能撑起大模型时代的基石。

分子动力学模拟

阶段：科学计算模拟

解决的问题：模拟数万乃至数百万个原子间的相互作用力（N体问题），每步都要计算所有原子对的距离与受力，计算量随原子数平方增长；GPU并行把每个原子的受力计算分配给一个线程，利用共享内存分块缓存邻居原子坐标，减少全局内存访问。

具体案例：药物研发中筛选候选分子与蛋白质的对接构象，传统CPU模拟需数周。使用CUDA优化的分子动力学软件（如GROMACS），在GPU集群上模拟百万原子体系达到每天微秒级的模拟速度，显著加速候选药物的虚拟筛选。每一个Warp都在“锁步切菜”，切的是原子间的相互作用。

实时视频图像处理

阶段：数据预处理与推理加速

解决的问题：4K视频每秒需处理上亿像素，每帧图像的滤波、增强、目标检测等操作具有天然像素级并行性；GPU按“像素—线程”映射，每个线程处理一个像素点，配合访存合并优化，达到实时处理帧率。

具体案例：直播平台对视频流做美颜、超分与背景替换。一张1080P图像约200万像素，CUDA核函数按<<<ceil(2000000/256), 256>>>启动，将每个像素的RGB值并行调整；通过让相邻线程访问相邻像素地址，访存合并让带宽利用率接近峰值，最终在毫秒级完成全流程处理，支撑起数以万计的同时在线直播。

Python 演示

python53 行
import mathimport osimport timefrom multiprocessing import Pool  def warp_reduce(chunk):    """模拟一个Warp（32个线程）锁步归约求和。"""    total = 0    for val in chunk:        total += val    return total  def block_reduce(array):    """模拟一个Block：把数据切成多个Warp大小的组并行归约。"""    warp_size = 32    chunks = [array[i:i + warp_size] for i in range(0, len(array), warp_size)]    with Pool(processes=os.cpu_count()) as pool:        partials = pool.map(warp_reduce, chunks)    return sum(partials)  def serial_sum(array):    """串行参考实现，模拟CPU单核逐个处理。"""    total = 0    for val in array:        total += val    return total  def main():    n = 1_000_000    data = list(range(1, n + 1))    expected = n * (n + 1) // 2     t0 = time.perf_counter()    s1 = serial_sum(data)    t1 = time.perf_counter()     t2 = time.perf_counter()    s2 = block_reduce(data)    t3 = time.perf_counter()     assert s1 == expected    assert s2 == expected    print(f"串行求和结果: {s1}, 耗时 {t1 - t0:.4f}s")    print(f"模拟GPU并行结果: {s2}, 耗时 {t3 - t2:.4f}s")    print(f"加速比: {(t1 - t0) / (t3 - t2):.2f}x")  if __name__ == "__main__":    main()
目的：用多进程池模拟 GPU 的并行归约求和，展示数据切分、并行计算、结果汇聚的完整流程，与 CUDA 的 Grid/Block/Warp 层次一一对应。
  1. chunks = [array[i:i + warp_size] for i in range(0, len(array), warp_size)]模拟 CUDA 中把数据按 Warp 大小切块，每块由一个"Warp"负责。
  2. with Pool(processes=os.cpu_count()) as pool:模拟 GPU 中多个 SM 并行执行多个 Block 任务。
  3. partials = pool.map(warp_reduce, chunks)每个进程独立完成一块数据的归约，对应各 Warp 锁步累加。
  4. expected = n * (n + 1) // 2用等差数列公式校验结果，确保并行实现没有算错。
  5. if __name__ == "__main__":多进程程序的安全入口，避免递归创建子进程。

常见误区

误区我启动了几百万个线程，GPU 肯定被充分利用了，性能应该非常好。

正解线程数超过硬件可并行的上限后，额外线程只会增加调度开销和资源竞争。GPU 每个 SM 能容纳的线程数受限于寄存器堆和共享内存容量，过量的 Block 会排队等待，甚至导致寄存器溢出到本地内存（实际是全局内存）而大幅减速。更关键的是，隐藏延迟靠的是有意义的 Warp 轮换，而不是盲目堆线程数量。就像后厨面积有限，塞进三千个帮厨只会互相撞手，出菜速度反而下降。

为何常见：人们直观地认为“人越多干活越快”，忽略了硬件资源的物理上限和调度代价。

误区每个线程有自己的寄存器，所以 Warp 里的 32 个线程应该可以各走各的分支，互不干扰。

正解Warp 以锁步方式执行——同一条指令同时广播给 32 个线程。如果出现 if/else 分支且不同线程条件不同，硬件会分两次执行：先让满足条件的线程走 if 路径，其余线程停等；再让剩下的走 else 路径。性能损失约等于分支路径的串行执行。正确做法是尽量让 Warp 内所有线程条件一致，或用数学运算替代分支。

为何常见：把 CPU 多线程的独立执行模型套用到 GPU 上，忽略了 SIMT 共享程序计数器的硬件机制。

误区共享内存和 CPU 的 L1 缓存一样，是硬件自动缓存热数据的，我不用管它。

正解共享内存是 GPU 上由程序员显式控制的片上存储。需要你在核函数里手动声明（如 __shared__ 变量）、手动保证同步（__syncthreads()），数据不会自动从全局内存搬进来。它的价值在于你能精确控制哪些数据留在“操作台调料架”上反复使用。若不加显式管理，编译器只会在寄存器层面做有限的自动优化。共享内存用得好，是性能神器；用不好，反而会因为同步开销和 Bank 冲突拖慢程序。

为何常见：将 CPU 成熟的缓存体系经验直接迁移到 GPU 编程，混淆了“硬件自动缓存”与“用户管理暂存”的本质区别。

误区GPU 算力这么强，把所有算法都扔给 GPU 跑不就完了吗？

正解CPU 擅长复杂流程控制、依赖性强、数据量小的任务；GPU 擅长计算密集、数据并行度高、分支少的任务。且数据从 CPU 到 GPU 需要经过 PCIe 搬运，搬运开销可能抵消全部性能收益。只有在“计算时间 > 搬运时间 + 启动开销”时，异构加速才有意义。工程上常把 GPU 比喻成“后厨帮工队”：让帮工队干只有一道菜的活，光喊集合排队就亏了。

为何常见：只看峰值算力指标，忽略了数据传输瓶颈和任务本身的可并行性，形成了“GPU 万能”的直觉。

面试题

Q1. SIMT 的中文含义是什么？

单选题 · easy

  • A.单指令多数据
  • B.单指令多线程← 正确答案
  • C.多指令多线程
  • D.单线程多指令

答案：B。SIMT 即 Single Instruction Multiple Threads，一条指令同时驱动多个线程执行，就像主厨一声口令全体帮厨齐动手。A 是 SIMD（单指令多数据），属于向量化数据并行，二者机制不同。

Q2. NVIDIA GPU 中一个 Warp 包含多少个线程？

单选题 · easy

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

答案：B。Warp 是 GPU 硬件调度和执行的最小单位，固定为 32 个线程，如同 32 人一队的后厨编组。这是 NVIDIA GPU 的硬件设计约定，程序员无法更改，但可以在启动配置时让 Block 大小为 32 的倍数以充分利用 Warp。

Q3. 当 Warp 内线程出现 if/else 分支发散时，硬件会如何执行？

单选题 · mid

  • A.每个线程独立走自己的分支
  • B.先执行一半线程的路径，再执行另一半← 正确答案
  • C.只执行 if 分支，忽略 else 分支
  • D.编译器自动合并分支，无性能损失

答案：B。Warp 锁步执行指令，遇到发散分支时硬件将路径串行化：先执行条件为真的线程的指令，其余停等；再调头执行条件为假的线程。就像一队帮厨有人要剥蒜，全队等剥完蒜才继续走。因此应避免 Warp 内分支不均匀。

Q4. 访存合并（Memory Coalescing）指的是什么？

单选题 · mid

  • A.相邻线程访问相邻的内存地址← 正确答案
  • B.多个线程共享同一份寄存器
  • C.把全局内存合并到共享内存
  • D.把多个核函数合并成一个

答案：A。访存合并是让同一个 Warp 的线程访问连续内存地址，硬件自动合并为一个或少数内存事务，减少访问次数。好比帮厨按取料单顺序取货，一趟推车全取回；地址跳跃则要跑多趟，带宽利用率显著下降。

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

单选题 · mid

  • A.容量比全局内存大，速度比全局内存慢
  • B.Block 内线程共享，速度远快于全局内存← 正确答案
  • C.所有线程（跨 Block）均可访问
  • D.由硬件自动缓存，无需程序员管理

答案：B。共享内存在 SM 片上，同一 Block 内的线程共享，速度比全局内存快一个数量级，但需要程序员用 __shared__ 显式声明并管理同步。它就像操作台上的调料架，只有本操作台的帮厨够得着，且要自己摆料。

Q6. 某 GPU 访存延迟为 400 周期，每个周期最多发射 1 条指令。要完全隐藏访存延迟，至少约需多少个线程？

单选题 · hard

  • A.40
  • B.200
  • C.400← 正确答案
  • D.4000

答案：C。延迟隐藏所需线程数 = 延迟周期 × 每周期发射指令数 = 400 × 1 = 400。相当于后厨需要备足约 13 个 Warp，一队人去冷库取料时，其余队伍能持续上灶台干活，不让任何一个灶台空转。

Q7. 下列属于 CUDA 线程组织层次的有哪些？

多选题 · mid

  • A.Grid← 正确答案
  • B.Block← 正确答案
  • C.Warp← 正确答案
  • D.SIMD

答案：ABC。CUDA 的逻辑层次是 Grid > Block > Thread，其中 Warp 是硬件隐藏的调度单位（32 线程锁步执行），也是实际运行时的层次。SIMD 是指令级数据并行方式，不是 CUDA 的线程组织层次，故不选。

Q8. GPU 内存层次中包含下列哪些类型？

多选题 · easy

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

答案：ABC。GPU 内存层次从快到慢、从小到大约为：寄存器、共享内存、全局内存，另有常量内存和纹理内存。硬盘属于 CPU 侧的外存层次，不在 GPU 内存层次内，故不选 D。

Q9. 以下哪些做法有助于提高 GPU 计算效率？

多选题 · hard

  • A.大幅增加每个 Block 的共享内存占用
  • B.提高 SM 上的活动 Warp 数量← 正确答案
  • C.允许同一 Warp 内线程频繁走不同分支
  • D.让 Warp 内线程访问连续内存地址← 正确答案

答案：BD。提高活动 Warp 数能增强延迟隐藏能力，访存合并能减少内存事务次数，二者都是经典优化手段。A 会降低 SM 可容纳的 Block 数、拉低占用率；C 的分支发散会让路径串行化，损害性能。

Q10. 启动配置为 kernel<<<64, 512>>>，某线程的 blockIdx.x = 3、threadIdx.x = 70，其全局线程 ID 是多少？

单选题 · extreme

  • A.256
  • B.512
  • C.1606← 正确答案
  • D.2176

答案：C。全局线程 ID = blockIdx.x × blockDim.x + threadIdx.x = 3 × 512 + 70 = 1606。这相当于“3 号操作台，每台 512 人，第 70 号帮厨”在整个后厨中的唯一编号，CUDA 用这个编号来决定每个线程负责哪块数据。