FEATURED · 精选文章

百度AI异构计算笔试复盘:从CUDA到系统设计的核心考点

发布时间 / 2026/8/31 20:09:38
来源 / 创域科博编辑部
栏目 / 资讯中心
百度AI异构计算笔试复盘:从CUDA到系统设计的核心考点 2018年那会儿AI异构计算这个岗位还不像现在这么泛滥百度在校招提前批里单独划出这么个方向其实很有信号意义。我当时手忙脚乱地投了简历拿到笔试题一看发现它考的东西跟普通后端完全不是一回事大量体系结构、CUDA、并行计算、深度学习框架底层设计的内容甚至还有手写算力估算的题。那场笔试让我第一次意识到AI异构计算工程师不是会调库就行而是要在CPU-GPU-FPGA这些不同计算单元之间找到最优解的人。这篇文章我就从自己的答题经历出发把当年那套笔试题背后真正想考察的知识点拆开说一说也顺便聊聊如果现在准备这类岗位应该往哪些方向使劲。1. 当年这道提前批笔试到底在筛什么人1.1 岗位定位与笔试风格提前批本身就有抢人的性质笔试题目会比正式批更偏筛选尖子。百度招AI异构计算工程师目标很明确不是招一个会写Python调框架的算法工程师也不是招一个只懂驱动和内核的系统工程师而是招能同时理解算法、框架、硬件三层逻辑的人。所以笔试题目没有一道是纯记忆题几乎全是给你一个场景你来算、来设计、来排查的题型。我印象里整张卷子大概包括选择题、简答题和一道大型综合设计题。选择题覆盖计算机体系结构、OS并发、CUDA编程模型、深度学习前向推理的数学基础简答题集中在GPU内存层次、kernel launch开销、共享内存使用等最后的大题则给了一个类似在异构设备上部署一个CNN模型的题目要求做方案设计和性能估算。这种结构现在看仍然很有代表性先确认你的知识面再确认你在真实工程约束下能不能做决策。1.2 考察范围全景图从笔试内容可以反推岗位能力模型我总结成四个层次硬件层CPU、GPU、FPGA、ASIC各自适合什么负载带宽、延迟、功耗、算力怎么算。编程层CUDA/OpenCL的线程模型、内存模型、同步原语、性能优化手段。框架层TensorFlow/PaddlePaddle等框架如何把计算图切分到不同设备算子如何注册和调度。算法层CNN/RNN等模型的计算量、参数量、中间结果shape变化怎么映射到硬件资源。这四个层次不是割裂的题目经常交叉考察。比如问一个卷积层在GPU上为什么比CPU快这种问题既涉及算法计算模式又涉及硬件并行能力还涉及数据排布方式。如果只背结论不思考底层原理很容易在变化问法上翻车。所以准备笔试时别急着刷LeetCode先把这个岗位的知识坐标系建立起来。2. 异构体系结构从懂概念到会计算2.1 CPUGPUFPGA共存时的关键指标异构计算的核心不是堆硬件而是让每个计算单元做自己最擅长的事。笔试里经常出现的一个选择题是CPU、GPU、FPGA、ASIC四种平台各自的特点是什么这种题表面是常识实际上很多人的理解是错位的。我当时总结了这样一张对照表笔试和面试都用得上平台优势劣势适合场景CPU控制逻辑强单线程响应快生态成熟并行单元少算力密度低逻辑分支复杂、任务调度、IO密集型GPU数千个核心并行吞吐高显存带宽大分支发散敏感单核频率低不适合延迟敏感矩阵运算、卷积、Transformer等规整大计算FPGA可重构流水线延迟极低功耗可控开发周期长时钟频率低需要硬件描述预处理、协议解析、固定流水线推理ASIC性能功耗比最优一次性成本高迭代慢大规模量产、固定算法场景光记表格不够笔试会算。我遇到一道题大概是给了GPU的SM数量、核心频率、每周期浮点操作数让算峰值算力。这类题的核心公式是峰值算力 核心数 × 频率 × 每周期浮点运算次数。但很多人会漏掉每周期浮点运算次数这个系数以为频率×核心数就是全部。实际上modern GPU一个线程束warp里的32个线程如果执行FMA乘加指令一次会算两个浮点操作乘和加所以总操作数要再乘2。严谨地算应该拆成SP数量 × 2 × 频率。这个细节如果没理解后面涉及算力对比的题目都会错。2.2 访存延迟与带宽的数学感觉异构计算里数据搬移往往比计算更贵。笔试常问的一个概念是算术强度arithmetic intensity定义是每个字节的数据搬移对应多少次浮点操作公式为AI 运算量(FLOPs) / 数据搬移量(Bytes)。判断一个kernel是计算密集型还是访存密集型就是拿这个值和硬件平台的机器平衡值machine balance对比。机器平衡值约等于峰值算力 / 内存带宽比如某GPU峰值算力是10 TFLOPS带宽是800 GB/s那机器平衡值就是12.5 FLOPs/Byte。如果kernel的算术强度小于这个值说明受带宽限制此时优化重点是减少访存如果大于这个值说明受计算限制此时优化重点是提高计算效率和ILP指令级并行。我当年笔试时遇到一道简答题为什么矩阵乘法在GPU上能实现很高的利用率而ElementWise操作很难这就是用算术强度解释的典型例子。矩阵乘法中每个输出元素需要读入两个K维向量计算量为 2K 次操作但数据量是 2K×4 字节假设浮点算术强度约K/4K1024时就是256 FLOPs/Byte远大于机器平衡值属于计算密集。ElementWise每个输出只读一个输入、写一个输出算术强度是0.25 FLOPs/Byte左右明显访存密集带宽决定上限。明白这个就知道为什么GPU做向量加法快不了多少了。我在实际答题时习惯先写公式再带数字因为阅卷人更看重思路。就算最终数字算错公式和变量定义清晰也能拿到大部分分。这个写法也是后来做性能分析时养成的习惯。3. CUDA与并行编程真正的分水岭题目3.1 线程模型与映射关系CUDA相关题目基本每年必考而且问法很多变。最基础的是线程层次grid、block、thread三者的关系以及blockIdx、threadIdx跟实际数据索引的映射。笔试里常见问题一维数组长度NblockSize256要用多少个block这题很简单就是(N blockSize - 1) / blockSize但很多人会在边界处理上翻车。更进阶的问题是二维数据比如图像怎么映射到二维block这个要分清blockIdx.x和blockIdx.y分别对应数据宽还是高以及线程tid blockIdx.x * blockDim.x threadIdx.x这种公式在二维时如何扩展。我记得有一道选择题问一个warp中32个线程执行非对齐访问的后果——这是考察对内存合并(memory coalescing)的理解。GPU全局内存访问是分段的比如128 bytes段如果32个线程读的数据地址完全连续且对齐就合并成少数几次事务如果地址分散就会拆成多次事务访存效率骤降。答题时要说清楚非连续地址会导致内存事务增多带宽利用率下降最好还能举一个例子说明如何通过__ldg或共享内存做重排来改善。3.2 归约、共享内存与Bank Conflict共享内存和归约是CUDA笔试的经典大题。我记得试卷里有一道让分析warp内归约为什么要用__syncthreads()的题目实际考察的是同步语义和race condition。正确的写法是每个线程先从全局内存读一部分数据到共享内存然后循环折叠sdata[i] sdata[i stride]每次折叠后要加__syncthreads()保证所有线程看到更新后的结果。但很多人会忘记最后一次同步或者把__syncthreads()放在if条件内导致死锁。这两类错误在笔试简答题里很常见。再往深一层是Bank Conflict。共享内存被分成32个bank每个bank同周期只能服务一个地址。如果warp内多个线程访问的地址映射到了同一个bank比如步长为32的访问硬件就要串行化性能退化最多到原来的1/32。笔试常考题一个float s[32][32]数组按列访问与按行访问哪个容易bank conflict答案是按列访问。因为float类型每行占32个float刚好占满32个bank按列访问时同一列的32个元素地址相差32个bank全部落在同一个bank上产生严重冲突。解决方案是给数组加padding比如定义成float s[32][33]让每行错开一个bank列访问就分散到不同bank了。这种细节只是看书不做题是体会不到的我当时也是笔试后认真跑了一次代码才理解透。3.3 原子操作与并发正确性笔试还会问到原子操作尤其是多block并发更新同一个全局变量的场景。比如直方图统计如果用globalHist[bin]非原子操作会产生race condition用atomicAdd可以保证正确但高并发下会变成串行瓶颈。回答优化思路时会提到每个block先有私有直方图最后再做原子归约这块其实暴露的是对局部性分段归约思想的理解。写答案时最好把共享内存的归约过程写出来包括初始化和同步让阅卷人知道你确实动手写过CUDA程序。我记得当时有一道题问当一个kernel函数里既有全局内存原子操作又有共享内存原子操作哪个开销更大为什么全局内存原子操作要经过L2延迟可能几百个周期共享内存原子操作基本在一个bank内完成延迟很低。但要注意如果共享内存原子操作也发生bank conflict性能同样会劣化。这种题没有标准死答案关键看你是不是真的理解每一层硬件的cost。4. 深度学习框架与算子实现背后的问题4.1 框架计算图的异构执行逻辑AI异构计算工程师离不开深度学习框架。笔试里有一个方向是问框架如何做异构切分比如一个模型有多个op一部分放在GPU上一部分放在CPU上框架怎么管理设备间依赖当时TensorFlow的数据流图已经把设备放置device placement做得比较清楚每个节点在建立时就可以指定设备图执行器根据设备信息决定哪些节点在哪个设备上运行。如果相邻op在不同设备上就需要通过Send/Recv节点把张量拷贝到对方设备这个拷贝过程本身就是性能热点。我记得有道简答题是如何减少CPU与GPU之间的数据拷贝次数最佳回答要提到一是用pin_memory锁页内存加速host到device的拷贝二是尽量让计算在同一个设备上连续执行三是利用CUDA Stream实现计算与拷贝重叠四是对于小数据量可以用Zero-Copy或统一内存Managed Memory避免显式拷贝。虽然2018年时Zero-Copy在部分场景有争议但能写出这些方案说明你对异构执行有系统认识。如果只答不要频繁拷贝这种废话基本就拿不到分。4.2 手写卷积/矩阵乘的优化思路大笔试里出现手写卷积优化题的概率不低。我记得有一段填空是关于im2col的把一个输入特征图[N, C, H, W]和一个卷积核[M, C, Kh, Kw]的卷积操作转换成矩阵乘法[N*C*Kh*Kw, H*W] × [M, C*Kh*Kw]^T。题目会让补充各个维度的数值这就要求你清楚im2col的索引顺序。写完im2col后往往还问这样转换有什么优缺点优点是把非规整的滑动窗口卷积变成高度优化的GEMM可以利用cuBLAS的kernel缺点是会产生大量冗余数据、内存膨胀尤其当Kh*Kw较大时。2018年那会儿很多推理框架已经在用类似方式现在Winograd、FFT等方法也经常考。答题时最好能画一个数据流示意用文字描述让阅卷人看到你对卷积计算过程的宏观把握。矩阵乘优化也是常客从最简单的三层循环到ikj循环优化再到分块tile、向量化FMA、重排数据为sgemm标准结构其实就是在考你对存储层次的理解。我记得笔试里有一道问为什么说循环顺序从ijk变成ikj能提高缓存命中率解释是ikj遍历时每次计算C[i][j] A[i][k]*B[k][j]A的行和B的列在内存中都是连续的一段但j内层循环时A[i][k]是标量可以放在寄存器里对B的访问是连续的这样局部性更好。而ijk循环中内层k遍历会导致两个矩阵都跳着访问缓存命中率低。能从这个角度答题说明你有cache-line意识。4.3 算子的访存和计算强度评估还有一种题型是给你一个具体算子让你估算它的FLOPs和访存量。比如一个全连接层输入[M, K]权重[K, N]输出[M, N]。计算量是2*M*N*KFLOPs参数量是K*Nfloat。如果输入从外部读入权重驻留在显存输出的写也要算进去那么总访存量近似为(M*K K*N M*N) * 4bytes。有了这两个数字就能算算术强度判断优化方向。我当时答这类题时习惯明确写出假设条件比如假设权重常驻显存不考虑多batch复用这样就算题目没说明阅卷人也会觉得你考虑周全。5. 性能分析与系统设计的实战题型5.1 算力估算与瓶颈定位笔试最后一道综合设计题通常是给定一个AI推理任务让你做异构计算方案。我遇到的题目大概是一个实时视频流处理系统每秒要处理30帧1080p图像跑一个轻量级检测模型问怎么分配到CPU和GPU上需要什么规格的硬件。这类题看起来开放其实有固定的拆解思路先算总计算量再算单帧延迟预算再根据算力需求选硬件最后设计流水线。以一个简单模型为例假设模型对单帧1080p图像做前向推理需要15 GMACs十亿次乘加那么单帧浮点操作数约30 GFLOPs乘加各算一次。每秒30帧就是900 GFLOPs即0.9 TFLOPS。如果用一个int8量化模型算力需求会降到四分之一左右但工程上还要考虑数据预处理、后处理、内存拷贝的开销。笔试要考察的正是这种算清楚账的能力。我当时是把每一步代价都列出来再讨论在GPU上做卷积、在CPU上做预处理和NMS的划分方案尽量体现出端到端思维。5.2 数据搬移与计算重叠综合题里几乎必问如何隐藏数据搬移带来的延迟。这题的通用答案是使用多个CUDA Stream把数据拷贝H2D、kernel执行、结果拷回D2H放在不同Stream上形成流水线。如果只有一份数据流程是拷贝数据到GPU执行kernel拷回结果。假设拷贝耗时A计算耗时B总耗时AB...如果分块成N块用流水线并行则总耗时约为max(A, B) 小项而不是sum(A)sum(B)。这就是计算与通信重叠的核心逻辑。笔试中如果只写用stream是不够的最好画出类似这样的时序描述阶段0: 预处理CPU数据准备第0块同时H2D拷贝第0块为空。阶段1: H2D拷贝第0块kernel计算第0块。阶段2: H2D拷贝第1块同时kernel计算第0块D2H拷贝第0块结果。我当年答题时没有画时序图现在用mermaid那时候手画草图可以但通过文字描述了重叠关系也拿到了不错的分数。建议大家平时自己实现一个vectorAdd的pipeline代码真正跑一遍面试时才能讲得清楚。5.3 方案取舍与软硬件协同综合题还会设置一些坑比如题目说现有服务器只有CPU但要求延迟小于10ms该怎么做这种不能直接回答必须买GPU而是要分情况分析如果模型小可以用CPU的AVX-512指令集做性能优化或者用OpenVINO/TVM这样的编译优化工具如果模型太大INT8量化剪枝后可能满足时延预算如果实时性要求极高可以考虑FPGA做流水线。给出原则算法、框架、硬件三方面的优化空间都要先评估而不是一上来就换硬件。笔试问这种题本质是考察你是否具备系统设计的全局视角。回答时一定要条理化先说约束时延、吞吐、功耗、成本再列可行方案最后拍板选型。我当时就因为没有先厘清在线推理还是离线批量处理导致方案设计偏了被同场同学提醒才意识到。这个习惯后来做AI服务部署时非常有用。6. 复盘与备考路线测的是知识更是工程直觉6.1 当时我踩过的坑准备这场笔试时我犯过几个错误大家不值得重复走。第一个错误是忽略了对系统栈的复习把大量时间花在机器学习算法推导上结果笔试考了很多GPU内存模型、OS调度、cache一致性相关的内容。AI异构计算工程师虽然要有算法背景但笔试更偏系统能力。第二个错误是只会写CUDA但不会估算性能导致综合题一筹莫展。后来我强迫自己每写一个kernel先估算理论耗时再和实测对比慢慢建立起了数字直觉。第三个坑是只看博客不动手。网上讲CUDA优化的文章很多但笔试里的细节题比如bank conflict、原子操作、stream overlap只有自己真正写代码遇到奇葩输出才能理解为什么会有这些优化手段。我记得有一次在共享内存里声明了一个二维数组忘了padding结果性能比全局内存版本还差调了一下午。后来再看到相关题我就能直接反应出数据结构该怎么设计。6.2 给后来者可执行的清单如果你现在正在准备AI异构计算相关的笔试或面试下面的清单可以直接照着做把计算机体系结构的核心章节过一遍重点是存储层次、并行计算、缓存一致性、虚拟内存。自己编写并跑通至少5个CUDA程序vectorAdd、matrixMul、reduce、conv或者im2colGEMM、streamOverlap。每个kendel记录性能数据用ncu或nvprof查看占用率、带宽利用率和bank conflict形成调优经验。学会手动计算FLOPs、内存带宽、算术强度以及简单模型的推理延迟预算。了解一个主流深度学习框架的异构执行流程比如TensorFlow的device placement、PaddlePaddle的Executor。做至少一个综合设计题从0到1规划一个AI推理系统的异构方案写清楚每秒处理帧数、时延预算、模型计算量、硬件选型。笔试其实只是门槛后面还有面试环节但很多人的差距在笔试已经拉开了。那些能把为什么讲透的人往往不是刷题刷得多而是真的在GPU编程和系统设计上积累过第一手经验。我自己后来带团队面试时也一样只要候选人能对共享内存的bank conflict展开讲出自己调优时的数据变化我就会认为他具备异构计算的工程底色。6.3 另一个容易被忽略的准备方向除了硬核的GPU知识笔试里偶尔会出现少量机器学习基础题比如问为什么ReLU比sigmoid更适合深层网络BatchNorm在训练和推理时有什么区别。别以为异构计算就不考这些毕竟你写kernel时也要对算子的数值稳定性负责。我当年遇到一道ReLU的简答题差点因为紧张而漏掉计算量小这个点后来冷静下来补上了。这种题不考难度考的就是你在跨领域之间的融会贯通能力。回看2018年那场提前批笔试它对我的价值不只是拿offer而是让我很早就意识到异构计算不是某个单一技能的延伸而是一个交叉领域。只有把硬件架构、并行计算、深度学习框架、系统性能分析这四块拼图都握在手里才能在AI基础设施这个方向上走得更远。如果你现在正打算投这类岗位别怕笔试题难顺着上面这些考点去补大概率能踩中七八成。
RELATED — 相关阅读

相关资讯

LATEST — 最新资讯

最新发布

TODAY — 本日精选

新闻

WEEKLY — 本周精选

新闻

MONTHLY — 本月精选

新闻