大厂面试官最爱问的CUDA知识点
博主介绍:程序喵大人
- 35 - 资深C/C++/Rust/Android/iOS客户端开发
- 10年大厂工作经验
- 嵌入式/人工智能/自动驾驶/音视频/游戏开发入门级选手
- 《C++20高级编程》《C++23高级编程》等多本书籍著译者
- 更多原创精品文章,首发gzh,见文末
- 👇👇记得订阅专栏,以防走丢👇👇
😉C++基础系列专栏
😃C语言基础系列专栏
🤣C++大佬养成攻略专栏
🤓C++训练营
👉🏻个人网站
“来,说说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个点说清楚,已经超过了大多数候选人。
码字不易,欢迎大家点赞,关注,评论,谢谢!
AtomGit 是由开放原子开源基金会联合 CSDN 等生态伙伴共同推出的新一代开源与人工智能协作平台。平台坚持“开放、中立、公益”的理念,把代码托管、模型共享、数据集托管、智能体开发体验和算力服务整合在一起,为开发者提供从开发、训练到部署的一站式体验。
更多推荐



所有评论(0)