ARTICLE DETAIL

资讯详情

深耕网站建设与运营推广的一线实战洞察。

2018百度AI异构计算工程师笔试题复盘:从体系结构到CUDA优化

2018百度AI异构计算工程师笔试题复盘:从体系结构到CUDA优化 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步长Spadding P那么输出特征图的高OH和宽OW分别是OH (H 2P - K) / S 1OW (W 2P - K) / S 1im2col之后输入矩阵的维度是(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显存H2DGPU执行forward计算loss反向传播更新参数可能需要把部分统计信息拷回主机D2H。这里面隐藏着一个笔试高频考点PCIe的传输带宽有限2018年主流是PCIe 3.0x16链路实际传输大约12-16 GB/s和GPU的显存带宽相比差了一个数量级。所以H2D/D2H的拷贝越少越好最好还能和kernel执行重叠。基于这个背景就会考stream的概念。同一个stream内的操作按序执行不同stream之间可以并发。拿训练来说可以把数据预取放到一个stream里把计算放到另一个stream里让数据传输和上一轮的kernel计算重叠起来这样GPU的空闲时间就少了。笔试如果给你一段代码问为什么GPU利用率低最常出现的原因就是每次迭代都同步等待数据传输完成没有用异步stream。4.2 算子融合launch开销到底有多大另一个必考方向是算子融合。2018年TensorFlow里有grapplerPaddlePaddle也在做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.xPyTorch才刚起步百度的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、卷积优化这几个经典算子真正写一遍、调一遍、量一遍性能数据。等你亲手动过一轮再回到卷面上每一个考点都会变得具体得多。
返回列表