博主介绍:程序喵大人

“来,说说CUDA的线程模型。”

这是面试官开场最常见的问法。多数候选人会顺利背出Thread、Block、Grid三层结构,感觉自己答得很好。然后面试官轻轻一追问:

“Warp是什么?一个Warp多少个线程?如果一个Warp里的线程发生分支,会怎么样?”

很多人在这里就卡住了。

图片

CUDA面试不是考你能不能写出一个Hello Kernel,而是考你对GPU底层执行模型、内存结构和性能特征的理解深度。这篇文章整理了面试官最爱追问的5个基础点,帮你在面试中答得稳。

1. 线程层级:不是背三个单词就行

CUDA的执行模型是一个层级结构:Grid → Block → Thread。但面试官想听的不是这个课本定义,而是你对底层执行单位的理解。

真正执行线程的是GPU上的SM(Streaming Multiprocessor)。一个GPU有多个SM,每个SM可以同时运行多个Block。而内部调度单位是Warp——由32个线程组成,这些线程以SIMT(单指令多线程)方式执行同一条指令。

关键来了:Warp内的所有线程必须执行同一条指令。如果发生分支(比如if-else,部分线程走true、部分走false),GPU会先执行true分支的线程,再执行false分支的线程,另一条件不满足的线程会被置为不活跃状态。这就是Warp Divergence——一个经典的性能杀手。

面试话术:不要只说"有Thread、Block、Grid三层",要补充"调度单位是Warp,32个线程一组,同一Warp内分支会导致性能下降"。

2. 内存层次:面试官最爱画这个图

面试官常坐在那里说:"你来画一下CUDA的内存模型。"这时候能不能立马在白板上画出来,直接决定结果。

CUDA有五层内存结构,从快到慢、从小到大排列:

内存类型 速度 容量 生命周期 可见性
Register 最快 极少(KB级) 线程 单线程
Shared Memory 很快 小(48KB/Block) Block Block内
Constant Memory 快(缓存) 64KB 程序 全局
Texture Memory 快(空间局部性) 较大 程序 全局
Global Memory 最慢 大(GB级) 程序 全局

面试中最常问的是:为什么 Shared Memory 比 Global Memory 快?

因为 Shared Memory 位于芯片内部,每个SM独享,延迟极低;而 Global Memory 是DRAM,距离计算单元远,延迟高。工程上常用的优化策略是:先把数据从Global拷贝到Shared,再在Shared上做计算。

面试话术:能举例说明"报告里常用tile算法,把矩阵分块加载到Shared Memory再计算"。

3. 内存合并访问与 Bank Conflict

这就是我说的"死在第三个问题上"。

面试官接着问:“说说什么是Coalesced Memory Access?如果不合并会怎么样?”

Coalesced Access是指一个Warp内的所有线程访问连续的内存地址。这样GPU可以用一次内存交易读取整个Warp需要的数据。如果访问是跳跃的、散乱的,则需要多次交易,速度直接好几倍差。

更深一层,面试官还会追问 Shared Memory 的 Bank Conflict。Shared Memory 被分成32个Bank(以Fermi架构为例),同一时刻如果多个线程访问同一个Bank,就会产生冲突,访问会被串行化。避开方法包括:让线程访问不同的Bank,或者使用padding增加行偏移。

面试话术:“连续访问可以合并成一次交易;Bank Conflict会使得访问串行化,我们可以通过增加padding避开”。

4. Occupancy:隐藏延迟的关键

面试官不会放过一个经典问题:“什么是Occupancy?越高越好吗?”

Occupancy = 实际活动Warp数 / 最大可能Warp数。它描述了SM的利用率。

关键认知:并不是100%就最好。过高的Occupancy可能意味着每个线程使用更少的Register和Shared Memory,反而可能降低性能。真正重要的是隐藏延迟——当一个Warp在等待内存访问时,SM可以切换到另一个Warp执行。如果只有一个Warp,它就只能干等。

影响Occupancy的三大因素:每个Block的线程数、每个线程的Register数量、每个Block的Shared Memory使用量。面试时不需要背数字,但要能说出这三个因素。

面试话术:“Occupancy描述了SM利用率,但不是越高越好;核心目的是隐藏内存延迟,充分利用Warp切换。”

5. CUDA Stream:异步执行与流水线

最后一个常问点是 Stream。面试官会问:“如果你要让CPU和GPU并行工作,怎么做?”

CUDA Stream 是一个FIFO指令队列。同一Stream内的所有操作按提交顺序执行。使用多个Stream,可以实现计算与数据传输重叠。

典型的模式是:Stream A在做计算的同时,Stream B在传输下一批数据。配合 cudaMemcpyAsync(注意是Async版本,非阻塞版)和固定的pinned host memory,可以最大化GPU利用率。

面试话术:“用多个Stream实现计算与数据传输重叠,让GPU不停机;需要配合cudaMemcpyAsync和pinned memory使用。”

面试速查表

概念 一句话核心
Warp 32线程一组,SIMT执行,分支会导致性能下降
Shared Memory 芯片内部,延迟低,需要注意Bank Conflict
Coalesced Access 连续线程访问连续地址,合并交易提升带宽
Occupancy SM利用率,不是越高越好,重在隐藏延迟
Stream 指令队列,多Stream实现计算与传输重叠

面试CUDA不要求你背下每个参数,但要能从线程层级入手,逐层展开到内存优化。把这5个点说清楚,已经超过了大多数候选人。

码字不易,欢迎大家点赞,关注,评论,谢谢!

Logo

AtomGit 是由开放原子开源基金会联合 CSDN 等生态伙伴共同推出的新一代开源与人工智能协作平台。平台坚持“开放、中立、公益”的理念,把代码托管、模型共享、数据集托管、智能体开发体验和算力服务整合在一起,为开发者提供从开发、训练到部署的一站式体验。

更多推荐