又到一年校招季整理旧资料时翻到当年收藏的百度2019校招AI异构计算工程师笔试题第三批想了想还是把这些题和我的解题思路完整写出来。百度那会儿是全国少数在秋招里单独设立异构计算岗位的大厂题目不是简单的八股文而是在认真考察候选人有没有在真实硬件上思考问题的习惯。第三批又是三批卷子里综合性最强的一批覆盖了计算机体系结构、并行计算理论、CUDA编程、深度学习推理优化还有嵌入式场景的系统设计题。这篇文章不只给答案我会把每道题背后的考察意图、解题推理链以及对应到实际工程中的坑全部拆开讲适合准备AI基础设施、AI芯片、推理引擎方向校招的同学参考也适合刚入行异构计算方向、想建立系统认知的工程师阅读。1. 这批试题到底在考什么从考点分布看异构计算岗的能力模型2019年的AI行业处在一个特殊节点。深度学习框架趋近成熟AI芯片开始批量落地互联网大厂的技术团队从“用GPU训练模型”逐步转向“自己优化算子、设计推理系统”。百度在那一届校招里把异构计算工程师单独列出来本身就说明这个岗位的技术栈已经足够垂直。笔试分三批前两批相对侧重基础编程和机器学习常识第三批的难度明显上了一个台阶——它默认你掌握了Linux、C、CUDA至少其中两项并在此基础上考察系统设计能力。从网上流传的考生回忆版来看第三批题目整体风格偏“工程导向”几乎没有纯背诵题。考点大致可以归为五类考点方向典型题目形态考察能力计算机体系结构CPU与GPU设计哲学对比、访存延迟与带宽计算对底层硬件的直觉理解并行计算理论Amdahl定律计算、并行效率分析理论模型与工程差距的认知并行编程实践向量加法、shared memory bank conflict、矩阵乘法优化真实写kernel的功力深度学习算子优化BatchNorm融合、1x1卷积访存分析推理引擎常见优化手段嵌入式与系统设计有限内存下的模型部署方案约束条件下的全局设计能力这套能力模型和普通软件开发岗的笔试有本质区别。普通研发岗考的是数据结构和算法基本功异构计算岗考的是你能否在“带宽、延迟、并行度、内存容量”这四个维度的约束下设计出高性能计算方案。笔试题目看起来是散点式的实际上全部绕着一个核心问题打转当计算从单核走向多核、从CPU走向GPU和专用加速器时你如何重新理解“快”这个字。第三批题目里还有一个隐藏的考察点做题速度之外的工程直觉。有些题不是算得出来就能得分而是在看你会不会用数量级估算快速定位瓶颈。这一点在后面的题目解析中我会重点强调。2. 体系结构题目拆解GPU的延迟隐藏与吞吐量思维2.1 GPU和CPU为什么走上了不同的设计路线这批笔试中有一道简答题字面意思是请解释GPU和CPU在延迟和吞吐量上的不同设计取舍以及为什么深度学习推理更适合在GPU上运行。这道题看着基础其实筛掉了很多人。大多数候选人能说出“CPU核心少但频率高、GPU核心多但频率低”但说不清楚这套差异背后的本质原因。CPU的核心设计目标是尽可能降低单条指令的延迟。为了做到这一点CPU把大量晶体管花在了分支预测、乱序执行、大容量缓存、推测执行这些微架构组件上。以一颗现代服务器CPU为例核心数通常在几十个以内但每个核心的指令级并行ILP挖掘能力非常强单线程性能是GPU每个线程的几十倍。同时CPU必须保证低延迟响应因为操作系统、数据库、网络服务这些负载充满了不可预测的分支和依赖程序无法提前预知下一条指令要访问哪块内存。GPU走的完全是另一条路。GPU把晶体管重点用在计算单元和线程调度器上核心数量动辄数千甚至过万。单个线程的执行效率很低但GPU通过同时维护大量线程来掩盖访存和指令延迟。当一个线程等待全局内存返回数据时调度器立刻切换执行其他线程。这个机制叫延迟隐藏latency hiding也可以理解成一种极端的TLP线程级并行。深度学习推理为什么天然适配GPU核心原因是计算模式和访存模式的匹配。卷积、矩阵乘这类算子里包含海量并行度和数据复用卷积核在多个通道上滑动时同一个权重会被大量输入像素复用。GPU的大规模并行线程配合较高的访存带宽恰好能把这种并行度和复用率吃满。相反如果让GPU跑一个分支密集、访存随机的单线程程序性能会非常难看因为吞吐量架构在低效分支和离散访存面前几乎没有用武之地。2.2 一道关于延迟、带宽和在飞行字节数的量化题第三批里有一道计算题极具代表性考的是GPU全局内存延迟大约400到600个时钟周期的情况下需要多少个并发访存请求才能把显存带宽打满。这类题在面试中出现的频率非常高因为它考察的不是背参数的能力而是能不能用吞吐量视角看问题。基本计算模型是Little定律的硬件版本在飞行中的数据量 带宽 × 延迟。假设一张GPU的显存带宽是1TB/s访存延迟是500ns那么在任意时刻必须在内存子系统里“在飞行”的数据量是1TB/s × 500ns 500KB。也就是说如果活跃的并发访存数据量少于500KB访存流水线就喂不满带宽就一定有浪费。换算成请求数量就更直观。一个float4宽度为16字节500KB除以16字节约等于32000个访存请求同时在飞。GPU正是通过极高级别的线程并发来维持这个数量的。SM能容纳的线程和访存指令越多延迟隐藏能力越强。这也是为什么在选CUDA block大小时要考虑占用率occupancy一个block里开128个线程也许能跑但开512个线程往往更容易把访存流水线压满。这道题的工程意义在什么时候体现做算子优化时如果我发现某个kernel的算术强度太低比如纯访存型算子第一步就是去测它的实际带宽利用率。利用率上不去优先怀疑并发度是否足够再怀疑访存模式是否产生了过多的内存事务。这不只是笔试考点每次用Nsight Compute分析kernel时都在用。3. Amdahl定律与并行效率题理论加速比与工程现实的差距3.1 一道典型的Amdahl定律计算题笔试里出现了一道看起来非常“教科书”的题某个应用程序有10%的代码只能串行执行剩余90%可以完美并行。请求解它在64核处理器上能达到的理论加速比并讨论如果串行比例扩大到20%加速比会如何变化。套用公式加速比 S 1 / (串行比例 并行比例 / 核心数)。10%串行、64核时S 1 / (0.1 0.9/64) ≈ 8.5。也就是说即便有64个核心最理想也只能获得不到一个数量级的加速。如果串行比例变成20%S 1 / (0.2 0.8/64) ≈ 4.7加速比直接腰斩还多。这道题的标准答案不难但能拿高分的人会在最后补一个关键认知Amdahl定律是一个悲观模型它假设问题规模固定。如果问题规模随核心数增大而增大对应的加速比更适合用Gustafson定律描述。在实际AI场景里模型规模和数据集通常随着算力增强而扩大所以Gustafson定律往往比Amdahl定律更贴近真实。3.2 串行比例为什么是面试官精心埋的陷阱笔试里那个“10%串行比例”很容易让人觉得10%没什么大不了。但实际上恰恰是这部分主导了扩展性的上限。工程中真正的瓶颈也往往来自这些不起眼的串行片段。我在实际项目里遇到过多次类似情况。把一个数据预处理流程从单核改成多核最初以为只要把for循环改成OpenMP并行就行结果实测8核加速比只有2.4。原因就是每个线程完成后要往同一个日志缓冲区写行日志写入是有锁的随着线程数增多锁竞争越来越严重热点全部集中到串行区。这跟Amdahl定律的串行比例上升是同一个问题只是工程表现更加隐蔽。多核场景下串行区通常来自这几个地方内存分配器尤其是带全局锁的分配器、共享计数器和统计结构、日志与埋点、任务队列的入队出队操作、负载不均衡导致的收尾同步。面试中如果能在Amdahl定律题之后自己主动引出这些场景考官会觉得这个人是真正做过并行系统的人而不是只会背公式。3.3 负载不均衡与多核扩展性的实际案例笔试中还有一道扩展题问的是“在多核处理器上获得线性加速比为什么在现实中几乎不可能”要求列出至少三种限制因素。参考答案里通常包括串行代码段、同步开销、负载不均衡、缓存一致性和内存带宽竞争。每一条展开都能对应一个8核升16核收益不明显的真实原因。其中负载不均衡是我最想强调的。并行任务切分如果不均匀整体执行时间由最慢的那个线程决定。比如一个任务被分成16份其中一份的执行时间是其他的2倍那并行部分的有效利用率就损失在等待上这等价于凭空增加了串行比例。在做算子切分时我习惯对样本按计算量排序后再切块而不是简单按索引均分。AI推理场景里输入样本的序列长度往往差异巨大如果每个batch按固定数量分配有的线程处理长序列有的处理短序列整个kernel的执行时间就被长序列拖住。这时用动态任务窃取work stealing或者按token数分配batch才能让扩展性回到正轨。4. CUDA编程题从向量加法到Shared Memory与bank conflict4.1 向量加法的代码题陷阱第三批笔试有一道CUDA向量加法的编程题要求写一个完整的kernel把两个float数组相加。看起来是CUDA入门第一课但题目里故意加了一个要求说明block大小和grid大小的选取依据。这道题考察的不只是会不会写三行代码而是有没有真正理解线程配置和执行效率的关系。基础的kernel长这样__global__ void vectorAdd(const float* a, const float* b, float* c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { c[idx] a[idx] b[idx]; } }调用时block大小选多少这个问题学问很大。选太小的block比如32每个SM能容纳的block数量有限可能无法隐藏访存延迟。选太大的block比如1024又容易在block尾部出现大量线程因边界判断而空转。工程上常见的选择是128到512之间的值具体要结合SM的最大线程数和寄存器占用来确定。边界判断导致的分支发散divergence也值得玩味。虽然条件判断idx n几乎不会造成实际性能损失但一个严谨的kernel应该优先保证数组长度是block大小的整数倍。如果长度不能整除可以用grid-stride loop来减少尾部线程的浪费__global__ void vectorAddStride(const float* a, const float* b, float* c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; int stride gridDim.x * blockDim.x; for (; idx n; idx stride) { c[idx] a[idx] b[idx]; } }vectorAdd这类入门题在实际工作中没什么用但它背后的线程配置、边界处理、grid-stride策略在写任何数据并行kernel时都用得上。笔试出这道题的目的就是看候选人有没有从“能跑”往“跑得快”走的意识。4.2 Shared Memory与bank conflict比数据结构还常考第三批笔试中最有区分度的题目之一是一道关于shared memory bank conflict的分析题。题目大致是假设有一个长度为N的float数组存放在共享内存中线程t需要访问数组的第2t个元素即以步长2访问请分析这会发生几路bank conflict并提出改进方案。共享内存的硬件结构是连续的32个bank每个bank的宽度是4字节。因此bank号通常按地址计算地址偏移除以4再对32取模。float类型恰好4字节所以第i个float落在bank (i % 32)上。线程t访问第2t个元素即float序号2tbank号为(2t) % 32。当t0时bank0t1时bank2t16时bank32%32bank0。这时线程0和线程16落在同一个bank上线程1和17落在bank2上以此类推。同一个bank在同一个周期内只能为一个线程服务所以要分两个周期来完成这批访问这就是2路bank conflict。这个知识点为什么在笔试里反复出现因为bank conflict直接决定了shared memory能否跑满带宽。很多CUDA新手把shared memory当成一个“更快的全局内存”以为把数据拷贝进去就自动变快完全没意识到它的内部访问规则。如果不同线程同时访问同一bank实际吞吐量会被显式折减。改进方案不止一个要看具体访问模式。对于步长2的访问最直接的手段是改变数组的存储布局在每行数据末尾填充一个floatpadding让相邻行的起始bank错开从而避免行间冲突。另一个通用思路是把对数组的访问改成按float4向量化让一个线程一次性读取4个float既减少访存指令数量又改变了bank分布规律。还有一类方案是使用共享内存的广播机制如果多个线程访问的是完全相同的地址硬件能在一个周期内完成广播这不算冲突。4.3 矩阵乘法的Tiling优化思路题矩阵乘法优化是异构计算笔试的“压轴常客”。第三批里出了一道没有要求写完整代码、只要求“描述如何利用shared memory优化矩阵乘法”的题这其实比写代码更难因为它要求候选人脑子里有一张完整的kernel执行图。基础版本是每个线程计算C矩阵的一个元素内层循环要读A的一行和B的一列对全局内存的访问量是O(N^3)性能受制于访存带宽。Tiling优化的核心思想是把大矩阵切块每次把A的一个tile和B的一个tile搬进shared memory在片内完成计算从而大幅减少对全局内存的重复访问。我在实际优化时会这样安排设tile大小为32x32或16x16每个block加载一个tile的A和一个tile的B到共享内存线程块内的256个线程分别负责tile内不同的输出元素。计算每个输出元素时遍历k维度上的一长条数据但每次内层循环的A和B数据都直接从shared memory读取全局内存只被精准访问两次。如果再加上double buffering把下一块tile的数据在计算当前tile时提前搬运访存和计算还能进一步重叠。这类题目最终考察的是对访存层次结构的理解和利用而不是某个具体优化trick。无论是CPU上的循环分块还是GPU上的shared memory tiling底层逻辑都是同一句话数据在寄存器里跑得最快在缓存和共享内存里次之在全局内存里最慢算法设计要尽量提升数据在被慢速存储访问前的复用次数。5. 深度学习推理优化题BatchNorm融合与1x1卷积的访存瓶颈5.1 BatchNorm在推理阶段如何融合到卷积层第三批笔试出现了三道和深度学习推理算子优化相关的题其中一道是说明BatchNorm在推理阶段如何合并进前一个卷积层并推导合并公式。BatchNorm在训练阶段依赖batch内统计量需要维护均值和方差还有可学习参数gamma和beta。推理阶段不再使用batch统计而是使用训练完毕后固定的全局均值和方差。一个重要事实是BatchNorm在推理时对每个通道执行的是线性变换所以它完全可以合并到相邻卷积层的权重和偏置里。卷积层输出是y W * x bBatchNorm的推理变换是z gamma * (y - mean) / sqrt(var eps) beta把y代入可以得到z (gamma * W / sqrt(var eps)) * x (gamma * (b - mean) / sqrt(var eps) beta)也就是说新的卷积权重是W乘以scale新的偏置是b乘以scale再平移。在推理引擎实现时只需要在模型加载阶段对每一层做一次参数改写就能完全消除BatchNorm在推理时的额外计算。这在网络层数较深时可以节省不少访存和计算开销。我从事推理引擎优化后才发现这类“参数折叠”的手段不止BatchNorm融合这一种。卷积后面的ReLU可以融合进上一个kernel的epilogue阶段LayerNorm的仿射变换可以折叠进前面的线性层INT8量化中的per-channel scale也可以提前乘到权重里。推理引擎优化的第一课永远是把数学上能合并的线性变换全部合并减少kernel数量而不是一上来就去魔改CUDA代码。5.2 1x1卷积为什么是访存密集型算子另一道题问1x1卷积的计算量不大但为什么在GPU上性能往往不理想请从访存特征角度分析。1x1卷积本质上是跨通道的线性组合对每个空间位置的各个通道做矩阵向量乘。它的计算量是N×H×W×Cin×Cout看起来不小但在现代GPU上真正的瓶颈往往是数据搬运。每个输入像素要被Cout个输出通道使用这意味着输入数据的复用率取决于Cout的大小如果Cout不大每个输入元素只被读几次就丢弃算术强度每次访存对应的浮点计算次数很低。低算术强度意味着kernel被访存瓶颈主导GPU的高算力完全使不上劲。优化这类算子的核心思路是提高数据复用比如让一个线程块处理多个空间位置把输入块在shared memory里缓存下来使多个输出通道共享同一份输入数据。另一个方向是融合相邻算子比如把1x1卷积后面的ReLU、BN、残差相加都放进同一个kernel的末尾处理避免中间结果写到全局内存再读出来。5.3 从题目到推理引擎的工程实践笔试题里有一道开放题让设计一个高效的卷积kernel。这道题没有唯一答案但高分答案通常包含这些层次第一明确卷积的shape和硬件平台第二选择合适的数据布局NCHW、NHWC还是NC1HWC2第三考虑im2col加GEMM还是直接卷积或者Winograd第四说明tiling策略、向量化宽度和double buffering。Winograd是当年笔试的热门考点之一。它通过减小卷积乘法次数来降低计算量代价是增加变换阶段的访存和数值误差适合小卷积核3x3大通道数的场景。我做实际项目时对Winograd的态度是先写好标准GEMM版本和直接卷积版本作为baseline再根据profile数据判断值不值得上Winograd。因为Winograd的预处理需要额外的内存拷贝和格式转换如果数据布局没优化好节省的计算时间会被访存开销吃掉。这类题目对考生的真正要求不是背出一堆优化名词而是具备“选择一个优化手段之前先分析它在当前硬件和当前shape下的收益”的工程判断力。6. 嵌入式与边缘场景题有限内存下的模型部署方案6.1 当年那道让我印象深刻的系统设计题第三批笔试的压轴题是一道系统设计题假设要把一个图像分类模型部署到内存只有几十MB的嵌入式设备上可用的硬件只有一颗低功耗CPU加一个小的NPU加速器请给出完整的部署优化方案。这道题把之前所有考点知识串成了一条线。它没有限定用什么框架也没有限定模型结构考察的是候选人在真实约束条件下做技术决策的能力。我当时的答题思路是先分类讨论约束内存受限、算力受限、功耗受限、模型结构未知。然后按层次给出方案。第一层是模型层面的优化。优先检查模型结构是否过大考虑知识蒸馏把大模型换成小模型或者使用神经架构搜索出来的轻量网络。这一层的收益是数量级的比任何算子优化都明显。第二层是权重压缩。对模型做INT8量化可以把模型体积和带宽需求压缩到原来的四分之一。如果精度损失在可接受范围内甚至可以尝试INT4量化配合混合精度。剪枝则把不重要的通道和连接删掉进一步减少计算量和存储。值得注意的是量化不只是把权重从FP32换成INT8还需要校准激活值的分布范围以及处理量化敏感层比如检测头的输出层。第三层是运行时优化。算子融合减少中间张量的内存占用内存池复用避免动态分配和碎片计算图优化把可以并行执行的节点调度到CPU和NPU上同时运行。嵌入式设备上的内存池设计与服务器端完全不同服务器随便malloc没问题嵌入式上频繁分配释放容易出现碎片必须从一开始就规划好各个张量的生命周期。第四层是硬件适配。NPU对算子的支持往往有限不是所有层都能直接跑在NPU上。有些层需要切成CPU上的子图执行。这时候CPU和NPU之间的数据拷贝可能是最大的性能杀手所以要尽量减少CPU和NPU之间的转换次数尽量把大块数据一次性切到NPU侧处理。6.2 这类设计题的答题方法论这道题能拿高分的核心不是方案多全而是是否先讲清楚约束和优化空间。我在实际面试别人时也喜欢出这类开放题最怕听到的回答是一上来就背量化剪枝名词完全不管设备的内存上限和算力上限。正确的方法是先画一条技术链路模型结构简化、参数压缩、运行时优化、硬件映射。每往前走一步都要附带代价分析。比如量化会带来精度损失剪枝会改变模型结构导致NPU算子映射更复杂算子融合可能增加寄存器压力。方案不是越高级越好而是在约束下收益最大、代价可接受的那个。笔试中的这类题本质上是看候选人有没有系统思维。我曾经把一个模型从300MB压缩到20MB并在嵌入式设备上跑通靠的恰恰是每一层都做一点优化50MB来自INT8量化100MB来自通道剪枝剩余的多余部分通过修改数据布局和内存复用解决。没有哪一种技术是银弹组合起来才有意义。6.3 从2019到今天的边缘异构趋势回想2019年嵌入式AI部署还集中在手机和安防设备上MCU级别的机器学习更多是跑跑裸机上的轻量识别。这几年边缘场景对异构计算的需求变化也非常明显尤其是工业控制和实时通信场景MCU加NPU、CPU加实时内核的架构已经成为趋势这类架构强调的正是异构计算中的“异构”二字不同的计算单元执行不同类型的任务MCU负责精确到微秒的实时控制和协议解析NPU/加速器负责AI推理两边通过共享内存或高速总线通信。这和笔试里的GPU优化题目其实是同一个世界观异构计算不是把计算单元堆在一起而是让每个计算单元做自己最适合的事情并管理好数据在单元之间的流动。在嵌入式设备上这种设计还要额外考虑功耗和确定性问题难度比纯GPU场景更高。但核心的分析方法没有变先算清楚瓶颈在哪再决定哪些计算放哪个单元最后为数据通路做优化。7. 备考这批题给我留下的东西给后来者的三点建议7.1 与其刷题不如亲手写一遍CUDA并验证想法回头看这套笔试题里最难准备的不是某个具体知识点而是“对性能有没有直觉”。这种直觉只能靠亲手写kernel、跑profile工具、对比不同实现的数据来建立。准备这类笔试时比起刷几百道LeetCode不如把矩阵乘法从naive版一步步优化到tiled版本再用Nsight Compute看一眼实际带宽利用率。每做一个阶段记录一次性能数据分析一次变化原因。这个过程做上两三个算子对访存层次、bank conflict、occupancy这些概念的理解会远比看书深刻。7.2 把理论题和工程案例绑定记忆这套笔试题有一个特点理论题都能在真实工程里找到对应场景。Amdahl定律对应多核任务扩展性bank conflict对应算子访存效率BatchNorm融合对应推理引擎的图优化。备考时如果把每个知识点都绑定一个实际工程案例来记不仅记得牢面试时还能答出别人答不出的“为什么”。比如我后来做BERT推理优化时发现Self-Attention的计算里Q、K、V矩阵乘法和后面的softmax、注意力池化之间的数据流天然就适合算子融合因为中间结果的体积巨大但计算又非常简单。如果能从推理引擎的角度理解这套笔试题这一章前文的很多题目其实都找到了落地的影子。7.3 数值估算能力是隐藏加分项这批卷子中好几道题都在考察数量级估算能力比如访存延迟需要多少并发来掩盖、嵌入式设备能跑多大模型。这种能力只能在平时刻意练习看到任何数据都随手做一次粗略计算一个FP32模型多少权重带宽需要多少显存够不够估算完再和实测对比长期下来对硬件性能的直觉会敏锐很多。最后再分享一个很实用的经验准备异构计算方向笔试和面试尽量在简历里附上自己的算子优化记录——原来性能多少优化后多少中间遇到什么问题怎么解决的。面试官问技术细节时这些材料远比空谈理论更有说服力。这套笔试题虽然已经是好几年前的但今天再看它考察的那些底层能力和思维方式在AI基础设施方向依然完全适用。 百度2019校招AI异构计算工程师笔试真题解析与工程实践
又到一年校招季整理旧资料时翻到当年收藏的百度2019校招AI异构计算工程师笔试题第三批想了想还是把这些题和我的解题思路完整写出来。百度那会儿是全国少数在秋招里单独设立异构计算岗位的大厂题目不是简单的八股文而是在认真考察候选人有没有在真实硬件上思考问题的习惯。第三批又是三批卷子里综合性最强的一批覆盖了计算机体系结构、并行计算理论、CUDA编程、深度学习推理优化还有嵌入式场景的系统设计题。这篇文章不只给答案我会把每道题背后的考察意图、解题推理链以及对应到实际工程中的坑全部拆开讲适合准备AI基础设施、AI芯片、推理引擎方向校招的同学参考也适合刚入行异构计算方向、想建立系统认知的工程师阅读。1. 这批试题到底在考什么从考点分布看异构计算岗的能力模型2019年的AI行业处在一个特殊节点。深度学习框架趋近成熟AI芯片开始批量落地互联网大厂的技术团队从“用GPU训练模型”逐步转向“自己优化算子、设计推理系统”。百度在那一届校招里把异构计算工程师单独列出来本身就说明这个岗位的技术栈已经足够垂直。笔试分三批前两批相对侧重基础编程和机器学习常识第三批的难度明显上了一个台阶——它默认你掌握了Linux、C、CUDA至少其中两项并在此基础上考察系统设计能力。从网上流传的考生回忆版来看第三批题目整体风格偏“工程导向”几乎没有纯背诵题。考点大致可以归为五类考点方向典型题目形态考察能力计算机体系结构CPU与GPU设计哲学对比、访存延迟与带宽计算对底层硬件的直觉理解并行计算理论Amdahl定律计算、并行效率分析理论模型与工程差距的认知并行编程实践向量加法、shared memory bank conflict、矩阵乘法优化真实写kernel的功力深度学习算子优化BatchNorm融合、1x1卷积访存分析推理引擎常见优化手段嵌入式与系统设计有限内存下的模型部署方案约束条件下的全局设计能力这套能力模型和普通软件开发岗的笔试有本质区别。普通研发岗考的是数据结构和算法基本功异构计算岗考的是你能否在“带宽、延迟、并行度、内存容量”这四个维度的约束下设计出高性能计算方案。笔试题目看起来是散点式的实际上全部绕着一个核心问题打转当计算从单核走向多核、从CPU走向GPU和专用加速器时你如何重新理解“快”这个字。第三批题目里还有一个隐藏的考察点做题速度之外的工程直觉。有些题不是算得出来就能得分而是在看你会不会用数量级估算快速定位瓶颈。这一点在后面的题目解析中我会重点强调。2. 体系结构题目拆解GPU的延迟隐藏与吞吐量思维2.1 GPU和CPU为什么走上了不同的设计路线这批笔试中有一道简答题字面意思是请解释GPU和CPU在延迟和吞吐量上的不同设计取舍以及为什么深度学习推理更适合在GPU上运行。这道题看着基础其实筛掉了很多人。大多数候选人能说出“CPU核心少但频率高、GPU核心多但频率低”但说不清楚这套差异背后的本质原因。CPU的核心设计目标是尽可能降低单条指令的延迟。为了做到这一点CPU把大量晶体管花在了分支预测、乱序执行、大容量缓存、推测执行这些微架构组件上。以一颗现代服务器CPU为例核心数通常在几十个以内但每个核心的指令级并行ILP挖掘能力非常强单线程性能是GPU每个线程的几十倍。同时CPU必须保证低延迟响应因为操作系统、数据库、网络服务这些负载充满了不可预测的分支和依赖程序无法提前预知下一条指令要访问哪块内存。GPU走的完全是另一条路。GPU把晶体管重点用在计算单元和线程调度器上核心数量动辄数千甚至过万。单个线程的执行效率很低但GPU通过同时维护大量线程来掩盖访存和指令延迟。当一个线程等待全局内存返回数据时调度器立刻切换执行其他线程。这个机制叫延迟隐藏latency hiding也可以理解成一种极端的TLP线程级并行。深度学习推理为什么天然适配GPU核心原因是计算模式和访存模式的匹配。卷积、矩阵乘这类算子里包含海量并行度和数据复用卷积核在多个通道上滑动时同一个权重会被大量输入像素复用。GPU的大规模并行线程配合较高的访存带宽恰好能把这种并行度和复用率吃满。相反如果让GPU跑一个分支密集、访存随机的单线程程序性能会非常难看因为吞吐量架构在低效分支和离散访存面前几乎没有用武之地。2.2 一道关于延迟、带宽和在飞行字节数的量化题第三批里有一道计算题极具代表性考的是GPU全局内存延迟大约400到600个时钟周期的情况下需要多少个并发访存请求才能把显存带宽打满。这类题在面试中出现的频率非常高因为它考察的不是背参数的能力而是能不能用吞吐量视角看问题。基本计算模型是Little定律的硬件版本在飞行中的数据量 带宽 × 延迟。假设一张GPU的显存带宽是1TB/s访存延迟是500ns那么在任意时刻必须在内存子系统里“在飞行”的数据量是1TB/s × 500ns 500KB。也就是说如果活跃的并发访存数据量少于500KB访存流水线就喂不满带宽就一定有浪费。换算成请求数量就更直观。一个float4宽度为16字节500KB除以16字节约等于32000个访存请求同时在飞。GPU正是通过极高级别的线程并发来维持这个数量的。SM能容纳的线程和访存指令越多延迟隐藏能力越强。这也是为什么在选CUDA block大小时要考虑占用率occupancy一个block里开128个线程也许能跑但开512个线程往往更容易把访存流水线压满。这道题的工程意义在什么时候体现做算子优化时如果我发现某个kernel的算术强度太低比如纯访存型算子第一步就是去测它的实际带宽利用率。利用率上不去优先怀疑并发度是否足够再怀疑访存模式是否产生了过多的内存事务。这不只是笔试考点每次用Nsight Compute分析kernel时都在用。3. Amdahl定律与并行效率题理论加速比与工程现实的差距3.1 一道典型的Amdahl定律计算题笔试里出现了一道看起来非常“教科书”的题某个应用程序有10%的代码只能串行执行剩余90%可以完美并行。请求解它在64核处理器上能达到的理论加速比并讨论如果串行比例扩大到20%加速比会如何变化。套用公式加速比 S 1 / (串行比例 并行比例 / 核心数)。10%串行、64核时S 1 / (0.1 0.9/64) ≈ 8.5。也就是说即便有64个核心最理想也只能获得不到一个数量级的加速。如果串行比例变成20%S 1 / (0.2 0.8/64) ≈ 4.7加速比直接腰斩还多。这道题的标准答案不难但能拿高分的人会在最后补一个关键认知Amdahl定律是一个悲观模型它假设问题规模固定。如果问题规模随核心数增大而增大对应的加速比更适合用Gustafson定律描述。在实际AI场景里模型规模和数据集通常随着算力增强而扩大所以Gustafson定律往往比Amdahl定律更贴近真实。3.2 串行比例为什么是面试官精心埋的陷阱笔试里那个“10%串行比例”很容易让人觉得10%没什么大不了。但实际上恰恰是这部分主导了扩展性的上限。工程中真正的瓶颈也往往来自这些不起眼的串行片段。我在实际项目里遇到过多次类似情况。把一个数据预处理流程从单核改成多核最初以为只要把for循环改成OpenMP并行就行结果实测8核加速比只有2.4。原因就是每个线程完成后要往同一个日志缓冲区写行日志写入是有锁的随着线程数增多锁竞争越来越严重热点全部集中到串行区。这跟Amdahl定律的串行比例上升是同一个问题只是工程表现更加隐蔽。多核场景下串行区通常来自这几个地方内存分配器尤其是带全局锁的分配器、共享计数器和统计结构、日志与埋点、任务队列的入队出队操作、负载不均衡导致的收尾同步。面试中如果能在Amdahl定律题之后自己主动引出这些场景考官会觉得这个人是真正做过并行系统的人而不是只会背公式。3.3 负载不均衡与多核扩展性的实际案例笔试中还有一道扩展题问的是“在多核处理器上获得线性加速比为什么在现实中几乎不可能”要求列出至少三种限制因素。参考答案里通常包括串行代码段、同步开销、负载不均衡、缓存一致性和内存带宽竞争。每一条展开都能对应一个8核升16核收益不明显的真实原因。其中负载不均衡是我最想强调的。并行任务切分如果不均匀整体执行时间由最慢的那个线程决定。比如一个任务被分成16份其中一份的执行时间是其他的2倍那并行部分的有效利用率就损失在等待上这等价于凭空增加了串行比例。在做算子切分时我习惯对样本按计算量排序后再切块而不是简单按索引均分。AI推理场景里输入样本的序列长度往往差异巨大如果每个batch按固定数量分配有的线程处理长序列有的处理短序列整个kernel的执行时间就被长序列拖住。这时用动态任务窃取work stealing或者按token数分配batch才能让扩展性回到正轨。4. CUDA编程题从向量加法到Shared Memory与bank conflict4.1 向量加法的代码题陷阱第三批笔试有一道CUDA向量加法的编程题要求写一个完整的kernel把两个float数组相加。看起来是CUDA入门第一课但题目里故意加了一个要求说明block大小和grid大小的选取依据。这道题考察的不只是会不会写三行代码而是有没有真正理解线程配置和执行效率的关系。基础的kernel长这样__global__ void vectorAdd(const float* a, const float* b, float* c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { c[idx] a[idx] b[idx]; } }调用时block大小选多少这个问题学问很大。选太小的block比如32每个SM能容纳的block数量有限可能无法隐藏访存延迟。选太大的block比如1024又容易在block尾部出现大量线程因边界判断而空转。工程上常见的选择是128到512之间的值具体要结合SM的最大线程数和寄存器占用来确定。边界判断导致的分支发散divergence也值得玩味。虽然条件判断idx n几乎不会造成实际性能损失但一个严谨的kernel应该优先保证数组长度是block大小的整数倍。如果长度不能整除可以用grid-stride loop来减少尾部线程的浪费__global__ void vectorAddStride(const float* a, const float* b, float* c, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; int stride gridDim.x * blockDim.x; for (; idx n; idx stride) { c[idx] a[idx] b[idx]; } }vectorAdd这类入门题在实际工作中没什么用但它背后的线程配置、边界处理、grid-stride策略在写任何数据并行kernel时都用得上。笔试出这道题的目的就是看候选人有没有从“能跑”往“跑得快”走的意识。4.2 Shared Memory与bank conflict比数据结构还常考第三批笔试中最有区分度的题目之一是一道关于shared memory bank conflict的分析题。题目大致是假设有一个长度为N的float数组存放在共享内存中线程t需要访问数组的第2t个元素即以步长2访问请分析这会发生几路bank conflict并提出改进方案。共享内存的硬件结构是连续的32个bank每个bank的宽度是4字节。因此bank号通常按地址计算地址偏移除以4再对32取模。float类型恰好4字节所以第i个float落在bank (i % 32)上。线程t访问第2t个元素即float序号2tbank号为(2t) % 32。当t0时bank0t1时bank2t16时bank32%32bank0。这时线程0和线程16落在同一个bank上线程1和17落在bank2上以此类推。同一个bank在同一个周期内只能为一个线程服务所以要分两个周期来完成这批访问这就是2路bank conflict。这个知识点为什么在笔试里反复出现因为bank conflict直接决定了shared memory能否跑满带宽。很多CUDA新手把shared memory当成一个“更快的全局内存”以为把数据拷贝进去就自动变快完全没意识到它的内部访问规则。如果不同线程同时访问同一bank实际吞吐量会被显式折减。改进方案不止一个要看具体访问模式。对于步长2的访问最直接的手段是改变数组的存储布局在每行数据末尾填充一个floatpadding让相邻行的起始bank错开从而避免行间冲突。另一个通用思路是把对数组的访问改成按float4向量化让一个线程一次性读取4个float既减少访存指令数量又改变了bank分布规律。还有一类方案是使用共享内存的广播机制如果多个线程访问的是完全相同的地址硬件能在一个周期内完成广播这不算冲突。4.3 矩阵乘法的Tiling优化思路题矩阵乘法优化是异构计算笔试的“压轴常客”。第三批里出了一道没有要求写完整代码、只要求“描述如何利用shared memory优化矩阵乘法”的题这其实比写代码更难因为它要求候选人脑子里有一张完整的kernel执行图。基础版本是每个线程计算C矩阵的一个元素内层循环要读A的一行和B的一列对全局内存的访问量是O(N^3)性能受制于访存带宽。Tiling优化的核心思想是把大矩阵切块每次把A的一个tile和B的一个tile搬进shared memory在片内完成计算从而大幅减少对全局内存的重复访问。我在实际优化时会这样安排设tile大小为32x32或16x16每个block加载一个tile的A和一个tile的B到共享内存线程块内的256个线程分别负责tile内不同的输出元素。计算每个输出元素时遍历k维度上的一长条数据但每次内层循环的A和B数据都直接从shared memory读取全局内存只被精准访问两次。如果再加上double buffering把下一块tile的数据在计算当前tile时提前搬运访存和计算还能进一步重叠。这类题目最终考察的是对访存层次结构的理解和利用而不是某个具体优化trick。无论是CPU上的循环分块还是GPU上的shared memory tiling底层逻辑都是同一句话数据在寄存器里跑得最快在缓存和共享内存里次之在全局内存里最慢算法设计要尽量提升数据在被慢速存储访问前的复用次数。5. 深度学习推理优化题BatchNorm融合与1x1卷积的访存瓶颈5.1 BatchNorm在推理阶段如何融合到卷积层第三批笔试出现了三道和深度学习推理算子优化相关的题其中一道是说明BatchNorm在推理阶段如何合并进前一个卷积层并推导合并公式。BatchNorm在训练阶段依赖batch内统计量需要维护均值和方差还有可学习参数gamma和beta。推理阶段不再使用batch统计而是使用训练完毕后固定的全局均值和方差。一个重要事实是BatchNorm在推理时对每个通道执行的是线性变换所以它完全可以合并到相邻卷积层的权重和偏置里。卷积层输出是y W * x bBatchNorm的推理变换是z gamma * (y - mean) / sqrt(var eps) beta把y代入可以得到z (gamma * W / sqrt(var eps)) * x (gamma * (b - mean) / sqrt(var eps) beta)也就是说新的卷积权重是W乘以scale新的偏置是b乘以scale再平移。在推理引擎实现时只需要在模型加载阶段对每一层做一次参数改写就能完全消除BatchNorm在推理时的额外计算。这在网络层数较深时可以节省不少访存和计算开销。我从事推理引擎优化后才发现这类“参数折叠”的手段不止BatchNorm融合这一种。卷积后面的ReLU可以融合进上一个kernel的epilogue阶段LayerNorm的仿射变换可以折叠进前面的线性层INT8量化中的per-channel scale也可以提前乘到权重里。推理引擎优化的第一课永远是把数学上能合并的线性变换全部合并减少kernel数量而不是一上来就去魔改CUDA代码。5.2 1x1卷积为什么是访存密集型算子另一道题问1x1卷积的计算量不大但为什么在GPU上性能往往不理想请从访存特征角度分析。1x1卷积本质上是跨通道的线性组合对每个空间位置的各个通道做矩阵向量乘。它的计算量是N×H×W×Cin×Cout看起来不小但在现代GPU上真正的瓶颈往往是数据搬运。每个输入像素要被Cout个输出通道使用这意味着输入数据的复用率取决于Cout的大小如果Cout不大每个输入元素只被读几次就丢弃算术强度每次访存对应的浮点计算次数很低。低算术强度意味着kernel被访存瓶颈主导GPU的高算力完全使不上劲。优化这类算子的核心思路是提高数据复用比如让一个线程块处理多个空间位置把输入块在shared memory里缓存下来使多个输出通道共享同一份输入数据。另一个方向是融合相邻算子比如把1x1卷积后面的ReLU、BN、残差相加都放进同一个kernel的末尾处理避免中间结果写到全局内存再读出来。5.3 从题目到推理引擎的工程实践笔试题里有一道开放题让设计一个高效的卷积kernel。这道题没有唯一答案但高分答案通常包含这些层次第一明确卷积的shape和硬件平台第二选择合适的数据布局NCHW、NHWC还是NC1HWC2第三考虑im2col加GEMM还是直接卷积或者Winograd第四说明tiling策略、向量化宽度和double buffering。Winograd是当年笔试的热门考点之一。它通过减小卷积乘法次数来降低计算量代价是增加变换阶段的访存和数值误差适合小卷积核3x3大通道数的场景。我做实际项目时对Winograd的态度是先写好标准GEMM版本和直接卷积版本作为baseline再根据profile数据判断值不值得上Winograd。因为Winograd的预处理需要额外的内存拷贝和格式转换如果数据布局没优化好节省的计算时间会被访存开销吃掉。这类题目对考生的真正要求不是背出一堆优化名词而是具备“选择一个优化手段之前先分析它在当前硬件和当前shape下的收益”的工程判断力。6. 嵌入式与边缘场景题有限内存下的模型部署方案6.1 当年那道让我印象深刻的系统设计题第三批笔试的压轴题是一道系统设计题假设要把一个图像分类模型部署到内存只有几十MB的嵌入式设备上可用的硬件只有一颗低功耗CPU加一个小的NPU加速器请给出完整的部署优化方案。这道题把之前所有考点知识串成了一条线。它没有限定用什么框架也没有限定模型结构考察的是候选人在真实约束条件下做技术决策的能力。我当时的答题思路是先分类讨论约束内存受限、算力受限、功耗受限、模型结构未知。然后按层次给出方案。第一层是模型层面的优化。优先检查模型结构是否过大考虑知识蒸馏把大模型换成小模型或者使用神经架构搜索出来的轻量网络。这一层的收益是数量级的比任何算子优化都明显。第二层是权重压缩。对模型做INT8量化可以把模型体积和带宽需求压缩到原来的四分之一。如果精度损失在可接受范围内甚至可以尝试INT4量化配合混合精度。剪枝则把不重要的通道和连接删掉进一步减少计算量和存储。值得注意的是量化不只是把权重从FP32换成INT8还需要校准激活值的分布范围以及处理量化敏感层比如检测头的输出层。第三层是运行时优化。算子融合减少中间张量的内存占用内存池复用避免动态分配和碎片计算图优化把可以并行执行的节点调度到CPU和NPU上同时运行。嵌入式设备上的内存池设计与服务器端完全不同服务器随便malloc没问题嵌入式上频繁分配释放容易出现碎片必须从一开始就规划好各个张量的生命周期。第四层是硬件适配。NPU对算子的支持往往有限不是所有层都能直接跑在NPU上。有些层需要切成CPU上的子图执行。这时候CPU和NPU之间的数据拷贝可能是最大的性能杀手所以要尽量减少CPU和NPU之间的转换次数尽量把大块数据一次性切到NPU侧处理。6.2 这类设计题的答题方法论这道题能拿高分的核心不是方案多全而是是否先讲清楚约束和优化空间。我在实际面试别人时也喜欢出这类开放题最怕听到的回答是一上来就背量化剪枝名词完全不管设备的内存上限和算力上限。正确的方法是先画一条技术链路模型结构简化、参数压缩、运行时优化、硬件映射。每往前走一步都要附带代价分析。比如量化会带来精度损失剪枝会改变模型结构导致NPU算子映射更复杂算子融合可能增加寄存器压力。方案不是越高级越好而是在约束下收益最大、代价可接受的那个。笔试中的这类题本质上是看候选人有没有系统思维。我曾经把一个模型从300MB压缩到20MB并在嵌入式设备上跑通靠的恰恰是每一层都做一点优化50MB来自INT8量化100MB来自通道剪枝剩余的多余部分通过修改数据布局和内存复用解决。没有哪一种技术是银弹组合起来才有意义。6.3 从2019到今天的边缘异构趋势回想2019年嵌入式AI部署还集中在手机和安防设备上MCU级别的机器学习更多是跑跑裸机上的轻量识别。这几年边缘场景对异构计算的需求变化也非常明显尤其是工业控制和实时通信场景MCU加NPU、CPU加实时内核的架构已经成为趋势这类架构强调的正是异构计算中的“异构”二字不同的计算单元执行不同类型的任务MCU负责精确到微秒的实时控制和协议解析NPU/加速器负责AI推理两边通过共享内存或高速总线通信。这和笔试里的GPU优化题目其实是同一个世界观异构计算不是把计算单元堆在一起而是让每个计算单元做自己最适合的事情并管理好数据在单元之间的流动。在嵌入式设备上这种设计还要额外考虑功耗和确定性问题难度比纯GPU场景更高。但核心的分析方法没有变先算清楚瓶颈在哪再决定哪些计算放哪个单元最后为数据通路做优化。7. 备考这批题给我留下的东西给后来者的三点建议7.1 与其刷题不如亲手写一遍CUDA并验证想法回头看这套笔试题里最难准备的不是某个具体知识点而是“对性能有没有直觉”。这种直觉只能靠亲手写kernel、跑profile工具、对比不同实现的数据来建立。准备这类笔试时比起刷几百道LeetCode不如把矩阵乘法从naive版一步步优化到tiled版本再用Nsight Compute看一眼实际带宽利用率。每做一个阶段记录一次性能数据分析一次变化原因。这个过程做上两三个算子对访存层次、bank conflict、occupancy这些概念的理解会远比看书深刻。7.2 把理论题和工程案例绑定记忆这套笔试题有一个特点理论题都能在真实工程里找到对应场景。Amdahl定律对应多核任务扩展性bank conflict对应算子访存效率BatchNorm融合对应推理引擎的图优化。备考时如果把每个知识点都绑定一个实际工程案例来记不仅记得牢面试时还能答出别人答不出的“为什么”。比如我后来做BERT推理优化时发现Self-Attention的计算里Q、K、V矩阵乘法和后面的softmax、注意力池化之间的数据流天然就适合算子融合因为中间结果的体积巨大但计算又非常简单。如果能从推理引擎的角度理解这套笔试题这一章前文的很多题目其实都找到了落地的影子。7.3 数值估算能力是隐藏加分项这批卷子中好几道题都在考察数量级估算能力比如访存延迟需要多少并发来掩盖、嵌入式设备能跑多大模型。这种能力只能在平时刻意练习看到任何数据都随手做一次粗略计算一个FP32模型多少权重带宽需要多少显存够不够估算完再和实测对比长期下来对硬件性能的直觉会敏锐很多。最后再分享一个很实用的经验准备异构计算方向笔试和面试尽量在简历里附上自己的算子优化记录——原来性能多少优化后多少中间遇到什么问题怎么解决的。面试官问技术细节时这些材料远比空谈理论更有说服力。这套笔试题虽然已经是好几年前的但今天再看它考察的那些底层能力和思维方式在AI基础设施方向依然完全适用。