1. 英伟达暑期实习笔试到底在考什么聊英伟达暑期实习笔试之前先得把一件事说清楚这家公司的笔试跟国内大多数互联网大厂的套路完全不一样。国内很多公司的笔试题偏向八股文和LeetCode中等难度偏上的算法题刷够题量基本能应付。但英伟达的笔试尤其是涉及深度学习、GPU计算方向的岗位题目设计明显更偏向“底层理解工程直觉”的组合拳。我带过几届准备英伟达实习的学生也跟几位已经进去的朋友聊过大家的共识是笔试的区分度不在你刷了多少题而在你对并行计算模型和深度学习底层机制的理解深度。先说说笔试的基本盘。英伟达暑期实习的笔试通常分两部分一部分是通用编程能力另一部分是方向相关的专业题。通用编程部分不会出特别刁钻的算法但会考察你对C内存模型、指针操作、多线程基础的理解——这跟GPU编程的底层逻辑是一脉相承的。专业题部分则根据你投递的岗位方向有较大差异比如深度学习方向会考CNN的计算量分析、反向传播的链式法则推导、优化器的选择逻辑GPU计算方向会考线程层次结构、内存层次、warp调度机制系统方向可能涉及驱动模型和中断处理。为什么英伟达要这么设计因为GPU编程本质上是一个“资源受限下的并行调度”问题。你在写CUDA kernel的时候脑子里必须同时装着一件事这个计算任务怎么拆成几千个线程每个线程怎么访问内存线程之间怎么同步。这种思维方式跟传统的串行编程完全不同所以笔试题目会刻意考察你在这方面的直觉。我见过很多同学准备英伟达笔试的方式是疯狂刷LeetCode结果笔试一出来就懵了——因为题目根本不是那个路子。比如有一道经典题是给出一段CUDA代码问shared memory的bank conflict发生在哪里怎么优化。这种题你刷一万道LeetCode也遇不到但如果你理解shared memory的bank划分机制看一眼就能定位问题。所以准备英伟达笔试的第一步是把复习重心从算法题转移到并行计算基础和深度学习原理上。具体来说你需要搞清楚这几件事GPU的线程层次结构thread、block、grid是怎么组织的warp是什么概念shared memory和global memory的访问延迟差多少CNN的卷积操作在GPU上怎么并行化反向传播中哪些步骤可以并行哪些必须串行。这些内容在后面我会逐一展开。还有一点值得注意英伟达的笔试题目往往会给出一个具体的工程场景让你分析性能瓶颈或者选择方案。比如“给定一个矩阵乘法任务在什么样的矩阵尺寸下使用tiling技术能带来最大收益”——这种题没有标准答案考的是你的分析框架和权衡能力。你得能从内存带宽、计算强度、occupancy这几个维度去拆解问题而不是背一个结论。2. 从一道矩阵乘法题看GPU线程层次的设计逻辑矩阵乘法是GPU计算里最经典的入门案例也是英伟达笔试出现频率最高的题型之一。但笔试不会只让你写一个naive的矩阵乘法kernel而是会围绕它做各种变体让你分析性能、让你优化、让你解释为什么某种写法更快。要答好这类题你得先把GPU的线程层次结构吃透。2.1 thread、block、grid三层结构到底怎么理解CUDA的线程组织分三层thread是最小的执行单元block是一组thread的集合grid是一组block的集合。这个结构看起来简单但很多人在笔试里栽跟头是因为没有理解这三层结构跟硬件执行单元的映射关系。打个比方grid就像是一个大工程项目block是项目里的各个施工队thread是施工队里的工人。每个施工队有自己的临时仓库shared memory工人之间可以快速交换材料但不同施工队之间要交换材料就得走公司的大仓库global memory速度慢很多。这个类比的关键在于同一个block内的thread可以通过shared memory高效通信不同block之间只能通过global memory通信而且不能保证执行顺序。笔试里经常考的一个点是为什么block的大小通常设置为128或256个thread这跟硬件的一个关键参数有关——warp size。在英伟达的GPU架构中一个warp固定包含32个thread这是硬件调度的基本单位。如果你把block大小设成100那硬件会把它拆成4个warp3232324最后一个warp只有4个thread在干活28个thread的资源被浪费了。所以block大小最好是32的整数倍而且通常选128或256这样既能保证warp满载又不会因为block太大导致register压力过高。2.2 矩阵乘法中tiling技术的来龙去脉Naive的矩阵乘法kernel是这样的每个thread负责计算输出矩阵C的一个元素它需要读取A的一整行和B的一整列做点积。这个写法的问题在于A的每一行会被读取N次N是输出矩阵的列数B的每一列会被读取M次。global memory的访问量是O(MNK)级别的而实际计算量也是O(MNK)计算强度只有1 FLOP/byte左右远远达不到GPU的计算峰值。Tiling技术的核心思想是把矩阵分块让每个block负责计算一个tile的输出block内的thread协作把A和B的对应tile加载到shared memory里然后从shared memory里读取数据做计算。这样global memory的访问量降到了O(MNK/tile_size)级别计算强度提升了tile_size倍。笔试里常见的问法是给定矩阵尺寸和tile大小计算shared memory的使用量判断会不会超出硬件限制。比如tile大小是16x16数据类型是float4字节那么A的tile和B的tile各需要161641024字节总共2048字节。英伟达GPU的shared memory每个block通常有48KB到164KB不等取决于架构所以这个用量完全没问题。但如果tile大小是64x64那就需要64644*232KB在一些老架构上就可能超限了。还有一个容易忽略的点tiling之后每个thread的计算量增加了。Naive版本里每个thread算一个输出元素tiling版本里每个thread通常算一个tile内的多个元素。这意味着register的使用量会增加如果register不够occupancy就会下降。笔试里如果考到“为什么tile不是越大越好”你就得从shared memory容量、register压力、occupancy这三个角度去回答。2.3 笔试中关于线程层次的常见陷阱题有一类题特别阴险给出一段kernel代码问“这段代码有没有问题”。问题往往出在线程同步上。比如在tiling的矩阵乘法中加载完shared memory之后必须调用__syncthreads()否则某些thread可能还没写完shared memory其他thread就开始读了。但__syncthreads()只能同步同一个block内的thread如果你在代码里跨block做了某种假设那就错了。另一个常见陷阱是线程索引越界。当矩阵尺寸不是tile大小的整数倍时边界上的thread会访问到矩阵之外的内存。笔试里可能会让你写出边界检查的代码或者问你如果不做边界检查会发生什么。答案不是“程序崩溃”这么简单——在GPU上越界访问可能读到垃圾数据导致计算结果错误但不报错这种bug最难排查。我个人的经验是准备这类题目的时候不要只看书上的标准答案要自己动手写一遍kernel用cuda-memcheck跑一下看看边界情况下的行为。笔试考的是你的工程直觉而工程直觉只能从实践中来。3. 深度学习方向的专业题CNN计算量与反向传播如果你投的是英伟达深度学习方向的实习笔试里一定会出现CNN相关的计算题。这类题目不会让你从头推导卷积公式而是会给你一个具体的网络结构让你计算参数量、FLOPs、内存占用或者分析某个操作在GPU上的并行性。3.1 卷积层的参数量和FLOPs怎么快速估算先给一个最常用的估算公式。对于一个卷积层输入特征图尺寸为H_in × W_in × C_in卷积核大小为K × K输出通道数为C_out步长为Spadding为P那么输出特征图尺寸H_out (H_in 2P - K) / S 1W_out同理参数量K × K × C_in × C_out C_out偏置FLOPs乘加各算一次2 × K × K × C_in × C_out × H_out × W_out笔试里经常考的一个对比是3x3卷积和5x5卷积的参数量差多少。答案是25/9≈2.78倍。但如果你用两个3x3卷积堆叠来代替一个5x5卷积感受野是一样的参数量却是2×918比25小而且中间多了一个非线性激活。这就是VGG网络设计的核心思想之一。笔试如果考到“为什么VGG用3x3卷积堆叠”你就得从这个角度回答。还有一个常考的点是1x1卷积的作用。很多人觉得1x1卷积没有意义因为它不改变感受野。但实际上1x1卷积可以用来做通道数的升降维从而减少计算量。比如一个3x3卷积前面加一个1x1卷积把通道数从256降到64再做3x3卷积升回256总FLOPs可能比直接做3x3卷积还少。这种“瓶颈结构”在ResNet和MobileNet里都有应用。3.2 反向传播中哪些步骤可以并行英伟达笔试里有一类题是让你分析反向传播的计算图问哪些操作可以并行、哪些必须串行。这个问题的核心在于理解数据依赖关系。以卷积层的反向传播为例它需要计算三个东西对输入feature map的梯度、对卷积核权重的梯度、对偏置的梯度。其中对输入梯度的计算是一个“转置卷积”操作对权重梯度的计算是一个“输入和输出梯度的卷积”操作。这两个计算之间没有依赖关系可以并行。但对偏置的梯度需要把输出梯度的所有空间位置加起来这是一个reduction操作需要同步。在GPU上实现的时候reduction操作通常用树形归约来做时间复杂度是O(log N)。笔试里可能会让你写出树形归约的伪代码或者问你在warp级别怎么做归约。warp级别的归约可以用__shfl_down_sync()指令来实现不需要shared memory速度更快。这个知识点在GPU计算方向的笔试里也经常出现。还有一个容易忽略的点是batch normalization的反向传播。BN的反向传播需要用到整个batch的统计量所以它天然是一个跨样本的操作。在GPU上如果batch size很大BN的反向传播会成为性能瓶颈因为它需要做全局的reduction。笔试里如果考到“怎么优化BN的反向传播”你可以从“用warp-level reduction减少同步开销”或者“用persistent kernel避免多次kernel launch”的角度回答。3.3 优化器的选择逻辑与笔试中的陷阱英伟达笔试里偶尔会考优化器相关的题目比如“SGD和Adam的区别是什么什么情况下用哪个”。这种题看起来是送分题但答不好也容易丢分。关键是要说出具体的场景和权衡而不是背定义。SGD的更新规则是θ θ - lr * g简单直接但需要仔细调学习率而且容易陷入鞍点。Adam用了一阶矩和二阶矩的指数移动平均自适应地调整每个参数的学习率收敛更快但可能泛化不如SGD。笔试里如果问“为什么CV任务里SGD动量比Adam更常用”你得从“Adam的自适应学习率可能导致某些参数更新过大破坏预训练权重的结构”这个角度回答。还有一个陷阱是weight decay和L2正则化的区别。在SGD里weight decay等价于L2正则化但在Adam里因为学习率是自适应的weight decay和L2正则化不等价。这就是AdamW被提出来的原因。笔试里如果考到“AdamW和Adam的区别”你得能说清楚这个细节。我个人的建议是准备深度学习方向的笔试时不要只盯着CNN和反向传播还要把优化器、正则化、初始化这些基础但容易忽略的点过一遍。英伟达的面试官很喜欢问“为什么”类的问题比如“为什么He初始化适合ReLU”你得能从方差保持的角度解释清楚。4. GPU计算方向warp、occupancy与内存层次GPU计算方向的笔试是英伟达所有方向里最“硬核”的因为它直接考察你对硬件执行模型的理解。这个方向的题目往往不需要你写完整的代码而是给你一段代码或者一个场景让你分析性能瓶颈在哪里、怎么优化。4.1 warp divergence为什么是性能杀手Warp是GPU调度的基本单位一个warp里的32个thread必须执行相同的指令。如果代码里有分支比如if (threadIdx.x 16) { ... } else { ... }那么warp里的thread会分成两组先执行if分支再执行else分支两组都执行完之后再继续。这就是warp divergence它会让实际执行时间变成原来的两倍。笔试里经常考的一个场景是在一个循环里做条件判断问怎么消除divergence。常见的优化手段是把分支改成算术运算。比如if (x 0) y x; else y 0;可以改成y max(x, 0);这样就没有分支了。另一个手段是重新组织数据布局让同一个warp里的thread走相同的分支。比如在图像处理里如果按行处理会导致边界上的warp divergence可以改成按列处理或者用padding来避免。还有一个更隐蔽的divergence来源是循环次数不一致。如果warp里的thread循环次数不同那么循环次数少的thread会等待循环次数多的thread造成浪费。笔试里如果考到“怎么处理不规则循环”你可以从“用warp-level的ballot指令统计活跃thread”或者“用work-stealing的方式动态分配任务”的角度回答。4.2 occupancy到底怎么算为什么不是越高越好Occupancy是指GPU上实际活跃的warp数量与最大可能活跃warp数量的比值。它受三个因素限制register数量、shared memory大小、block大小。笔试里经常给你一组参数让你算occupancy。举个例子假设一个GPU有64K个register per SM每个block有256个thread每个thread用32个registershared memory每个block用8KBSM的shared memory总量是96KB。那么register限制64K / (256 * 32) 8个blockshared memory限制96KB / 8KB 12个block硬件限制通常一个SM最多驻留16或32个block所以实际能驻留的block数是min(8, 12, 硬件限制) 8个block也就是8 * 256 2048个thread。如果SM的最大thread数是2048那occupancy就是100%。但occupancy不是越高越好。有时候降低occupancy反而能提升性能因为每个thread可以用更多的register减少register spill到local memory的情况。笔试里如果问“什么情况下低occupancy反而更快”你可以从“计算密集型任务需要更多register来缓存中间结果”这个角度回答。4.3 shared memory的bank conflict怎么排查Shared memory被划分成32个bank每个bank宽度4字节。如果warp里的两个thread访问同一个bank的不同地址就会发生bank conflict访问被串行化。笔试里经常给出一段访问shared memory的代码让你判断有没有bank conflict。比如shared[threadIdx.x]这种访问模式thread 0访问bank 0thread 1访问bank 1...thread 31访问bank 31没有冲突。但如果是shared[threadIdx.x * 2]thread 0访问bank 0thread 1访问bank 2...thread 16访问bank 32也就是bank 0这就跟thread 0冲突了。这种访问模式会导致2-way bank conflict。避免bank conflict的常见手段是padding。比如把shared memory数组的大小从32改成33这样shared[threadIdx.x * 2]的访问模式就变成了thread 0访问bank 0thread 1访问bank 2...thread 16访问bank 33也就是bank 1不再跟thread 0冲突。这个技巧在矩阵转置的kernel里特别常用。笔试里如果考到“怎么优化矩阵转置的shared memory访问”你就得写出padding的代码并解释为什么padding能消除bank conflict。这个知识点看起来小但能体现你对硬件细节的掌握程度。5. 笔试之外的准备从刷题到动手实践聊完具体的知识点再说说准备策略。我见过太多同学把英伟达笔试当成国内互联网大厂的笔试来准备疯狂刷LeetCode结果笔试一出来发现题目完全不是那个路子。英伟达的笔试更看重你对并行计算模型和深度学习底层机制的理解而这些理解只能从动手实践中来。5.1 用CUDA写几个kernel比刷一百道题有用如果你投的是GPU计算方向我强烈建议你在笔试前至少亲手写这几个kernel向量加法、矩阵乘法naive和tiling两个版本、归约求和、矩阵转置。写完之后用nvprof或者nsight compute跑一下看看occupancy、内存吞吐、计算吞吐这些指标。你会发现很多书上讲的“优化技巧”在实际跑的时候效果不一样。比如书上说tiling能大幅提升矩阵乘法的性能但你实际跑的时候可能会发现当矩阵尺寸比较小的时候tiling反而更慢因为shared memory的加载和同步有开销。这种“理论跟实践不一致”的地方恰恰是笔试里容易考的。面试官想看到的是你有自己的判断而不是背结论。如果你手头没有GPU可以用在线的GPU平台跑或者用nvcc的模拟器模式。虽然模拟器不能反映真实的性能但至少能验证代码的正确性。我个人的经验是写kernel的时候一定要用cuda-memcheck检查内存错误很多bug在模拟器里看不出来但在真实GPU上会暴露。5.2 深度学习方向要把PyTorch的底层实现过一遍如果你投的是深度学习方向光会用PyTorch搭网络是不够的你得知道PyTorch底层是怎么实现的。比如nn.Conv2d在GPU上是怎么调用的autograd是怎么构建计算图的DataLoader的多进程是怎么工作的。这些知识在笔试里可能不会直接考但在面试里一定会问。我建议你找一个简单的网络比如LeNet或者ResNet的一个block用PyTorch的profiler跑一下看看每个操作的时间占比。你会发现卷积操作通常占大头但batch normalization和激活函数也不容忽视。如果你能说出“在这个网络里卷积占了70%的时间BN占了15%”面试官会觉得你真的动手分析过。还有一个容易被忽略的点是混合精度训练。英伟达的GPU对FP16和TF32有硬件加速笔试里可能会考“混合精度训练的原理是什么为什么能加速”。你得能说清楚FP16的表示范围、loss scaling的作用、以及哪些操作必须用FP32。这个知识点在英伟达的笔试里出现频率很高因为它是英伟达GPU的核心卖点之一。5.3 时间分配和答题策略英伟达的笔试时间通常比较紧题目数量不多但每道题的分值很高。我的建议是先做你最有把握的题把能拿的分都拿到再回头啃难题。因为笔试的评分往往是按点给分的一道难题你可能只能拿到部分分但一道简单题你拿满分总分反而更高。还有一点写清楚你的推导过程。英伟达的笔试不像LeetCode那样只看最终答案它更看重你的分析思路。即使最后答案错了如果你的推导过程合理也可能拿到大部分分。所以答题的时候不要只写一个数字要把你的假设、公式、计算步骤都写出来。我见过一个同学笔试的时候遇到一道矩阵乘法的性能分析题他最后答案算错了但他把内存访问量、计算量、计算强度的推导过程都写清楚了最后还是过了。面试官后来跟他说他们看中的是分析框架不是最终数字。6. 几个容易踩的坑和我的个人体会最后说几个准备英伟达笔试时容易踩的坑这些都是我从自己和身边人的经历里总结出来的。第一个坑忽视C基础。英伟达的笔试里有一部分是C相关的题目比如指针、引用、虚函数、模板。这些题目不难但如果你长时间不写C很容易忘。我建议笔试前把C Primer的指针和内存管理章节过一遍特别是new/delete、malloc/free的区别以及RAII的概念。这些在GPU编程里也很重要因为你要手动管理device memory。第二个坑不熟悉Linux环境。英伟达的笔试通常是在Linux环境下进行的你需要会用gcc、g、make、gdb这些工具。如果你平时只用Windows或者Mac建议提前装一个Linux虚拟机或者用WSL熟悉一下。笔试里可能会让你写一个Makefile或者用gdb调试一段代码这些技能不是临时能补上的。第三个坑对GPU架构的演进不了解。英伟达的GPU架构从Fermi到Hopper每一代都有变化。比如Volta引入了独立的Tensor CoreAmpere引入了TF32和稀疏化支持Hopper引入了DPX指令和Thread Block Cluster。笔试里可能会考“某一代架构的新特性是什么”如果你只了解老架构就会答不上来。我建议你把最近三代架构的whitepaper翻一翻至少知道每代的核心卖点。第四个坑面试和笔试脱节。很多人笔试过了之后面试就放松了结果面试被问得更深。英伟达的面试通常会让你现场写代码或者白板推导一个算法。如果你笔试是靠刷题过的面试很容易露馅。所以准备笔试的时候就要用面试的标准要求自己——不仅要会做题还要能讲清楚为什么这么做。我个人的体会是准备英伟达笔试的过程其实是一个重新理解计算机体系结构的过程。你会被迫去思考数据在内存和计算单元之间怎么流动线程怎么调度延迟怎么隐藏。这些思考不仅对笔试有用对你以后写任何高性能代码都有用。所以不要把它当成一个应试任务把它当成一次系统学习的机会。还有一个实用的小技巧笔试前一周把英伟达官方的CUDA C Programming Guide的目录过一遍看看哪些章节你还不熟悉。这个文档写得非常好很多笔试题目都能在里面找到影子。特别是“Performance Guidelines”那一章几乎涵盖了所有性能相关的考点。最后再分享一个我自己的习惯每次做完一道题不管对错我都会问自己“如果我是出题人我会怎么改这道题来考更深的知识点”。这个习惯让我在笔试里遇到变体题的时候不会慌因为我已经提前想过各种可能的变体了。英伟达的笔试题目往往就是几个核心知识点的排列组合你把核心知识点吃透了怎么变都不怕。 英伟达暑期实习笔试全解析：并行计算与深度学习底层机制
1. 英伟达暑期实习笔试到底在考什么聊英伟达暑期实习笔试之前先得把一件事说清楚这家公司的笔试跟国内大多数互联网大厂的套路完全不一样。国内很多公司的笔试题偏向八股文和LeetCode中等难度偏上的算法题刷够题量基本能应付。但英伟达的笔试尤其是涉及深度学习、GPU计算方向的岗位题目设计明显更偏向“底层理解工程直觉”的组合拳。我带过几届准备英伟达实习的学生也跟几位已经进去的朋友聊过大家的共识是笔试的区分度不在你刷了多少题而在你对并行计算模型和深度学习底层机制的理解深度。先说说笔试的基本盘。英伟达暑期实习的笔试通常分两部分一部分是通用编程能力另一部分是方向相关的专业题。通用编程部分不会出特别刁钻的算法但会考察你对C内存模型、指针操作、多线程基础的理解——这跟GPU编程的底层逻辑是一脉相承的。专业题部分则根据你投递的岗位方向有较大差异比如深度学习方向会考CNN的计算量分析、反向传播的链式法则推导、优化器的选择逻辑GPU计算方向会考线程层次结构、内存层次、warp调度机制系统方向可能涉及驱动模型和中断处理。为什么英伟达要这么设计因为GPU编程本质上是一个“资源受限下的并行调度”问题。你在写CUDA kernel的时候脑子里必须同时装着一件事这个计算任务怎么拆成几千个线程每个线程怎么访问内存线程之间怎么同步。这种思维方式跟传统的串行编程完全不同所以笔试题目会刻意考察你在这方面的直觉。我见过很多同学准备英伟达笔试的方式是疯狂刷LeetCode结果笔试一出来就懵了——因为题目根本不是那个路子。比如有一道经典题是给出一段CUDA代码问shared memory的bank conflict发生在哪里怎么优化。这种题你刷一万道LeetCode也遇不到但如果你理解shared memory的bank划分机制看一眼就能定位问题。所以准备英伟达笔试的第一步是把复习重心从算法题转移到并行计算基础和深度学习原理上。具体来说你需要搞清楚这几件事GPU的线程层次结构thread、block、grid是怎么组织的warp是什么概念shared memory和global memory的访问延迟差多少CNN的卷积操作在GPU上怎么并行化反向传播中哪些步骤可以并行哪些必须串行。这些内容在后面我会逐一展开。还有一点值得注意英伟达的笔试题目往往会给出一个具体的工程场景让你分析性能瓶颈或者选择方案。比如“给定一个矩阵乘法任务在什么样的矩阵尺寸下使用tiling技术能带来最大收益”——这种题没有标准答案考的是你的分析框架和权衡能力。你得能从内存带宽、计算强度、occupancy这几个维度去拆解问题而不是背一个结论。2. 从一道矩阵乘法题看GPU线程层次的设计逻辑矩阵乘法是GPU计算里最经典的入门案例也是英伟达笔试出现频率最高的题型之一。但笔试不会只让你写一个naive的矩阵乘法kernel而是会围绕它做各种变体让你分析性能、让你优化、让你解释为什么某种写法更快。要答好这类题你得先把GPU的线程层次结构吃透。2.1 thread、block、grid三层结构到底怎么理解CUDA的线程组织分三层thread是最小的执行单元block是一组thread的集合grid是一组block的集合。这个结构看起来简单但很多人在笔试里栽跟头是因为没有理解这三层结构跟硬件执行单元的映射关系。打个比方grid就像是一个大工程项目block是项目里的各个施工队thread是施工队里的工人。每个施工队有自己的临时仓库shared memory工人之间可以快速交换材料但不同施工队之间要交换材料就得走公司的大仓库global memory速度慢很多。这个类比的关键在于同一个block内的thread可以通过shared memory高效通信不同block之间只能通过global memory通信而且不能保证执行顺序。笔试里经常考的一个点是为什么block的大小通常设置为128或256个thread这跟硬件的一个关键参数有关——warp size。在英伟达的GPU架构中一个warp固定包含32个thread这是硬件调度的基本单位。如果你把block大小设成100那硬件会把它拆成4个warp3232324最后一个warp只有4个thread在干活28个thread的资源被浪费了。所以block大小最好是32的整数倍而且通常选128或256这样既能保证warp满载又不会因为block太大导致register压力过高。2.2 矩阵乘法中tiling技术的来龙去脉Naive的矩阵乘法kernel是这样的每个thread负责计算输出矩阵C的一个元素它需要读取A的一整行和B的一整列做点积。这个写法的问题在于A的每一行会被读取N次N是输出矩阵的列数B的每一列会被读取M次。global memory的访问量是O(MNK)级别的而实际计算量也是O(MNK)计算强度只有1 FLOP/byte左右远远达不到GPU的计算峰值。Tiling技术的核心思想是把矩阵分块让每个block负责计算一个tile的输出block内的thread协作把A和B的对应tile加载到shared memory里然后从shared memory里读取数据做计算。这样global memory的访问量降到了O(MNK/tile_size)级别计算强度提升了tile_size倍。笔试里常见的问法是给定矩阵尺寸和tile大小计算shared memory的使用量判断会不会超出硬件限制。比如tile大小是16x16数据类型是float4字节那么A的tile和B的tile各需要161641024字节总共2048字节。英伟达GPU的shared memory每个block通常有48KB到164KB不等取决于架构所以这个用量完全没问题。但如果tile大小是64x64那就需要64644*232KB在一些老架构上就可能超限了。还有一个容易忽略的点tiling之后每个thread的计算量增加了。Naive版本里每个thread算一个输出元素tiling版本里每个thread通常算一个tile内的多个元素。这意味着register的使用量会增加如果register不够occupancy就会下降。笔试里如果考到“为什么tile不是越大越好”你就得从shared memory容量、register压力、occupancy这三个角度去回答。2.3 笔试中关于线程层次的常见陷阱题有一类题特别阴险给出一段kernel代码问“这段代码有没有问题”。问题往往出在线程同步上。比如在tiling的矩阵乘法中加载完shared memory之后必须调用__syncthreads()否则某些thread可能还没写完shared memory其他thread就开始读了。但__syncthreads()只能同步同一个block内的thread如果你在代码里跨block做了某种假设那就错了。另一个常见陷阱是线程索引越界。当矩阵尺寸不是tile大小的整数倍时边界上的thread会访问到矩阵之外的内存。笔试里可能会让你写出边界检查的代码或者问你如果不做边界检查会发生什么。答案不是“程序崩溃”这么简单——在GPU上越界访问可能读到垃圾数据导致计算结果错误但不报错这种bug最难排查。我个人的经验是准备这类题目的时候不要只看书上的标准答案要自己动手写一遍kernel用cuda-memcheck跑一下看看边界情况下的行为。笔试考的是你的工程直觉而工程直觉只能从实践中来。3. 深度学习方向的专业题CNN计算量与反向传播如果你投的是英伟达深度学习方向的实习笔试里一定会出现CNN相关的计算题。这类题目不会让你从头推导卷积公式而是会给你一个具体的网络结构让你计算参数量、FLOPs、内存占用或者分析某个操作在GPU上的并行性。3.1 卷积层的参数量和FLOPs怎么快速估算先给一个最常用的估算公式。对于一个卷积层输入特征图尺寸为H_in × W_in × C_in卷积核大小为K × K输出通道数为C_out步长为Spadding为P那么输出特征图尺寸H_out (H_in 2P - K) / S 1W_out同理参数量K × K × C_in × C_out C_out偏置FLOPs乘加各算一次2 × K × K × C_in × C_out × H_out × W_out笔试里经常考的一个对比是3x3卷积和5x5卷积的参数量差多少。答案是25/9≈2.78倍。但如果你用两个3x3卷积堆叠来代替一个5x5卷积感受野是一样的参数量却是2×918比25小而且中间多了一个非线性激活。这就是VGG网络设计的核心思想之一。笔试如果考到“为什么VGG用3x3卷积堆叠”你就得从这个角度回答。还有一个常考的点是1x1卷积的作用。很多人觉得1x1卷积没有意义因为它不改变感受野。但实际上1x1卷积可以用来做通道数的升降维从而减少计算量。比如一个3x3卷积前面加一个1x1卷积把通道数从256降到64再做3x3卷积升回256总FLOPs可能比直接做3x3卷积还少。这种“瓶颈结构”在ResNet和MobileNet里都有应用。3.2 反向传播中哪些步骤可以并行英伟达笔试里有一类题是让你分析反向传播的计算图问哪些操作可以并行、哪些必须串行。这个问题的核心在于理解数据依赖关系。以卷积层的反向传播为例它需要计算三个东西对输入feature map的梯度、对卷积核权重的梯度、对偏置的梯度。其中对输入梯度的计算是一个“转置卷积”操作对权重梯度的计算是一个“输入和输出梯度的卷积”操作。这两个计算之间没有依赖关系可以并行。但对偏置的梯度需要把输出梯度的所有空间位置加起来这是一个reduction操作需要同步。在GPU上实现的时候reduction操作通常用树形归约来做时间复杂度是O(log N)。笔试里可能会让你写出树形归约的伪代码或者问你在warp级别怎么做归约。warp级别的归约可以用__shfl_down_sync()指令来实现不需要shared memory速度更快。这个知识点在GPU计算方向的笔试里也经常出现。还有一个容易忽略的点是batch normalization的反向传播。BN的反向传播需要用到整个batch的统计量所以它天然是一个跨样本的操作。在GPU上如果batch size很大BN的反向传播会成为性能瓶颈因为它需要做全局的reduction。笔试里如果考到“怎么优化BN的反向传播”你可以从“用warp-level reduction减少同步开销”或者“用persistent kernel避免多次kernel launch”的角度回答。3.3 优化器的选择逻辑与笔试中的陷阱英伟达笔试里偶尔会考优化器相关的题目比如“SGD和Adam的区别是什么什么情况下用哪个”。这种题看起来是送分题但答不好也容易丢分。关键是要说出具体的场景和权衡而不是背定义。SGD的更新规则是θ θ - lr * g简单直接但需要仔细调学习率而且容易陷入鞍点。Adam用了一阶矩和二阶矩的指数移动平均自适应地调整每个参数的学习率收敛更快但可能泛化不如SGD。笔试里如果问“为什么CV任务里SGD动量比Adam更常用”你得从“Adam的自适应学习率可能导致某些参数更新过大破坏预训练权重的结构”这个角度回答。还有一个陷阱是weight decay和L2正则化的区别。在SGD里weight decay等价于L2正则化但在Adam里因为学习率是自适应的weight decay和L2正则化不等价。这就是AdamW被提出来的原因。笔试里如果考到“AdamW和Adam的区别”你得能说清楚这个细节。我个人的建议是准备深度学习方向的笔试时不要只盯着CNN和反向传播还要把优化器、正则化、初始化这些基础但容易忽略的点过一遍。英伟达的面试官很喜欢问“为什么”类的问题比如“为什么He初始化适合ReLU”你得能从方差保持的角度解释清楚。4. GPU计算方向warp、occupancy与内存层次GPU计算方向的笔试是英伟达所有方向里最“硬核”的因为它直接考察你对硬件执行模型的理解。这个方向的题目往往不需要你写完整的代码而是给你一段代码或者一个场景让你分析性能瓶颈在哪里、怎么优化。4.1 warp divergence为什么是性能杀手Warp是GPU调度的基本单位一个warp里的32个thread必须执行相同的指令。如果代码里有分支比如if (threadIdx.x 16) { ... } else { ... }那么warp里的thread会分成两组先执行if分支再执行else分支两组都执行完之后再继续。这就是warp divergence它会让实际执行时间变成原来的两倍。笔试里经常考的一个场景是在一个循环里做条件判断问怎么消除divergence。常见的优化手段是把分支改成算术运算。比如if (x 0) y x; else y 0;可以改成y max(x, 0);这样就没有分支了。另一个手段是重新组织数据布局让同一个warp里的thread走相同的分支。比如在图像处理里如果按行处理会导致边界上的warp divergence可以改成按列处理或者用padding来避免。还有一个更隐蔽的divergence来源是循环次数不一致。如果warp里的thread循环次数不同那么循环次数少的thread会等待循环次数多的thread造成浪费。笔试里如果考到“怎么处理不规则循环”你可以从“用warp-level的ballot指令统计活跃thread”或者“用work-stealing的方式动态分配任务”的角度回答。4.2 occupancy到底怎么算为什么不是越高越好Occupancy是指GPU上实际活跃的warp数量与最大可能活跃warp数量的比值。它受三个因素限制register数量、shared memory大小、block大小。笔试里经常给你一组参数让你算occupancy。举个例子假设一个GPU有64K个register per SM每个block有256个thread每个thread用32个registershared memory每个block用8KBSM的shared memory总量是96KB。那么register限制64K / (256 * 32) 8个blockshared memory限制96KB / 8KB 12个block硬件限制通常一个SM最多驻留16或32个block所以实际能驻留的block数是min(8, 12, 硬件限制) 8个block也就是8 * 256 2048个thread。如果SM的最大thread数是2048那occupancy就是100%。但occupancy不是越高越好。有时候降低occupancy反而能提升性能因为每个thread可以用更多的register减少register spill到local memory的情况。笔试里如果问“什么情况下低occupancy反而更快”你可以从“计算密集型任务需要更多register来缓存中间结果”这个角度回答。4.3 shared memory的bank conflict怎么排查Shared memory被划分成32个bank每个bank宽度4字节。如果warp里的两个thread访问同一个bank的不同地址就会发生bank conflict访问被串行化。笔试里经常给出一段访问shared memory的代码让你判断有没有bank conflict。比如shared[threadIdx.x]这种访问模式thread 0访问bank 0thread 1访问bank 1...thread 31访问bank 31没有冲突。但如果是shared[threadIdx.x * 2]thread 0访问bank 0thread 1访问bank 2...thread 16访问bank 32也就是bank 0这就跟thread 0冲突了。这种访问模式会导致2-way bank conflict。避免bank conflict的常见手段是padding。比如把shared memory数组的大小从32改成33这样shared[threadIdx.x * 2]的访问模式就变成了thread 0访问bank 0thread 1访问bank 2...thread 16访问bank 33也就是bank 1不再跟thread 0冲突。这个技巧在矩阵转置的kernel里特别常用。笔试里如果考到“怎么优化矩阵转置的shared memory访问”你就得写出padding的代码并解释为什么padding能消除bank conflict。这个知识点看起来小但能体现你对硬件细节的掌握程度。5. 笔试之外的准备从刷题到动手实践聊完具体的知识点再说说准备策略。我见过太多同学把英伟达笔试当成国内互联网大厂的笔试来准备疯狂刷LeetCode结果笔试一出来发现题目完全不是那个路子。英伟达的笔试更看重你对并行计算模型和深度学习底层机制的理解而这些理解只能从动手实践中来。5.1 用CUDA写几个kernel比刷一百道题有用如果你投的是GPU计算方向我强烈建议你在笔试前至少亲手写这几个kernel向量加法、矩阵乘法naive和tiling两个版本、归约求和、矩阵转置。写完之后用nvprof或者nsight compute跑一下看看occupancy、内存吞吐、计算吞吐这些指标。你会发现很多书上讲的“优化技巧”在实际跑的时候效果不一样。比如书上说tiling能大幅提升矩阵乘法的性能但你实际跑的时候可能会发现当矩阵尺寸比较小的时候tiling反而更慢因为shared memory的加载和同步有开销。这种“理论跟实践不一致”的地方恰恰是笔试里容易考的。面试官想看到的是你有自己的判断而不是背结论。如果你手头没有GPU可以用在线的GPU平台跑或者用nvcc的模拟器模式。虽然模拟器不能反映真实的性能但至少能验证代码的正确性。我个人的经验是写kernel的时候一定要用cuda-memcheck检查内存错误很多bug在模拟器里看不出来但在真实GPU上会暴露。5.2 深度学习方向要把PyTorch的底层实现过一遍如果你投的是深度学习方向光会用PyTorch搭网络是不够的你得知道PyTorch底层是怎么实现的。比如nn.Conv2d在GPU上是怎么调用的autograd是怎么构建计算图的DataLoader的多进程是怎么工作的。这些知识在笔试里可能不会直接考但在面试里一定会问。我建议你找一个简单的网络比如LeNet或者ResNet的一个block用PyTorch的profiler跑一下看看每个操作的时间占比。你会发现卷积操作通常占大头但batch normalization和激活函数也不容忽视。如果你能说出“在这个网络里卷积占了70%的时间BN占了15%”面试官会觉得你真的动手分析过。还有一个容易被忽略的点是混合精度训练。英伟达的GPU对FP16和TF32有硬件加速笔试里可能会考“混合精度训练的原理是什么为什么能加速”。你得能说清楚FP16的表示范围、loss scaling的作用、以及哪些操作必须用FP32。这个知识点在英伟达的笔试里出现频率很高因为它是英伟达GPU的核心卖点之一。5.3 时间分配和答题策略英伟达的笔试时间通常比较紧题目数量不多但每道题的分值很高。我的建议是先做你最有把握的题把能拿的分都拿到再回头啃难题。因为笔试的评分往往是按点给分的一道难题你可能只能拿到部分分但一道简单题你拿满分总分反而更高。还有一点写清楚你的推导过程。英伟达的笔试不像LeetCode那样只看最终答案它更看重你的分析思路。即使最后答案错了如果你的推导过程合理也可能拿到大部分分。所以答题的时候不要只写一个数字要把你的假设、公式、计算步骤都写出来。我见过一个同学笔试的时候遇到一道矩阵乘法的性能分析题他最后答案算错了但他把内存访问量、计算量、计算强度的推导过程都写清楚了最后还是过了。面试官后来跟他说他们看中的是分析框架不是最终数字。6. 几个容易踩的坑和我的个人体会最后说几个准备英伟达笔试时容易踩的坑这些都是我从自己和身边人的经历里总结出来的。第一个坑忽视C基础。英伟达的笔试里有一部分是C相关的题目比如指针、引用、虚函数、模板。这些题目不难但如果你长时间不写C很容易忘。我建议笔试前把C Primer的指针和内存管理章节过一遍特别是new/delete、malloc/free的区别以及RAII的概念。这些在GPU编程里也很重要因为你要手动管理device memory。第二个坑不熟悉Linux环境。英伟达的笔试通常是在Linux环境下进行的你需要会用gcc、g、make、gdb这些工具。如果你平时只用Windows或者Mac建议提前装一个Linux虚拟机或者用WSL熟悉一下。笔试里可能会让你写一个Makefile或者用gdb调试一段代码这些技能不是临时能补上的。第三个坑对GPU架构的演进不了解。英伟达的GPU架构从Fermi到Hopper每一代都有变化。比如Volta引入了独立的Tensor CoreAmpere引入了TF32和稀疏化支持Hopper引入了DPX指令和Thread Block Cluster。笔试里可能会考“某一代架构的新特性是什么”如果你只了解老架构就会答不上来。我建议你把最近三代架构的whitepaper翻一翻至少知道每代的核心卖点。第四个坑面试和笔试脱节。很多人笔试过了之后面试就放松了结果面试被问得更深。英伟达的面试通常会让你现场写代码或者白板推导一个算法。如果你笔试是靠刷题过的面试很容易露馅。所以准备笔试的时候就要用面试的标准要求自己——不仅要会做题还要能讲清楚为什么这么做。我个人的体会是准备英伟达笔试的过程其实是一个重新理解计算机体系结构的过程。你会被迫去思考数据在内存和计算单元之间怎么流动线程怎么调度延迟怎么隐藏。这些思考不仅对笔试有用对你以后写任何高性能代码都有用。所以不要把它当成一个应试任务把它当成一次系统学习的机会。还有一个实用的小技巧笔试前一周把英伟达官方的CUDA C Programming Guide的目录过一遍看看哪些章节你还不熟悉。这个文档写得非常好很多笔试题目都能在里面找到影子。特别是“Performance Guidelines”那一章几乎涵盖了所有性能相关的考点。最后再分享一个我自己的习惯每次做完一道题不管对错我都会问自己“如果我是出题人我会怎么改这道题来考更深的知识点”。这个习惯让我在笔试里遇到变体题的时候不会慌因为我已经提前想过各种可能的变体了。英伟达的笔试题目往往就是几个核心知识点的排列组合你把核心知识点吃透了怎么变都不怕。