2019年秋天百度校招官网上出现了“AI异构计算工程师”这个岗位。当时国内能把异构计算从底层做到系统级的公司没几家百度算是最早一批把GPU集群、FPGA加速、自研AI芯片三条线都铺开的团队。这个岗位的笔试题和普通互联网公司那种刷题能过的算法题完全是两码事它直接考你懂不懂硬件、懂不懂并行、懂不懂怎么让模型在芯片上跑得快。我当年完整做过第一批题后来也参与过类似岗位的面试出题。回头看这套题的价值早就超出了“校招考试”本身它其实是AI基础架构领域最实用的一张能力地图直到今天依然值得想往高性能计算、AI基础设施方向走的同学认真过一遍。1. 笔试全景AI异构计算工程师究竟在考什么1.1 岗位画像百度AI异构计算工程师做什么很多人看到岗位名字里的“AI”和“异构计算”第一反应是“这是个算法岗”实际上完全不是。它更接近系统工程师、性能优化工程师、底层软件工程师的组合体。2019年百度AI异构计算工程师的核心工作大致分四块一是在GPU、FPGA、CPU上实现和优化深度学习算子二是做推理引擎的底层加速让训练好的模型在线上以最低延迟、最高吞吐跑起来三是训练侧的加速包括分布式通信、混合精度、显存优化等四是软硬协同设计和芯片团队一起让自研AI芯片能真正跑起来。这个岗位画像决定了笔试题的风格。它不考你怎么训模型不考SVM、LSTM的公式推导它考的是计算机体系结构里那些“看起来偏底层”的知识存储层次怎么工作、Cache命中率怎么影响性能、线程之间怎么同步、访存带宽和计算吞吐的关系。这些知识在普通后端开发岗位上可能一辈子用不到但在异构计算领域它们是每天的日常。我记得当时拿到试卷的第一感受是没有一道题是“背出来”的每道题都在逼你通过计算和分析去判断。这和百度做AI基础架构的思路一脉相承——他们需要的人不是会调包调参的工程师而是能在芯片手册、性能分析报告和框架源码之间自由穿梭的人。1.2 题型结构与时间分配虽然没有拿到官方原始试卷但根据当时参加笔试的同学反馈和各类面经整理第一批笔试题的题型结构大致可以归纳为三类我用一张表格说明考察方向和答题策略题型大致数量考察内容建议时间客观选择题20题左右体系结构、操作系统、数据结构、并行计算概念30分钟编程/代码分析题2-3题CUDA/OpenCL kernel编写、代码性能瓶颈分析60分钟系统设计/问答题2题左右软硬协同任务拆分、访存优化方案设计60分钟整场笔试时间大概两个半小时题量不算大但每道题都需要深入思考。我见过有同学在选择题上花太多时间结果最后一道系统设计题只写了三行字。这里我建议的答题顺序是先快速扫一遍所有题目把系统设计题留到大块时间集中输出客观题控制在每题1分半以内编程题如果卡住超过20分钟就先跳过去做后面的。这套笔试题的难度梯度很明显。客观题里有一部分是送分题比如“下列哪个不是GPU的存储类型”这些是给基础扎实的同学准备的真正的分水岭在编程题和系统设计题它们直接考察你有没有真正写过并行程序、有没有在性能调优时踩过坑。纸上谈兵的人和实际操作过的人写出来的答案一眼就能分辨。2. 核心考点拆解从体系结构到并行编程2.1 异构架构的底层逻辑CPU、GPU、FPGA、ASIC异构计算这个词听起来高大上本质却很朴素用不同特点的处理器去处理不同特点的任务。我把这四种处理器打一个比方CPU是一位全能厨师什么菜都会做但一次只能同时处理几道菜GPU是一千个只会切菜配菜的帮厨单个不如大厨但架不住人多适合处理大量重复劳动FPGA是一条可以按需重新搭的流水线今天做红烧肉明天做火锅都能改但搭流水线需要时间ASIC是一台只能做一道菜的专业机器做其他菜就不行但做这道菜的速度和成本无人能比。在2019年这个时间节点深度学习的训练阶段几乎被GPU垄断因为训练本质上是海量矩阵运算正好是GPU最擅长的场景推理阶段则是百花齐放CPU、GPU、FPGA、ASIC都有部署。百度的做法是训练用GPU集群搜索推荐场景的线上推理大量用FPGA做低延迟加速同时投入研发昆仑芯片ASIC路线来做更极致的能效比。这套布局要求工程师既能写GPU程序又理解FPGA的流水线思想还要能从指令集层面思考ASIC的架构取舍。笔试题里关于这部分的考察通常不会直接问你“什么是异构计算”而是给一个具体场景让你判断“用哪种硬件合适”。比如给一个自动驾驶的物体检测任务要求极低延迟、有限功耗、模型可能频繁更新这时候FPGA往往比ASIC更合适因为模型更新后FPGA可以重新配置ASIC则要重新流片。这些题目没有唯一标准答案关键是展示你的分析过程是否考虑了计算特征、数据带宽、延迟约束、部署周期这几个维度。2.2 GPU编程模型的核心考点CUDA与OpenCLGPU编程是这套笔试题的绝对重点。CUDA和OpenCL的编程模型很相似都是“主机端设备端”的模式CPU负责控制流和串行逻辑GPU负责大规模并行计算。笔试常考的核心概念包括线程层次、内存层次、同步机制三个方面。线程层次是一个三层结构grid网格、block线程块、thread线程。你可以把它理解成一家工厂整个工厂是grid车间是block工人是thread。同一车间里的工人可以高效地共享物资shared memory通过一个广播系统沟通__syncthreads不同车间的工人之间没有直接沟通渠道只能通过工厂的公共仓库global memory传递物资成本很高。这个层次模型直接决定了程序怎么写、怎么优化。内存层次是GPU编程里最容易出题的地方也是理解和性能相关的核心。一个线程有自己私有的寄存器register和局部内存local memory一个线程块共享一块shared memory所有线程都能访问global memory还有只读的constant memory和texture memory。笔试里经常让你判断“哪种内存访问最快”或者分析“某个kernel为什么慢”。答案是访问速度从快到慢大致是寄存器、shared memory、global memory。寄存器最快但数量有限shared memory次之但要处理bank conflict问题global memory最慢但容量最大。同步机制也是必考点。__syncthreads()用于同一个线程块内的线程同步常见的使用场景是先把数据从global memory搬到shared memory然后等待所有线程都搬完再开始用shared memory里的数据计算。如果漏掉这一步就会出现有些线程读到旧数据、有些线程读到新数据的问题。原子操作atomicAdd等用于跨线程的安全累加但笔试中经常强调能用并行归约就别用原子操作因为原子操作会串行化。这些知识点看似细节实际在真实GPU编程中全是高频踩坑点。2.3 性能分析的核心思想访存密集还是计算密集异构计算的终极目标是让程序尽可能跑得快而“快”的前提是知道瓶颈在哪。笔试中必考的一个方法是Roofline模型它用一条屋顶线来描述一个硬件平台的计算能力上限和访存带宽上限。任何程序的实际性能要么被计算量限制compute-bound要么被访存量限制memory-bound就像水桶的短板效应。我来演示一道典型的笔试题。假设一个GPU的峰值算力是10 TFLOPS访存带宽是200 GB/s某个kernel的总计算量是100 GFLOPS总访存量是20 GB。问这个kernel是计算密集还是访存密集瓶颈在哪优化方向是什么。计算过程是这样的算术强度Arithmetic Intensity 计算量 / 访存量 100G / 20G 5 FLOP/Byte。机器的转折点Ridge Point 峰值算力 / 带宽 10T / 200G 50 FLOP/Byte。因为5小于50说明这个kernel的算术强度低于机器的转折点它跑不满算力被带宽卡住了属于访存密集。优化方向就是减少访存量比如用shared memory做数据复用、使用向量化访存、合并小请求为大批量读写。这类题目考察的不是单纯的计算能力而是你是否理解了“性能优化的本质是让数据和计算的位置尽量靠近”。在笔试考场里能又快又准地写出这个分析过程的同学往往都有过真实性能调优经验。因为只有被带宽瓶颈折磨过的人才会对“算术强度”这个概念有肌肉记忆。3. 典型真题还原与解题思路3.1 矩阵乘法的GPU实现从朴素到Shared Memory Tiling矩阵乘法是GPU编程的“hello world”也是这套笔试题编程题的常客。题目通常会这样出给出一个矩阵乘法C A x B要求用CUDA实现并分析性能瓶颈、给出优化方案。如果只写过CPU代码或者只是听说过GPU大概率会写出一个朴素版本__global__ void matmul_naive(float *A, float *B, float *C, int N) { int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; if (row N col N) { float sum 0.0f; for (int k 0; k N; k) { sum A[row * N k] * B[k * N col]; } C[row * N col] sum; } }这个版本的性能很差原因非常典型内层循环每次都要从global memory读一个A的元素和一个B的元素一个线程计算一个输出元素需要2N次全局访存。整个kernel的访存量是O(N^3)级别而计算量也是O(N^3)算术强度大约是1 FLOP/Byte左右远远低于现代GPU的转折点性能被访存带宽死死卡住。优化的核心思路是数据复用。每个A的元素会被同一行的N个线程重复使用每个B的元素会被同一列的N个线程重复使用。如果我们把一个线程块负责计算一个TILE×TILE的输出子块就可以先把对应的A子块和B子块加载到shared memory里让块内所有线程反复使用这些数据。这样访存量从每次内层循环都读全局内存降为每个TILE块只读一次全局内存。这个技术叫Shared Memory Tiling是GPU矩阵优化中最基础也最有效的手段。优化后的kernel结构大致是每个线程块负责M_TILE×N_TILE的输出块先把A的一个M_TILE×K_TILE分块和B的一个K_TILE×N_TILE分块搬进shared memory用__syncthreads()同步然后块内线程计算局部累加循环直到处理完整个K维度。再进一步还可以用float4向量化访存来提高带宽利用率用double buffer双缓冲让数据搬运和计算重叠。笔试时即使不要求写出完整代码也要把这个优化思路写清楚这才是核心得分点。3.2 系统设计题Embedding层如何做异构加速2019年百度正值搜索和信息流业务高速发展期推荐系统是重头戏Embedding层的大规模稀疏特征处理自然成了系统设计题的常见素材。题目会这样出一个推荐系统的CTR模型Embedding表有数十亿参数存储在CPU内存中线上推理时CPU成为瓶颈请你设计一个异构加速方案。这题关键是不能一上来就说“把Embedding全放GPU显存”而是要先分析问题的特征。Embedding查询的特点是访存密集、随机性强、单个查询只读取一行向量、但并发查询量巨大。GPU虽然算力强但随机小数据量访问恰恰是最不擅长的地方因为GPU的访存优势在于合并访问coalesced access多个线程访问连续地址才能跑满带宽。把数十亿参数的Embedding表全部塞进显存也不现实2019年主流的GPU显存也就16GB到32GB而Embedding表动辄上百GB。更合理的方案是分层存储加批量预取。把最热门的Embedding高频特征对应的行缓存在GPU显存或L2 Cache中寒门特征留在CPU内存中通过批量异步预取pipeline把下一批查询需要的Embedding提前搬到GPU端。这个思路的核心是“理解数据的访问分布”用热度分析决定数据放哪一层。另一个思路是CPU-FPGA协同FPGA做Embedding查询的专用流水线把随机访存的最坏情况变成流水线的固定延迟同时用片上BRAM缓存最热门的特征。两种方案都能说关键是把访存特征分析展示出来再给出对应的工程取舍。我当年见过一个高分回答它的亮点在于算了一笔账假设查询QPS是10万每查询平均查20个Embedding每次Embedding向量是128维float512字节那么网络/PCIe需要承载的带宽是10万×20×512字节1.024 GB/s。这个带宽虽然很大但还在PCIe 3.0 x16约16GB/s的能力范围内所以“数据要不要全部搬到GPU”这个问题的答案不是简单的行不行而是在具体业务指标下能不能承受这个带宽和延迟。这种量级意识正是出题人想看到的。3.3 选择题与判断题并行编程的几个高频易错点客观题里最坑人的几类我单独拿来说。第一类是shared memory的bank conflict问题。一个线程块内的线程访问shared memory时会被分成32个bank如果多个线程同时访问同一个bank的不同地址就会发生冲突导致访问被串行化。笔试里经常给一个访问模式让你判断有没有bank conflict比如thread t访问address[t]一般没有冲突连续访问但thread t访问address[t*2]就有冲突偶数线程打到同一个bank的相邻位置实际上访问同一bank的地址会冲突。第二类是判断一个kernel为什么慢。常见的错误答案包括“因为线程数不够”而实际上更多的瓶颈是分支发散divergence。GPU以warp32个线程为调度单位如果同一个warp里的线程走了不同的if分支这些分支会被串行执行性能直接打折。笔试里考这个点是想看你对warp执行模型有没有概念。要注意不是所有if都导致发散关键看同一warp内的线程是否彼此间分支不一致。第三类是volatile关键字的作用。在CUDA里volatile告诉编译器这个变量可能在外部被修改禁止对它做缓存优化。典型场景是CPU和GPU之间有共享的通信标志位。笔试考察点是不加volatile编译器可能把变量优化到寄存器里导致CPU端修改的值GPU始终看不到。这个知识点很细但有真实调试经验的工程师都没少吃它的亏。第四类是__syncthreads()放在if分支里的问题。假如一个线程块里的某些线程进入if另一些不进入而__syncthreads()在if里面那么不在if里的线程会跳过同步点在if里的线程会卡住等待程序直接死锁。正确做法是保证所有线程都能执行到同一个同步点。这种题看起来简单但每年都有不少人在考场上栽跟头因为没写过死锁的代码就不容易记住这个坑。4. 常见失分点与备考经验4.1 最容易丢分的四个地方结合我和几届参加过百度笔试的同学交流发现丢分点高度集中在四个方面。第一个是只写代码不做性能分析。编程题不是让你把功能实现了就行题目明确问“瓶颈在哪”“怎么优化”但很多人上来就甩一段代码没有任何推导。正确做法是先给出baseline再分析访存量、算术强度、占用率occupancy最后给出优化方案和预期收益。哪怕优化方案不够完善只要分析过程在就能拿到大部分分数。第二个是忽略数据的局部性。很多同学写GPU代码时默认GPU的global memory和CPU内存一样快这是一个致命的误解。笔试中的矩阵乘法题如果不提用shared memory做数据复用基本就告别高分了。要理解数据局部性最好的方式是实际去跑一次性能剖析工具看看访存相关指标占了多少体会一次“代码看起来没问题但性能上不去”的绝望感。第三个是同步语义不清。什么时候需要__syncthreads()什么时候需要原子操作什么时候用锁很多人分不清楚。笔试中经常给一段有坑的并行代码问你“哪里会导致错误”。这类题的正解往往就一句话线程间存在数据依赖时没有同步。但这句话背后需要你理解并行程序里“先读后写”“先写后读”的问题没有实际写过多线程/多线程块程序很难答得准。第四个是没接触过任何性能分析工具。2019年NVIDIA主推的性能分析工具是nvprof后面变成了NVIDIA Nsight Computencu。笔试题里可能会给一段性能分析输出让你判断瓶颈是访存、计算还是延迟。如果从来没看过真实的profiler报告这道题基本靠蒙。体验过一遍真实性能分析报告你才能把GPU kernel里的各种性能指标和硬件行为真正联系起来。4.2 备考路线从零基础到能力覆盖如果你现在想冲百度这类AI异构计算岗位或者只是想在AI基础架构方向打基础我建议按这个路线走。第一层是计算机体系结构基础。推荐先读《深入理解计算机系统》CSAPP的存储器层次、流水线、并行章节再读《计算机体系结构量化研究方法》里的GPU部分。读的时候不要去背概念而是想一想“如果我要在GPU上把一段代码跑快哪些硬件知识能帮上忙”。这一层是地基花一个月时间比较合适。第二层是CUDA编程实战。教材方面《CUDA C编程权威指南》和NVIDIA官方编程指南都值得认真过但更重要的是动手。建议自己实现这几个经典例子向量加法、矩阵乘法从朴素到shared memory tiling再到向量化、归约reduce、卷积至少理解im2col思路。每个例子都要配合性能分析工具去看指标理解为什么优化前后差异那么大。这个阶段可以穿插看OpenCL因为它和CUDA概念几乎一一对应学起来会很快。第三层是看深度学习框架的底层源码。TensorFlow、PyTorch、PaddlePaddle的算子实现里都有大量GPU kernel代码去读最常用的卷积、GEMM、广播、拼接的算子实现再对照优化手段你就能理解工业级代码和教学代码的差距。如果有余力可以把一个简单算子的手写GPU实现和框架自带的实现做性能对比这个对比本身就能让你学到很多。整个备考周期建议安排3到6个月每天保持两小时以上强度。笔试不是临时抱佛脚能过的它需要的是持续积累的“手感”。5. 从笔试到实际工作这套题到底在训练什么能力5.1 笔试题和日常工作的对照关系很多人以为校招笔试是“应试教育”考完就忘。但AI异构计算这套笔试每一个考点在工作里都有对应。矩阵乘法优化对应的是全连接层、Attention里的QK^T、卷积的底层GEMM实现你在笔试里写的shared memory tiling工作后就是深度学习框架算子库里的标准套路。Embedding系统设计题对应的是大规模稀疏训练和推理的日常。做推荐系统的人每天都要和特征表打交道如何缓存、如何分片、如何降低带宽消耗和笔试里那道题几乎一模一样。如果你的方案扎实入职后很快就能上手这类任务。同步与死锁的问题对应多卡分布式训练里的通信同步、AllReduce、参数服务器设计。单卡GPU上遇到的__syncthreads()问题在多卡环境里会放大成分布式系统的死锁和性能问题。理解本地同步是理解全局同步的基础。性能分析工具的使用更是日常工作的一部分。给一个Kernel做优化第一步永远是profiling用数据说话而不是用直觉。这套笔试的本质是在筛选一种思维模式先量化分析再动手优化最后验证收益。5.2 我的备考与工作体会以我个人经验来说这套笔试题最值钱的地方在于它帮你把AI异构计算的学习路线划出了重点。很多人想入行AI基础设施却被海量的技术名词淹没不知道先学什么。这套题等于给你画了一张地图体系结构是起点并行编程是手段性能分析是标尺框架源码是落地场景。我后来在面试候选人的时候特别喜欢问一个和笔试有关的问题你做过的最耗时的一次性能优化是什么怎么定位瓶颈最后怎么解决的。回答得好的人往往都能讲出一段折腾的经历而不是只报一个“我用了共享内存”的三字答案。这种折腾的过程比任何知识点都重要。如果你正在准备类似的岗位我给你的建议是不要只刷题。去借一张GPU卡把矩阵乘法一步步优化到访存带宽的90%以上去跑一次卷积的profiling报告去调试一次死锁和bank conflict。这些实打实的经历会在笔试考场上、面试交流中、甚至入职后的头三个月成为你最有底气的资本。