尧图建网站 尧图建网站 YAOTU WEB BUILD 免费咨询
ARTICLE DETAIL

资讯详情

深耕网站建设与建站编程的一线实战洞察。

百度AI异构计算笔试题复盘:从体系结构到CUDA优化的能力图谱

百度AI异构计算笔试题复盘:从体系结构到CUDA优化的能力图谱 2018年秋招那会儿百度在校招官网挂出了AI异构计算工程师的岗位同时还放出了两批笔试题。当时这个方向在校招里算是相当小众的存在大部分人还在投算法岗写模型真正愿意去啃GPU底层、研究算力基础设施的人并不多。我刷到第二批这套题的时候第一反应是终于有公司愿意把异构计算从概念变成一个可以筛选人的工程标准了。今天回看这套题再结合这几年做AI基础设施、推理引擎优化的经历我觉得有必要把这套题背后的考察逻辑完整复盘一遍给正在准备类似方向的同学做参考。提示本文不逐题贴出原始题目官方题目不便二次传播重点拆解这套笔试背后实际考察的知识体系、题型思路和复习路径这些内容比背答案更有迁移价值。1. 2018年百度AI异构计算笔试出现的底层原因1.1 为什么是2018年为什么是异构计算2018年前后是深度学习训练规模快速膨胀的时间窗口。图像分类、目标检测、语音识别这些任务的数据集和模型规模都在成倍增长CPU已经明显撑不住训练和线上推理的算力需求。而英伟达的GPU凭借CUDA生态和Tensor Core的早期布局成了深度学习训练的事实标准。百度这时候单独开设AI异构计算工程师岗位本质上是为大规模GPU集群、飞桨框架底层算子、线上推理引擎做人才储备。这篇笔试题的意义不在于考了多少道题而在于它划定了一个AI异构计算工程师的能力边界你既要有计算机体系结构的底子又得懂并行编程、编译优化还得对深度学习算子的计算模式有概念。这是一个典型的需要向下钻取的岗位笔试的第二批题目明显比第一批更侧重工程细节和底层原理。1.2 异构计算工程师和普通AI工程师的分工差异很多同学分不清算法工程师和异构计算工程师的区别。算法工程师的核心任务是设计模型结构、调整训练策略、评估效果指标异构计算工程师的核心任务是让模型在特定硬件上跑得更快、更省显存、吞吐更高。一个是做什么一个是怎么跑得更快。从笔试的考察范围就能看出这种分工差异能力维度算法工程师异构计算工程师模型设计重轻体系结构轻重并行编程可选必选性能优化低优先级核心工作编译原理基本不用需要理解框架源码用到哪读到哪需要深读理解了这层分工你就知道笔试题目为什么长成那样——它不是考你会不会用深度学习框架而是考你有没有能力去写框架底层、改算子实现、压榨硬件性能。2. 笔试覆盖的知识地图四个绕不开的硬骨头复盘下来这套笔试真正想考察的底层能力可以归成四个板块。这四个板块几乎覆盖了异构计算工程师日常工作的全部知识底盘。2.1 计算机体系结构从CPU到GPU的硬件直觉体系结构是这套笔试的地基。题目里大量出现关于存储层次、Cache一致性、总线带宽、内存访问延迟的内容。你如果只是会调API、写Python看到这些题基本是懵的。核心需要掌握的体系结构知识包括存储层次寄存器、L1/L2 Cache、主存、显存、全局内存、共享内存每一级的容量、延迟、带宽差异。GPU典型的全局内存延迟在数百个时钟周期而共享内存延迟在几十个周期这种数量级差异直接决定了算子的设计策略。Cache一致性协议MESI协议的状态转换、写回/写分配策略。GPU里虽然不像CPU那样有复杂的多核Cache一致性需求但理解这套机制对你理解CPU端的并发数据同步非常有帮助。PCIE与异构互联CPU和GPU之间的数据要经PCIE总线搬运带宽远低于显存带宽。笔试中会考察数据搬运时间远大于计算时间的场景引导你思考如何减少主机端与设备端的拷贝。指令级并行流水线、分支预测、向量化指令。GPU的SIMT执行模型本质上就是在指令级并行之上的大规模扩展。备考这个板块我建议把《计算机组成与设计硬件/软件接口》和《深入理解计算机系统》CSAPP的存储层次章节反复读。CSAPP的三章和五章如果吃透了概念辨析和计算推导题基本不会丢分。2.2 并行编程模型线程、同步与通信这一板块直接对应异构计算最核心的异构两个字。笔试会考察CUDA或OpenCL的编程模型、线程组织方式、同步原语、原子操作和内存模型。需要建立清晰的认知框架线程层次Grid、Block、Thread的逻辑层级以及它们如何映射到SM、Warp。理解Warp是SIMT执行的基本调度单位一个Warp内32个线程执行同一条指令分支发散会带来性能惩罚。内存层次寄存器、共享内存、全局内存、常量内存、纹理内存的用途和访问特性。共享内存是片上的带宽极高但容量有限需要程序员手动管理全局内存是大容量、高延迟。同步与并发__syncthreads()的块内同步语义、原子操作的性能影响、死锁的产生条件。通信原语规约reduction、扫描scan、分散scatter、聚集gather这些并行模式的实现思路。一个很典型的笔试考察点是让你写出一个线程块内多线程协作的规约求和代码或者问你如果将100万个整数的累加任务分配给1024个线程应该如何组织线程结构、如何处理最后的结果合并。这类题考察的不只是语法而是你对线程并发模型和数据依赖的理解深度。2.3 编译与底层优化算法效率的另一半异构计算工程师写代码不是写完就行而是写完要能看出生成的指令在硬件上怎么跑。笔试里会出现关于编译器优化、循环变换、访存局部性优化的问题。这个板块的核心内容包括循环优化循环展开loop unrolling、循环交换loop interchange、循环分块loop tiling。矩阵乘法访存优化中循环分块是最基础也是最高频的优化手段。访存局部性时间局部性和空间局部性。为什么访问二维数组时按行访问比按列访问快得多——背后的Cache line机制是什么。编译优化选项-O2、-O3、-funroll-loops、-marchnative等选项的差异和适用范围。算法复杂度与常量因子笔试会提醒你同一个算法复杂度下常数优化指令数减少、分支优化、内存对齐往往决定了实际性能。我在实际优化经历中逐步意识到编译优化这块知识是很多科班出身但没做过底层开发的同学的盲区笔试中出现这类题目就是想把只会调参的人和能真正写推理引擎的人区分开。2.4 深度学习算子的计算模式与框架原理既然岗位名字带了AI两个字笔试难免要考察你对深度学习算子计算模式的理解。这里的考察不是让你推导反向传播公式而是考察你能否判断这个算子适合在GPU上怎么执行。常见的考察点**矩阵乘法GEMM**的访存与计算比、如何用分块来提升访存复用。卷积操作转换为矩阵乘法im2col的原理与代价im2col会扩大显存占用但能利用高度优化的GEMM内核。BatchNorm、ReLU、Pooling这类elementwise或reduce类算子的并行度特征。Tensor Core / FMA指令的基本概念2018年时Tensor Core已经出现在V100上笔试可能会涉及混合精度训练的基础认知。如果你只会用PyTorch或飞桨调算子不知道算子底层的计算模式这部分的题只能靠猜。反过来如果你认真写过一次CUDA的卷积或GEMM算子这部分就是送分题。3. 三类高频题型与具体解题思路还原虽然不能原样复述题目但根据我的复盘和同期同学的反馈第二批笔试题基本可以分成概念辨析、计算推导、代码优化三类。下面分别给出各自的典型形态和解题方法你要做的就是按这个方法去练同类题。3.1 概念辨析题内存一致性与缓存同步这类题一般会给出一段多线程或异构系统下的读写场景问你某个现象产生的原因或者让你判断几种同步机制的区别。典型形态场景两个CPU核心同时读写一个共享变量其中一个核心的修改没有被另一个核心立即看到。问原因。变体GPU kernel中多个Block同时往全局内存的同一地址做累加得到的结果不确定。问该如何解决。解题思路是抓住三点第一理解可见性问题的根源。现代CPU和GPU都有多级缓存写操作可能停留在缓存中没有立刻刷到内存其他计算单元读到的就是旧值。解决思路是加内存屏障或使用原子操作。第二区分原子操作和锁的区别。原子操作是硬件级别保证某个操作的不可拆分性而锁是软件层面实现临界区的互斥。GPU里原子操作的开销较高大量冲突时性能会急剧下降。第三规范答题先说原因再说解决方案最后补充方案的代价。例如使用原子操作可以保证结果正确但高并发下会产生严重的竞争冲突建议先在各Block内用共享内存做局部规约再对局部结果使用原子操作。这种答题结构会显得对工程有实际判断。3.2 计算推导题带宽、延迟与加速比这类题是区分度最高的。给出一组硬件参数如CPU时钟频率、GPU核心数、显存带宽、PCIE带宽要求你估算某个算子的理论计算时间或访存时间判断瓶颈在哪里。典型形态参数GPU有5120个CUDA核心核心频率1.5GHz显存带宽900GB/s计算一个矩阵乘法C A x B矩阵维度2048x2048类型为FP32。问题计算这个算子的理论计算时间、访存时间判断是计算密集型还是访存密集型。解题步骤要牢记第一步算计算量。一个N x N的矩阵乘法需要大约2N³次浮点运算乘加各N³次但FMA指令可以把乘和加合并所以通常按N³次FMA或2N³次FLOPs来估算具体看你用什么口径。N2048时计算量约为2 x 2048³ 17.18 GFLOPs。第二步算理论峰值。5120个核心 x 1.5GHz x 2FMA算两次浮点计算≈ 15.36 TFLOPs。第三步算理论计算时间17.18GFLOPs / 15.36TFLOPs ≈ 1.12ms。第四步算访存量。FP32矩阵每个元素4字节A、B、C加起来是3 x 2048 x 2048 x 4字节 ≈ 50.33MB。按900GB/s算访存时间约0.056ms。第五步对比结论计算时间远大于访存时间所以这是计算密集型算子优化重心应该放在提升计算效率、减少指令流水停滞而不是拼命压访存。这类题的关键是一定要把单位换算清楚。TB和GB、ms和us之间差三个数量级错一个单位全盘皆输。我当年笔试时就在单位上吃过亏后来凡是算峰值性能一律先统一到FLOPs和Bytes再代入公式。3.3 代码优化题算子实现与性能陷阱代码题一般不会让你在笔试环境里写出完整可编译的CUDA程序更多是给一段有明显问题的伪代码或简短C代码让你指出性能瓶颈并给出改进方案。典型的低效代码形态二维矩阵按列遍历空间局部性差改进方式是按行遍历或调整循环顺序。全局内存访问不对齐导致事务放大改进方式是确保数据类型与访问长度对齐、使用float4向量化加载。每个线程只处理一个元素启动开销太大改进方式是每个线程处理多个元素grid-stride loop减少线程启动和调度开销。使用atomicAdd在高冲突场景下做全局累加改进方式是先共享内存局部规约。缺少共享内存复用每个线程反复访问全局内存中同一个数据改进方式是先协作加载到共享内存。笔试答代码优化题不要只写应该使用共享内存要说明为什么共享内存快、快多少、使用后还需要考虑哪些问题如bank conflict、同步开销。这样能体现出你真正理解优化的本质而不是背了几个专有名词。4. 从笔试到Offer一份可执行的备考路线如果你现在正打算投递AI异构计算、高性能计算或推理引擎方向的岗位这条路线是我实测下来比较高效的路径总周期可以压缩到两到三周。4.1 第一周体系结构和并行编程原理打通不要一上来就刷LeetCode先把《深入理解计算机系统》的存储层次、汇编基础、并发那几章读完再配合《CUDA C Programming Guide》的前半部分搞清楚线程模型和内存模型。这个阶段的目标不是记住所有细节而是建立代码最终在硬件上如何执行的直觉。同步做的小练习跑一个简单的CUDA向量加法对比不同数组大小、不同Block/Thread配置下的耗时变化。用nvprof或ncu分析跑一次矩阵乘法的核函数观察内存吞吐和计算吞吐的数字。尝试把二维矩阵转置的CPU版本和GPU版本都写一遍体会访存连续性与转置带来的性能差异。4.2 第二周算子实战与典型算法手写这个阶段动手写三个经典算子矩阵乘法、规约求和、卷积im2col版本即可。这是异构计算笔试中最常见的三个代码原型也是面试中经常让你现场写的三个热点。矩阵乘法的优化过程建议按下面这条路线走每一步都记录性能数字朴素三重循环版本baseline交换循环顺序利用Cache行主序下i-k-j循环比i-j-k快很多循环分块tiling提升数据复用引入共享内存分块加载A和B向量化访存float4解决对齐问题处理bank conflict共享内存填充padding使用寄存器缓存、展开循环规约求和重点理解树形规约的思想和__syncthreads()的必要性。卷积算子重点理解im2col的空间换时间思路以及为什么在显存充足时im2col方案往往比直接卷积更容易达到高性能。4.3 第三周真题模拟与知识补盲第三周开始做套题模拟。你可以把笔试题目归类按概念题、计算题、代码题三个方向各找10题左右练手。练习时给自己限时概念题5分钟一题、计算题10分钟一题、代码题15分钟一题。模拟完一定要做复盘不光是复盘错题还要复盘哪些板块用时太长。如果体系结构题每题都花10分钟说明前面两章的基础没打牢回头重读而不是继续刷题。知识补盲的重点熟悉飞桨、PyTorch的目标检测和Transformer类模型的常见算子知道其中有哪些可以优化的点。了解TensorRT、XLA等推理编译器的基本思想算子融合、图优化、内存池。了解混合精度训练、bf16/fp16的特点以及为什么现代AI基础设施都在向低精度计算发展。5. 复盘之后那些笔试里踩过的坑和考场经验5.1 第一个坑轻视存储层次导致的连锁失误我当年复习时花了大量时间研究并行模型和GPU语法几乎没有认真推导过存储层次的延迟差异。结果遇到一道为什么GPU访存要尽量减少全局内存访问次数的题虽然知道要写共享内存更快但写不出具体的数据对比回答显得很空洞。后来我把各存储层次的数量级背到滚瓜烂熟全局内存延迟约400-800周期共享内存延迟约20-30周期寄存器基本是0周期。这种数量级认知让你的优化建议有说服力。笔试和面试中给出共享内存比全局内存快一个数量级以上这个量级判断比泛泛地说共享内存快要扎实得多。5.2 第二个坑单位换算出低级错误有一类计算推导题给出的显存带宽是900 GB/s核心频率是1.48 GHz问你矩阵乘法的理论耗时时我把GB当成了Gbit来算结果答案差了8倍。这种错误非常可惜因为你的思路完全正确仅仅是一个单位换算就丢掉了整道题的分数。我的建议是所有计算推导题动笔前先把所有数据写成标准形式。带宽写为字节每秒、容量写为字节、时间写为秒计算完再换算成毫秒或微秒。这样虽然慢一点但正确率会高很多。5.3 第三个坑代码题只写结论不写原因笔试里遇到过一道代码优化题代码里有个明显的循环内重复调用pow()函数的问题。我当时直接写了把 pow() 提到循环外这个结论没有解释为什么。但事实上编译器和硬件的优化原理是pow()是一个昂贵的数学库函数编译器一般不会自动将它提到循环外因为函数可能改变全局状态或有副作用手动提取之后循环体里只剩一次指数运算同时循环内余下的部分更容易触发指令级并行。如果我在回答里加上这层原因这道题就能从基本答对变成答得漂亮。笔试考的不只是结论更是你得出结论的推导能力和对硬件行为的理解深度。这一点在异构计算领域尤为明显。最后说点我的个人体会。从2018年到现在AI异构计算的面试考法一直在变工具越来越抽象、硬件越来越复杂但核心考的东西一直没变你能不能把一段代码放进真实的硬件执行模型里看出它快在哪、慢在哪。如果你正在准备这类岗位别只刷题库、背结论回到原理层面把体系结构、并行模型、编译优化这三块基础打牢任何形式的笔试题都难不住你。
返回列表