2018年那批秋招,我第一次在招聘页面上看到“AI异构计算工程师”这个岗位,第一反应是:这岗位到底考什么?是考深度学习模型,还是考数据结构?后来真把笔试题过了一遍才发现,这套题比想象中实在得多——它几乎不跟你玩什么花活,核心就是在考三件事:体系结构、GPU并行编程、还有深度学习框架的底层执行逻辑。说白了,这就是一个在AI应用和硬件之间搭桥的岗位,笔试要筛的,就是那些真正理解“异构”二字的人。
这篇文章我会围绕百度2018校招AI异构计算工程师笔试题(第一批)做一次完整的复盘拆解,结合当时的岗位定位,把背后涉及的异构计算核心原理、常考题型、答题思路和现场策略都展开讲清楚。不管你是准备AI基础设施岗、GPU开发岗,还是单纯想补一补异构计算的基本功,这篇内容应该都能给你一个比较完整的参考框架。
1. 从岗位JD反向拆解:异构计算笔试到底想筛什么人
1.1 异构计算工程师实际做什么
先说一个很多人忽略的事实:笔试内容从来不是凭空出题,它一定是从岗位职责里倒推出来的。2018年的“AI异构计算工程师”在百度做什么?核心就是围绕PaddlePaddle训练框架和推理引擎,做GPU算子优化、多卡训练加速、框架底层运行时优化,以及模型部署时对各种异构硬件的适配。
这里面包含几个关键能力:第一,看得懂体系结构,知道数据在CPU、GPU、CPU/GPU之间是怎么流动的;第二,写得出高性能算子,也就是得深入CUDA编程,而不是只会调库;第三,理解深度学习框架的执行机制,比如算子怎么调度、显存怎么管理、图优化怎么搞;第四,具备系统性能分析能力,能定位瓶颈、跑profiling、给出可验证的优化方案。
所以笔试的设计逻辑其实非常清晰:概念题考你的知识面,计算题考你的并行思维,设计题考你的工程框架。很多人把准备重点放在背深度学习概念上,结果一看到寄存器、warp、bank conflict这些词就懵了。这恰恰说明他们没读懂这个岗位的本质。
1.2 “第一批”的真实含义和题型分布
“第一批”意味着什么?意味着这套题是题库里最早放出来的一批,覆盖面通常比较完整,会兼顾基础知识和拔高题。从题型上看,大致可以分成四类:
- 选择题/填空题:考察概念是否清晰,比如线程层级、内存类型、访存带宽等
- 简答题/原理题:要求解释某种现象或机制的成因,比如GPU为什么能掩盖访存延迟、为什么GPU上的分支发散会影响性能
- 编程题/手写推导题:给出一个并行计算问题,要求写出kernel伪代码,或者推导某个算子的计算量、访存量、理论性能
- 综合设计题:描述一个实际优化场景,让你给出完整的优化思路和方案
这里我得提醒一句:笔试不是让你全做对的。按照这类岗位的筛选逻辑,出题人很清楚大多数人是做不完的。他们真正想看的是,当你面对一道不会的题时,能不能写下有效的推导思路和判断过程。哪怕一道题没做出来,但你在旁边写了“先分析访存模式,再看计算量,判断是否访存受限”这样的思路,也比交白卷强得多。阅卷的人能看到你的思维习惯。
2. CPU/GPU/CUDA执行模型:这套题的高频考点与典型错误
2.1 CPU和GPU的本质区别,怎么答才算到位
异构计算笔试里,CPU与GPU的区别几乎是必考的。很多人上来就写“GPU核多,CPU核少”,这只能拿一小半分。要真正答到位,得从体系结构的目标差异入手。
CPU的设计目标是优化延迟,也就是说,让单条指令尽可能快地执行完。所以你看到CPU里有乱序执行、分支预测、超大缓存,这些都是为了减少单线程的等待时间。而GPU的设计目标是优化吞吐,也就是单位时间内能处理多少条指令或多少个数据元素。它不追求单条指令快,而是追求同时运行海量线程,当一个线程在等内存数据时,硬件立刻切换到另一个线程继续计算,用并发把延迟藏起来。
这里面有一个很关键的细节:GPU之所以能容忍高延迟访存,靠的不是缓存,而是线程切换。一个线程访问全局内存可能要几百个周期,但SM上同时驻留着上千个线程,这些线程交替执行,算力和访存都在持续运转。只有把这个逻辑讲清楚,面试官才会觉得你是真懂架构,而不是在背诵。
另外,2018年NVIDIA主推的是Volta架构,V100上已经有Tensor Core了,FP16算力很强。笔试如果问到“为什么深度学习训练适合用GPU”,除了说并行度高,还可以补一点向量化计算和混合精度对吞吐的帮助,这会让你的答案更有时代感。
2.2 CUDA线程层级、warp和分支发散
CUDA的线程模型也是必考项。grid、block、thread三个层级的关系要能画出来,Block被调度到SM上执行,Block内的线程可以通过共享内存和同步机制协作。真正的考点在warp这个层面:GPU实际执行指令的单位是warp,一个warp包含32个线程,这32个线程在硬件上执行同一条指令。
基于warp的概念,有个非常经典的考察点:分支发散。如果同一个warp里的32个线程,有一半走了if分支,另一半走了else分支,硬件会怎么执行?答案是串行执行两个分支,先执行if分支,再执行else分支,性能在极端情况下退化为原来的二分之一。笔试里只要提到这个,就应该立刻想到warp是32线程宽这个硬件事实。
还有一个高频概念是占用率。一个SM上能同时驻留多少线程是有限制的,比如Maxwell之后很多架构是每个SM最多2048个线程。很多人误以为block开得越大越好、线程数开得越多越好,其实不然。如果每个线程用了太多寄存器,或者共享内存占用太大,都会限制SM上驻留的线程数,从而降低占用率,导致延迟隐藏能力变差。这种“看似并行实则串行等待”的细节,笔试里特别喜欢拿出来当迷惑选项。
2.3 内存层级与访存模式:笔试中的硬骨头
异构计算工程师的笔试,内存部分从来不会缺席。你要能按延迟从低到高排出:寄存器、共享内存、全局内存、局部内存、常量内存/纹理内存。其中寄存器是每个线程私有的,共享内存是block内共享的,全局内存是所有线程都能访问但延迟最高的那部分,而且全局内存的带宽是天生的瓶颈。
考察访存模式时,最经典的考点是合并访问。GPU访问全局内存时,以cache line或者segment为单位,一个warp的32个线程如果访问的是连续地址,硬件就能用较少的内存事务把数据一次性取回来。像float a = data[threadIdx.x]这种就是合并访问,效率高。如果改成float a = data[threadIdx.x * 2],32个线程访问的就是跨步地址,硬件可能要把一个事务拆成多个,带宽利用率下降。笔试遇到这类题,第一步永远是判断访存是否连续、是否对齐。
共享内存则要额外注意bank conflict。GPU把共享内存分成32个bank,每个bank每周期可以服务一个线程。如果同一个warp里多个线程同时访问同一个bank的不同地址,硬件会把这次访问串行化,这就是bank conflict。笔试里最阴的题就是归约时写法不对导致严重的bank conflict——这个问题在下一节我详细说,因为它其实有专门的答题套路。
3. 归约、卷积、带宽计算:手写推导题的踩分要点
3.1 并行归约:从log复杂度到shared memory优化
并行归约几乎是GPU编程笔试的常青树。核心问题很简单:给定一个长度为N的数组,怎么用GPU快速求和?
最直接的思路是分层归约:第一轮让256个线程各读一部分数据,得到256个局部和;第二轮再把256个局部和归约成1个。这里有个细节很关键,也是笔试爱考的点:归约的复杂度是O(logN),不是O(N),因为每一轮都把数据量减半。
笔试如果让写kernel伪代码,建议写成这种形式:先用一个grid-stride循环让每个线程累积多个元素的值,再用shared memory做block内的树状归约,最后用原子操作把各block的结果加到全局变量。这里面有三个踩分点:第一,你把数据划分逻辑写清楚;第二,你意识到block内的中间结果需要放在shared memory里,避免重复访问全局内存;第三,你处理了同步问题,在shared memory数据写完和读之前,要加__syncthreads()。
另一个容易在归约题里出现的坑就是数值稳定性。GPU并行归约改变了浮点数的累加顺序,比如1e20 + 1.0 - 1e20这种场景,结果可能变成0。遇到这种题,你要能说出:并行归约结果可能和串行结果有差异,可以通过Kahan求和或者更高精度累加器来缓解。这不是考你多会用API,而是考你对浮点数本质的理解。
3.2 im2col与卷积矩阵化:这道题怎么算不丢分
2018年的深度学习笔试,几乎离不开卷积的底层实现。最经典的是im2col:把卷积操作展开成矩阵乘法,然后调用cuBLAS里的GEMM来加速。笔试常常会给你一个具体的输入尺寸,让你算im2col之后的矩阵维度。
我们来记一个通用公式。假设输入是H x W x C,卷积核是K x K x M,步长S,padding P,那么输出特征图的高OH和宽OW分别是:
OH = (H + 2P - K) / S + 1OW = (W + 2P - K) / S + 1
im2col之后,输入矩阵的维度是(K*K*C) x (OH*OW),权重矩阵的维度是M x (K*K*C),两个矩阵相乘得到M x (OH*OW),再reshape成M x OH x OW就是输出特征图。
笔试除了让算维度,还会问一个问题:im2col有什么缺点?答案是它把数据复制了一份或几份,内存开销很大。比如3x3卷积,im2col之后数据量大约变成原来的9倍,这对带宽和显存都是压力。所以后续才有了Winograd之类的替代算法,减少乘法次数,也尽量控制额外存储。你能在答案里把这个取舍逻辑讲出来,比单纯背公式更值分。
3.3 理论峰值与Roofline:给一道计算题做个标准姿势
有一类计算题长这样:给你一块GPU的理论算力是X TFLOPS,内存带宽是Y GB/s,然后给你一个算子的数据量和计算量,问这个算子的耗时下限是多少。
这种题的核心是算算术强度。算术强度 = 总操作数 / 总数据字节数。如果算术强度高于平台的平衡点,也就是算力/带宽的比值,那这个算子就是计算受限的;否则就是访存受限。访存受限的情况下,理论耗时大约等于数据量除以带宽,你费再大劲优化计算都没用,优化访存才有价值。
举个具体的例子,V100的FP32算力大约是15.7 TFLOPS,内存带宽约900 GB/s,那么平衡点大约在17.4 FLOP/Byte左右。一个elementwise操作,比如给一个1GB的数组加一个标量,算力需求是1GB*1次浮点操作 = 1G FLOP,但数据量是1GB,算术强度只有1 FLOP/Byte,远低于17.4,所以它妥妥是访存受限。笔试里只要看到elementwise、copy、norm这类操作,基本可以立刻判断:它优化到接近带宽上限就是极限了。
我建议在推导这类题时,把单位换算写清楚:TFLOPS是每秒10^12次浮点操作,GB是10^9字节。很多人计算时把GB当成2^30字节,换算出来的数字差得离谱,这种低级错误千万别犯。你可以在考卷侧边把公式列一列,哪怕最后数字算错,思路分也能保住。
4. 站在框架侧看异构执行:显存、stream与算子融合怎么考
4.1 数据在哪、计算在哪:训练迭代里的异步机制
异构计算工程师不是只写kernel,他得懂深度学习框架怎么把任务分发到异构设备上。笔试里有一种送分题:单机训练一次迭代的完整流程是怎样的?你要能答出:CPU端准备数据,把数据从主机内存拷贝到GPU显存(H2D),GPU执行forward,计算loss,反向传播,更新参数,可能需要把部分统计信息拷回主机(D2H)。
这里面隐藏着一个笔试高频考点:PCIe的传输带宽有限,2018年主流是PCIe 3.0,x16链路实际传输大约12-16 GB/s,和GPU的显存带宽相比差了一个数量级。所以H2D/D2H的拷贝越少越好,最好还能和kernel执行重叠。
基于这个背景,就会考stream的概念。同一个stream内的操作按序执行,不同stream之间可以并发。拿训练来说,可以把数据预取放到一个stream里,把计算放到另一个stream里,让数据传输和上一轮的kernel计算重叠起来,这样GPU的空闲时间就少了。笔试如果给你一段代码,问为什么GPU利用率低,最常出现的原因就是:每次迭代都同步等待数据传输完成,没有用异步stream。
4.2 算子融合:launch开销到底有多大
另一个必考方向是算子融合。2018年TensorFlow里有grappler,PaddlePaddle也在做subgraph的算子融合优化。笔试不要求你熟悉具体框架API,但要求你理解融合的目的。
最核心的理由就是kernel launch有开销。一次kernel的启动,CPU要把参数打包、提交命令、GPU端还要做上下文切换,整个过程可能耗时几微秒。如果kernel本身就是2微秒的轻量运算,那launch开销占比就非常夸张了。10个小的elementwise算子每个2微秒,加上launch开销可能总共70-100微秒;如果融合成一个kernel,总时间可能就12微秒左右,性能提升好几倍。
融合还有个附带的好处:减少中间结果的显存写入和读取。elementwise算子如果不融合,每个算子都要把中间结果写到全局内存,下一个算子再读回来,一次融合就能省掉大量访存。如果一个feature map是128MB,每省一次读写就是256MB的流量,这个收益在计算中可以直接量化。笔试遇到“为什么要做算子融合”这类题,建议从launch开销、减少中间访存、以及可能的精度优化三个层面来答。
4.3 显存估算与复用:一道很接地气的应用题
2018年前后,AI从业者对显存的第一反应都是“不够用”。笔试很可能会给一个网络结构,让你估算训练时的显存占用。这时你要能列出显存的几个大块:模型参数、梯度、优化器状态(比如Adam的动量项)、中间激活值、以及临时buffer。
其中最容易爆的是中间激活值。假设某一层输出是128个batch、256个通道、56x56分辨率,且是FP32,那张激活图的大小就是128*256*56*56*4 = 大约1.6GB。如果网络有几十层,每一层的激活值都留着用于反向传播,那显存压力是很大的。
针对显存问题,框架侧的经典手段是内存复用:不同算子产生的中间张量生命周期不同,框架可以维护一个显存池,把不再使用的显存块分配给新的张量使用。笔试题如果问“怎么减少显存占用”,可以从这几个方向答:增加内存复用、混精度训练、激活重计算(不保存所有中间结果,反传时按需重算)、优化batch size。这样答,既覆盖了系统层思路,也覆盖了算法层思路。
5. 开放题才是拉分项:瓶颈定位与方案设计的答题框架
5.1 一道典型开放题的解题路径
综合设计题往往是笔试最后的大题,也是拉分的关键。常见题目是:某个模型在GPU上训练速度慢,GPU利用率低,让你分析原因并给出优化方案。这种题没有标准答案,但阅卷人心里是有评分框架的。
我的建议是分四步走。第一步,先用profiling数据说话:打开nvprof或ncu,看每类kernel的耗时占比、H2D/D2H传输耗时、以及GPU空闲时间占比。第二步,按层归类,判断瓶颈出现在数据链路、算子本身还是通信环节。比如数据加载是CPU瓶颈,那就要看是不是用了多进程数据加载、有没有做预取和异步传输;如果是kernel太碎太小,那就要考虑算子融合;如果是kernel本身慢,那就得看是不是访存模式不理想、有没有bank conflict、有没有低占用率。第三步,给出优先级排序,先解决影响最大的那一项。第四步,说明验证方式:每做一项优化,都要重跑profiling对比收益,不能靠感觉。
这里我想重点强调一个新人容易犯的错误:一上来就说“我用TensorRT加速”或者“我换混合精度”。这些建议本身没错,但要是没有前面的定位分析,显得非常廉价。笔试的大题,本质上不是考你知不知道某个工具,而是考你会不会系统地分析一个性能问题。你把定位链路写出来,哪怕结论是保守的,分数也会比直接甩优化方案高。
5.2 给“不熟悉的问题”一个通用的分析起点
开放题还有一种考法,是给你一个你没见过的模型或算子,让你说优化思路。很多人拿到这种题就慌了,因为自己完全不了解这个模型。这时候得分的关键不是了解模型,而是掌握一套通用的分析起点。
我的分析框架通常是:先问这个算子的计算密集度高不高,是访存密集还是计算密集;再问数据规模多大、是否要频繁在CPU和GPU间搬运;接着看算子内部有没有可并行的维度、有没有依赖链、能不能把多个小算子合并。这四板斧几乎能覆盖90%的kernel。比如一个softmax算子,它既访存密集又带归约,优化重点就应该是避免多次遍历数据和归约部分的数据局部性;一个全连接层,本质就是GEMM,优化重点就是矩阵分块和共享内存的tile处理。
你完全可以提前准备一个“问题排查清单”,笔试时看到任何陌生问题,先按清单过一遍。这跟真正的工程排查一模一样——上来就乱优化的人,和先定位再动手的人,写出的答案一眼就能分辨。
5.3 当年技术栈与今天的差异
2018年那会儿,AI工程的主流栈还是TensorFlow 1.x,PyTorch才刚起步,百度的PaddlePaddle也远没有今天的生态。笔试里提到的很多细节,比如显存池设计、算子融合、stream并发,今天依然是PyTorch/CUDA优化里的核心话题,只是换了个壳。
现在的学生习惯直接写PyTorch的Module,很少自己写CUDA kernel,也不怎么看底层allocator的实现。但一个异构计算工程师需要的基本功,其实和2018年一样:内存模型、并发模型、线程同步、数据搬运、访存优化。所以如果你准备的是这一类岗位的面试,与其追新工具,不如把CUDA的经典优化手段反复练透。工具会变,硬件也不会一成不变,但底层的并行计算思维不会过时。
6. 笔试现场怎么排兵布阵:时间分配与考后复盘
6.1 拿到卷子先做信息分拣
笔试时间通常很紧张,尤其是这种覆盖了体系结构、CUDA编程、深度学习框架多个方向的综合卷。我的习惯是拿到卷子后的前5分钟不答题,先把所有题目扫一遍,做三件事:判断题型、标记难度、估算分值。
做好分拣之后,按照“会做的先拿分、半会不会的写思路、完全不会的留到最后”的顺序去推进。比如一道选择题问“哪种内存延迟最低”,这种送分题立刻做掉;一道im2col的计算推导,如果公式记得,那也算稳定拿分项;但如果一道开放题让你设计一套多卡训练方案,而你连基本的AllReduce都不太清楚,那就别死磕,先把确定的题做完,最后剩多少时间写多少思路。
这里有个细节:很多笔试题的填空/简答,留给空白的位置很大,并不代表你要写满,而是阅卷时想看到你的思考过程。即使题目不太会,把你理解的条件、公式、可能的方向写在旁边,也是一个专业习惯。
6.2 推导过程比标准答案更值钱
一道计算题,如果你只写结果,错了就是零分;但如果你把假设条件、公式、代入过程、逻辑链条一步步写出来,就算最后结果错了,阅卷人也能看出你的思路是合理的。
我在实际看卷的时候,最怕看到的是“算了一大堆,但没有标注每一步在干什么”。好的推导应该是可以当科普读的:先在开头写出已知条件,再列公式,再代入数字,最后给出结论。如果是访存受限还是计算受限的判断,还要把判断依据摆出来。这样做的好处是,即使中间有一次换算失误,前后逻辑还摆在那里,阅卷人也能给你相应的步骤分。
这里特别提醒:涉及公式推导时,仔细检查单位。字节和位容易混淆,GB和GiB容易混淆,TFLOPS里的T是10^12不是2^40。笔试讲求的是思路完整和答案靠谱,如果因为单位换算翻车,真的非常可惜。
6.3 考后复盘,才是这套题最大的价值
笔试结束不等于这件事就完了。我见过太多人考完之后只关心结果,完全不管错了什么。实际上,一套覆盖完整的笔试题就是你技术能力的体检报告,它暴露的每一个盲区,都是接下来一个月可以集中补的方向。
复盘时建议把错题按维度归类:如果错在GPU架构概念,就去读NVIDIA的架构白皮书,把线程调度、内存层级、合并访问画成图;如果错在CUDA编程题,就亲手去写一个归约或GEMM,跑一遍profiling,看和自己预想的一致不一致;如果错在框架底层问题,就去深挖一个算子从Python调用到GPU执行之间的完整链路。
至于训练路径,我个人的建议是找一本经典的并行编程教材,比如Kirk和Hwu那本《Programming Massively Parallel Processors》,配合NVIDIA官方的CUDA编程指南,一道题一道题地啃。等你有能力把一个简单的向量加法优化到接近显存带宽上限,再回头来看这套笔试题,你会明显感觉到那些曾经让你卡壳的考点,都变成了一个个具体的、可触摸的工程问题。
这些年我自己再看这类笔试卷,最大的感受是:题目本身其实不神秘,它更像一份压缩过的岗位说明书,把“异构计算工程师”日常需要的能力边界清清楚楚地画了个圈。笔试未必能100%反映一个人的实战水平,但它确实能筛选出那些具备系统化思维、遇到问题知道怎么解构的人。如果你正打算投这类岗位,最后给你一个实在的建议:别光刷题,找一个GPU环境,把归约、GEMM、卷积优化这几个经典算子真正写一遍、调一遍、量一遍性能数据。等你亲手动过一轮,再回到卷面上,每一个考点都会变得具体得多。