企业官网建设流程全解析

百度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，输出通道数为 N，stride=1，padding=1，要求写出这个卷积层一次的FLOPs和参数量。

这道题坑不在公式，在单位。很多人上来就写 FLOPs = H x W x N x C x 3 x 3，只算了乘法，忽略了加法。严格来说FLOPs（Floating 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，输出通道 N，stride=1，padding=p。那么输出特征图尺寸是 H_out x W_out。每个输出位置对应的输入窗口大小是 C x K x K，im2col后矩阵的列数就是输入窗口的个数，即 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，卷积核 11x11（AlexNet第一个卷积层），stride=4，padding=2。im2col之后，矩阵的规模膨胀得非常大，内存占用可能比输入本身大几十倍。笔试中有一问就是让你计算这种情况下im2col的内存占用和直接卷积的内存占用差异。这题如果没做过底层算子开发，很容易没有体感。

这也是为什么很多推理引擎在高性能实现里不会直接做完整im2col，而是用隐式im2col（implicit gemm），并不真正物理展开数据，而是通过索引计算在GEMM时直接访存原始输入，避免冗余拷贝。

2.3 考这道题的真实意图

我当时其实不理解为什么笔试要考这么细的卷积实现细节。后来做算子优化才明白，卷积是CNN推理里最核心的计算原语，对卷积的理解深度直接决定了你能不能写出高性能实现。

举个实际数据：在1080Ti上，一个ResNet50模型，如果卷积不使用im2col+GEMM或者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[i+16]这种操作模式，取决于具体写法。

这个细节如果没有真正在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 conflict：shared 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不是“谁快谁就多干活”

笔试中有一道题目很典型：在CPU+GPU异构系统上部署一个实时视频推理服务，输入视频流每秒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 流水线设计里的关键：重叠计算与拷贝

这道题延伸出来的经典知识点是：如何把拷贝和计算重叠起来。

一个最朴素的推理循环可能是这样：

1. CPU读取图片。
2. 拷贝到GPU（Host-to-Device, H2D）。
3. GPU推理。
4. 拷贝结果回CPU（Device-to-Host, D2H）。

这样每一步都是串行的，GPU在等待拷入的时候是空闲的，CPU在等待推理的时候也是空闲的。更好的做法是使用多个CUDA Stream，把数据切分成多个批次，让第i张图的拷贝和第i+1张图的推理重叠，形成软件流水线。

具体拆解成三段缓冲（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、ASIC（TPU/NPU），让你基于功耗、灵活性、开发成本、性能密度几个维度做选型，并说明理由。

这道题没有标准答案，核心是考察你是否理解体系结构设计上的经典权衡。我当时从三个角度看：

• GPU：灵活性和性能都很强，生态最成熟，适合通用深度学习训练和推理。但功耗高，不适合功耗受限的边缘设备。
• FPGA：硬件逻辑可重构，能定制数据通路，非常擅长某些低延迟、高吞吐的定点推理任务。但逻辑资源有限，如果算子特别复杂，往往做不过GPU；而且开发用Verilog/HLS语言，周期长。
• ASIC/NPU：性能和能效比理论最优，但不可编程或可编程性弱，一旦算法结构变化，硬件就废了。

在实践中取舍的核心其实是：你的部署规模有多大。大规模量产且算法确定，NPU是长期竞争最优解；任务类型多、迭代快，GPU或者GPU+CPU组合反而总成本更低。硬件选型不是选最强的，而是选最匹配业务的。

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字节，最后结构体整体对齐到最大成员对齐数（4），c后面再填充3字节，总共12字节。

这个知识点本身不难，但在异构计算的场景里有实际意义。GPU kernel里经常要定义参数结构体传给device端，如果结构体定义不合理，每个线程读取的时候内存带宽利用率会下降。更严重的是，如果结构体里含有指针且对齐方式不一致，可能导致aligned access异常。

C++11之后可以用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*N，FLOPs是计算量，它还要乘上输出特征图分辨率。这两个概念在模型压缩、架构选型时经常一起出现，一旦搞混，后续性能分析全错。

第二，对“内存带宽”只有概念没有数字。很多人知道内存带宽很重要，但不知道常见硬件的量级。DDR4内存大概几十GB/s，A100的HBM2e大概2TB/s，PCIe 3.0 x16大概16GB/s，PCIe 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基础设施岗位的面试风格变化不少，很多公司开始现场让候选人写一个简单算子的高性能实现，然后问各种极端情况下的性能问题。所以如果这套题你复盘得很透彻，那些新题型大概率也能触类旁通。

最后分享一个我后来带实习生时经常说的观点：笔试里的这些题目，不是考完就没用了。归约、卷积、访存优化、流水线重叠、内存对齐，你在真实项目里几乎每天都会遇到。把笔试当成一次体系化学习的机会，而不是一个必须过关的考试，收获会大得多。