百度2018校招AI异构计算工程师的笔试题放到今天来看依然很有嚼头。当年异构计算还没有现在这么热招这个岗位的厂商也不多百度算是布局比较早的一批。第三批题目整体风格偏工程实践不纸上谈兵很看重候选人对“从算法到硬件落地”这条链路有没有真实体感。我结合当年笔试后的复盘把这份题涉及的核心知识模块拆开揉碎讲一讲顺便把当时踩过的坑和后来面试时被追问到的细节一并补上希望能帮到正在准备AI基础设施、高性能计算、算子优化方向校招的同学。1. 这套题考的不是深度学习而是“深度学习之外的计算功底”先说结论百度AI异构计算工程师的笔试题和普通算法工程师的题完全是两个物种。普通算法岗考你模型理解、调参技巧、loss设计异构计算岗考的是你“能不能把算法高效跑起来”。第三批题目里几乎没有让你推导Transformer attention公式的题反而大量出现并行计算、访存优化、卷积计算量估算、缓存命中率分析这类偏底层的内容。我当时拿到卷子第一反应是这不像是招算法工程师更像招一个“懂算法的系统工程师”。后来入职后想明白了异构计算工程师的核心工作就是把算法工程师写的Python代码变成能在GPU/NPU上高效运行的算子库或者推理引擎。这件事的难点不在算法本身而在硬件边界。你得知道显存带宽多少、计算单元有多少个、访存和计算的比率是多少、哪些操作适合并行哪些适合串行。笔试题目对应的能力模型可以拆成四块并行计算基础线程模型、归约、扫描、访存合并、负载均衡。AI算法底子CNN卷积的计算过程、FLOPs估算、模型各层的计算访存比。体系结构认知CPU/GPU/NPU架构差异、缓存层级、数据传输开销。C/C工程能力内存管理、指针语义、编译链接、性能优化手段。这四块大概按 3:3:3:1 的比例出题。别小看那10%的C/C题往往是区分度最高的部分。因为算法基础可以靠刷题短时间突击但C功底、性能敏锐度这种东西没有真实项目经验很难装出来。2. 卷积的“手推”细节FLOPs计算、访存开销与im2col实现思路2.1 一个卷积层的计算量到底怎么算笔试中有一道让我印象深刻的题给定输入特征图 H x W x C卷积核 3 x 3输出通道数为 Nstride1padding1要求写出这个卷积层一次的FLOPs和参数量。这道题坑不在公式在单位。很多人上来就写 FLOPs H x W x N x C x 3 x 3只算了乘法忽略了加法。严格来说FLOPsFloating Point Operations应该把乘加都算上。因为卷积本质是乘累加操作一次乘加记作两次浮点运算。更准确的写法是乘法次数H x W x N x C x 3 x 3加法次数H x W x N x (C x 3 x 3 - 1)也就是把 C x 3 x 3 个乘积累加需要 C x 3 x 3 - 1 次加法总FLOPs H x W x N x (2 x C x 3 x 3 - 1)如果忽略偏置一般工程上习惯直接写成 2 x H x W x N x C x K x K因为当 C x K x K 远大于1时那个减1可以忽略不计。但如果笔试填空题要求精确值你就得把加法和减法分开写。我当时就是忽略了加法次数直接在前面乘了个2就算了。后面跟面试官复盘时他说其实2 x 乘加算的是MACs乘以2这里表达容易让人误解他们真正想看的是你有没有把“乘累加”这个最基本的运算单元想清楚。2.2 im2col是把卷积转矩阵乘的核心手段还有一道im2col相关的题给出输入张量和卷积核形状要求写出im2col之后矩阵的维度以及这种转换的内存开销。简单来说im2col是把卷积操作重构成大矩阵乘法。输入特征图每个和卷积核关联的窗口区域会被拉成一列或一行然后和卷积核展开的矩阵做GEMM。假设输入是 H x W x C卷积核是 K x K输出通道 Nstride1paddingp。那么输出特征图尺寸是 H_out x W_out。每个输出位置对应的输入窗口大小是 C x K x Kim2col后矩阵的列数就是输入窗口的个数即 H_out x W_out行数是 C x K x K。对比GEMM的维度。GEMM的另一个矩阵就是卷积核本身展开的矩阵维度是 N x (C x K x K)。所以这本质上是一个 (CKK) x (H_outW_out) 的矩阵乘以一个 N x (CK*K) 的矩阵。这里隐藏的一个关键问题是内存开销。im2col会引入大量的数据冗余因为相邻窗口之间有重叠区域这些重叠区域的像素会被复制多份。当输入通道多、卷积核大的时候冗余特别严重。一个典型的例子输入 224x224x3卷积核 11x11AlexNet第一个卷积层stride4padding2。im2col之后矩阵的规模膨胀得非常大内存占用可能比输入本身大几十倍。笔试中有一问就是让你计算这种情况下im2col的内存占用和直接卷积的内存占用差异。这题如果没做过底层算子开发很容易没有体感。这也是为什么很多推理引擎在高性能实现里不会直接做完整im2col而是用隐式im2colimplicit gemm并不真正物理展开数据而是通过索引计算在GEMM时直接访存原始输入避免冗余拷贝。2.3 考这道题的真实意图我当时其实不理解为什么笔试要考这么细的卷积实现细节。后来做算子优化才明白卷积是CNN推理里最核心的计算原语对卷积的理解深度直接决定了你能不能写出高性能实现。举个实际数据在1080Ti上一个ResNet50模型如果卷积不使用im2colGEMM或者Winograd等优化方式而是一层层朴素循环实现性能会差几十倍。你刷多少道LeetCode都不如把一个卷积优化到接近cuDNN的水平更能打动面试官。百度的笔试题不靠偏题怪题它靠的是“高频知识点的深挖”这也是一个很有效的筛选信号。3. GPU并行计算的典型题目归约求和、存储体冲突与warp发散3.1 并行归约一道被翻来覆去问的题笔试中出现了一道非常经典的并行归约题给一个长度为N的数组要求在GPU上并行计算所有元素的和写出kernel的伪代码并分析时间复杂度、访存模式以及存在的bank conflict。这道题看起来简单但答得完整的人很少。很多人上来就写一个split策略比如每个线程负责一部分数据做局部归约然后用原子操作把所有部分和加起来。这种写法能跑通但不是最优解。更高效的实现通常是树形归约。每个block加载一块数据到shared memory然后反复执行“相邻元素两两相加结果存回前半段”的操作直到只剩一个值。每次迭代的活跃线程数减半。这里的关键考点在于活跃线程数递减到一定程度以后会出现严重的warp发散比如只剩16个线程在干活但一个warp是32个线程那就有一半线程空闲。bank conflict的记忆如果两个线程访问shared memory的同一个bank不同地址硬件会把这些访问串行化导致性能下降。我当时写的是标准版tree reduction就是相邻归约。面试时被追问为什么不用stride 32的方式我回答后他才说你看看如果不用padding常见的写法在最后一轮会不会出现bank conflict。这才发现shared memory数组如果只按线程数来分配最后一轮的时候线程0访问address[0]线程1访问address[1]它们恰好在不同的bank不会冲突但如果往前看当32个线程同时在访存的时候线程0和线程16如果访问的是address[0]和address[16]而一个bank有32个存储单元address[0]和address[16]刚好落在不同的bank所以没问题。真正容易冲突的是线程i访问address[i]和address[i16]这种操作模式取决于具体写法。这个细节如果没有真正在CUDA上写过kernel、分析过profiler输出很难一次答对。3.2 一个可复用的标准归约kernel分享一个相对标准、面试时可以直接默写的版本基于CUDA它做了几层优化__global__ void reduce_sum_kernel(const float* input, float* output, int n) { __shared__ float sdata[256]; int tid threadIdx.x; int global_id blockIdx.x * blockDim.x tid; float sum 0.0f; // 每个线程做连续stride的访存保证合并访问 for (int i global_id; i n; i gridDim.x * blockDim.x) { sum input[i]; } sdata[tid] sum; __syncthreads(); // 树形归约 for (int stride blockDim.x / 2; stride 0; stride 1) { if (tid stride) { sdata[tid] sdata[tid stride]; } __syncthreads(); } if (tid 0) { output[blockIdx.x] sdata[0]; } }这个版本不够极致但足够展示基本功。有几个点可以在面试时主动讲解第一网格跨步循环grid-stride loop为什么不用global_id blockIdx.x * blockDim.x tid然后只处理一个元素因为数组长度可能比总线程数大很多一次性加载完会导致很多线程空转而且无法复用同一份数据。网格跨步循环能让每个线程连续处理多个元素而且相邻线程访问相邻内存访存是合并的。第二为什么用shared memory做中间缓存因为全局内存的延迟是几百个周期shared memory只有几十个周期。把每个线程的部分和存到shared memory然后从shared memory里归约能大幅减少全局内存访问。第三__syncthreads()为什么要保证每次归约迭代都同步因为所有线程都要先完成当前迭代的读操作才能让下一次迭代看到最新值不同步就是数据竞争。3.3 从归约题延伸出的性能优化思路笔试之后有一轮面试专门问归约的性能分析。让我给出一份简单的优化列表后来我把实际工作中总结的东西整理了一下对齐与合并tail effect尾效应是并行归约性能杀手。数组长度不是线程数的整数倍时最后一段数据的访存很不规整容易产生大量部分波前。预先对数组做padding填充到对齐长度或者让每个线程处理多个元素都能缓解。避免bank conflictshared memory数组可以采用float sdata[blockDim.x 1]的方式申请故意错开一个元素让某些访问模式不会落在同一个bank上。减少同步次数现代GPU上32个线程内的shuffle指令可以直接交换寄存器数据不需要经过shared memory也不再需要__syncthreads()。用__shfl_down_sync做warp内归约可以把同步开销降到最低。向量化访存使用float4一次读4个float能减少访存指令的数量提升内存带宽利用率。这个在带宽受限的算子里效果特别明显。这些经验不是笔试时能写完的但是在面试时主动提出来会明显拉开和普通候选人的差距。我一直建议准备这个岗位的同学不要只停留在“能做对题”的层面要往“把题做深”去准备。4. 异构计算架构题CPU与GPU的职责切分以及数据搬移的隐藏成本4.1 CPU和GPU不是“谁快谁就多干活”笔试中有一道题目很典型在CPUGPU异构系统上部署一个实时视频推理服务输入视频流每秒30帧每帧需要做预处理、推理、后处理。要求设计一个合理的流水线架构并指出各阶段的负载分配和潜在性能瓶颈。这类题没有标准答案但考察的是你是否理解异构计算的核心矛盾——数据搬移开销。很多人容易犯的错误是既然GPU快那就把所有计算都丢给GPU。但实际上GPU的强项是吞吐量受限的大规模并行计算而不适合处理逻辑分支复杂、串行依赖重的操作。举例来说视频解码decode用GPU不一定比CPU的专用硬件解码器更快而且解码出来的数据通常是NV12格式的YUV数据还需要转换成RGB再缩放、归一化这些操作如果全部放在GPU上每一步都要把中间数据留在显存里最后才能做推理显存占用会变得很大。同时一旦某个中间环节要在CPU上做比如常见的预处理库某些算子只有CPU版本那就得来回拷贝PCIe带宽瞬间成为瓶颈。我当时给的方案是CPU负责视频拉流和解码GPU负责预处理推理后处理GPU内部的数据用CUDA Stream和Pinned Memory实现异步拷贝。后来面试官说方案方向没问题但有一个点被我忽略了Pinned Memory并不是万能的如果用默认的cudaMemcpy是同步的必须用cudaMemcpyAsync才有意义而且pinned memory的分配开销比较大最好使用内存池来管理。4.2 流水线设计里的关键重叠计算与拷贝这道题延伸出来的经典知识点是如何把拷贝和计算重叠起来。一个最朴素的推理循环可能是这样CPU读取图片。拷贝到GPUHost-to-Device, H2D。GPU推理。拷贝结果回CPUDevice-to-Host, D2H。这样每一步都是串行的GPU在等待拷入的时候是空闲的CPU在等待推理的时候也是空闲的。更好的做法是使用多个CUDA Stream把数据切分成多个批次让第i张图的拷贝和第i1张图的推理重叠形成软件流水线。具体拆解成三段缓冲triple buffering是工程上最稳的做法Buffer A当前正在CPU上准备的数据。Buffer B正在从Host拷贝到Device的数据。Buffer C正在GPU上推理的数据。每一个时间切片里三个buffer各忙各的互不干扰只需要在切换阶段做同步。这个方式在真实项目里能显著提升吞吐通常能压掉30%-50%的端到端延迟。4.3 理论带宽和实际带宽的差距架构题往往还会让你估算数据传输时间。我记得有一道小题是这样一张1080p的RGB图像大约是 1920x1080x3 ≈ 6.2 MB如果PCIe 3.0 x16的理论带宽是16 GB/s问理想情况下传输需要多少时间又提示实际有效带宽通常只有理论值的50%-60%重新估算。理论上 6.2 MB / 16 GB/s ≈ 0.39 ms。实际按60%算的话大概是0.65 ms。这个量级看似很小但如果你的模型推理一次只需1 ms那数据搬移就占了将近40%的时间绝对不可忽略。后来看到很多做AI服务优化的同学经常把大量精力花在模型算子优化上却忽视了数据链路优化这其实就是没有建立“全链路视角”的问题。笔试里特意考这个就是为了筛选出有系统观的人。4.4 NPU/FPGA这类非GPU异构设备的选型陷阱还有一个选型类题目给出几种AI加速芯片包括GPU、FPGA、ASICTPU/NPU让你基于功耗、灵活性、开发成本、性能密度几个维度做选型并说明理由。这道题没有标准答案核心是考察你是否理解体系结构设计上的经典权衡。我当时从三个角度看GPU灵活性和性能都很强生态最成熟适合通用深度学习训练和推理。但功耗高不适合功耗受限的边缘设备。FPGA硬件逻辑可重构能定制数据通路非常擅长某些低延迟、高吞吐的定点推理任务。但逻辑资源有限如果算子特别复杂往往做不过GPU而且开发用Verilog/HLS语言周期长。ASIC/NPU性能和能效比理论最优但不可编程或可编程性弱一旦算法结构变化硬件就废了。在实践中取舍的核心其实是你的部署规模有多大。大规模量产且算法确定NPU是长期竞争最优解任务类型多、迭代快GPU或者GPUCPU组合反而总成本更低。硬件选型不是选最强的而是选最匹配业务的。5. 理论篇之外笔试中的C/C性能题与内存布局细节5.1 struct内存对齐看似送分踩坑的人无数第三批笔试里考了C/C struct内存对齐题目给定struct A { char a; int b; char c; };问这个结构体的大小是多少在默认对齐规则下。答案是12字节不是6字节。char占1字节然后为了int的4字节对齐在a后面会填充3字节b占4字节c占1字节最后结构体整体对齐到最大成员对齐数4c后面再填充3字节总共12字节。这个知识点本身不难但在异构计算的场景里有实际意义。GPU kernel里经常要定义参数结构体传给device端如果结构体定义不合理每个线程读取的时候内存带宽利用率会下降。更严重的是如果结构体里含有指针且对齐方式不一致可能导致aligned access异常。C11之后可以用alignas和alignof来控制对齐也可以用#pragma pack改变默认对齐但在GPU编程里我强烈建议不要随便pack除非你知道自己在做什么。shared memory的访存模式、全局内存的合并访问都跟对齐关系很大。5.2 指针与引用、值语义与移动语义的辨析还有一道C题考右值引用和移动语义给了一段代码要求分析发生了多少次拷贝、多少次移动。这种题对AI异构计算工程师来说不是可有可无的。原因是数据量大的Tensor、Buffer、模型权重在PC上拷贝一下可能没什么感觉但在内存紧张的推理引擎里多一次拷贝可能就是一次明显的延迟。现代C工程里所有深度学习框架的C前端底层都在大量使用移动语义比如PyTorch的Tensor、Llama.cpp的ggml_tensor、TensorRT的ITensor。理解拷贝和移动的区别是理解这些框架源代码怎么省内存的第一步。我当时笔试答得比较顺因为平时写C比较多。后来在面试中我还被追问了std::shared_ptr和std::unique_ptr的选择以及为什么深度学习框架的算子调度器一般用裸指针而不是智能指针传递张量数据。这个问题的本质是智能指针管理的是对象生命周期而张量数据本身的分配/释放逻辑往往被自定义allocator控制生命周期和内存分配两件事被解耦了裸指针加上显式的内存管理反而是更可控的方案。这种从具体题目延伸到工程理念的思考方式是我能过一面二面很重要的加分项。5.3 一层容易被忽略的细节pimpl模式与编译期解耦有一道附加题提到如何降低大型C项目的编译时间要求写出你的方案。我突然意识到很多校招生对编译期这个概念很模糊他们能理解运行时性能但不理解编译期性能也是工程率的重要一环。大型AI推理引擎动辄上百万行代码每次改一个头文件要重新编译几千个文件好几十分钟就没了。一个常用的解法是pimpl idiom也叫Pointer to Implementation。原理是把类的成员变量放到一个内部结构体里通过一个前置声明的指针来引用它这样头文件里不再依赖具体类型的定义因此修改内部实现时不需要重新编译外部头文件依赖者。这跟异构计算有什么关系算子库开发时头文件里大量包含CUDA runtime的头文件如果在公共API里不小心暴露了CUDA类型所有引用你这个API的模块都要强制装CUDA环境。用pimpl把CUDA类型藏到.cpp里能大大降低依赖耦合。这种问题在笔试中不常见但一旦出现就是筛选出那些真正读过大型开源项目源码的人。6. 从笔试到面试备考AI异构计算岗的实用反思6.1 高频失分点复盘先复盘一下我见过的、以及自己在笔试中踩过的典型失分点方便后来人对照自测第一卷面上直接把FLOPs和参数量搞混。参数量是CKK*NFLOPs是计算量它还要乘上输出特征图分辨率。这两个概念在模型压缩、架构选型时经常一起出现一旦搞混后续性能分析全错。第二对“内存带宽”只有概念没有数字。很多人知道内存带宽很重要但不知道常见硬件的量级。DDR4内存大概几十GB/sA100的HBM2e大概2TB/sPCIe 3.0 x16大概16GB/sPCIe 4.0 x16大概32GB/s这些数字要背熟。答题时随手做量级估算能显得你非常有“工程感”。第三不会分析访存密集型算子。笔试里如果给一个逐元素加法的算子让你分析是计算密集型还是访存密集型很多人只会说“GPU并行度高所以快”不会用算术强度arithmetic intensity来量化。算术强度 总FLOPs / 总访存字节数。如果这个值远小于当前硬件的FLOPs/带宽比就说明算子是访存受限的此时一味堆计算单元没有用得优先减少访存量。6.2 一个访存受限算子的优化案例具体回顾一下这类题的实际解法。假设要对一个1亿个float的数组做逐元素乘2总数据量是4亿字节400MB总计算量是1亿次乘法算下来算术强度 0.25 FLOPs/Byte。A100的理论算力是19.5TFLOPS带宽是2TB/s比值大约9.75FLOPs/Byte。我们的算子算术强度0.25远低于9.75说明瓶颈在带宽上。无论怎么优化计算指令性能上限都受制于400MB数据的搬运时间。此时真正能提速的方向是减少数据访问次数比如融合多个elementwise操作、提升每次搬运的字节效率比如向量化加载float4、或者用更高速的缓存层级。这类算术强度分析是AI异构计算笔试里一个非常高频的分析工具。理解了这套量化方法很多题就能一眼看穿本质。6.3 一些备考建议和资源清单如果你正在准备AI异构计算方向的校招我觉得可以按下面这个顺序来复习效率比较高并行编程入门完成一个简单的CUDA程序比如并行归约、矩阵乘法。建议对照NVIDIA官方文档把线程模型、共享内存、同步、原子操作这几个概念吃透。性能分析工具学会看Nsight Compute或者NVIDIA Nsight Systems报告。笔试不会让你现场用工具但面试官可能会问“你做过的算子优化是怎么衡量性能提升的”如果你能说出具体指标、实测数据可信度会高很多。体系结构课补课重点看缓存结构、虚拟内存、DMA传输、总线带宽。推荐《Computer Organization and Design》相关章节或者CSAPP深入了解计算机系统里的性能优化部分。框架源码阅读选一个深度学习推理引擎精读一个算子的实现比如OneFlow、TensorRT的开源版本、或者llama.cpp的GGML实现。重点观察它们是怎么管理内存的、怎么把算子拆成多个kernel的、怎么利用汇编级指令的。6.4 我对这套笔试题的总体印象回头再看百度2018校招AI异构计算工程师第三批笔试我的评价是题目覆盖面广难度梯度合理每一道题都能看出出题人希望招到什么样的人——不是背过几篇论文摘要的“名词党”而是真正动手写过CUDA kernel、分析过算子性能瓶颈、对系统全链路有认知的工程师。现在AI基础设施岗位的面试风格变化不少很多公司开始现场让候选人写一个简单算子的高性能实现然后问各种极端情况下的性能问题。所以如果这套题你复盘得很透彻那些新题型大概率也能触类旁通。最后分享一个我后来带实习生时经常说的观点笔试里的这些题目不是考完就没用了。归约、卷积、访存优化、流水线重叠、内存对齐你在真实项目里几乎每天都会遇到。把笔试当成一次体系化学习的机会而不是一个必须过关的考试收获会大得多。