企业官网建设流程全解析

我参加过2018年百度校招AI异构计算工程师的笔试，那批题让我印象很深。说实话，刚点开卷子的时候，还以为是选几道CUDA语法题，结果从头做到尾才发现，整套题根本不是考“会不会写kernel”，而是考“你懂不懂异构计算到底在解决什么问题”。那批题里覆盖了GPU架构、访存优化、算子融合、多卡通信，甚至还有一道关于工业场景里异构芯片选型的论述题，完全是在筛一个人有没有真正理解“CPU负责调度、加速器负责算力”这件事。

这批题对后来想投AI基础设施、模型部署、高性能计算相关岗位的同学很有参考价值。不需要你是个CUDA专家，但你需要有清晰的体系感。我下面按自己的记忆和复盘，把这套题的出题逻辑、核心考点、解题思路和一些容易踩的坑讲透。你如果正在准备异构计算、AI引擎优化这类岗位，这篇内容可以直接拿来当复习提纲。

1. 这张卷子背后，藏的其实是百度AI的“算力焦虑”

1.1 为什么百度在2018年就专门设“异构计算”岗

2018年这个时间点很关键。那时候深度学习框架已经普及，模型越做越深，显存越吃越紧。百度内部有大量搜索、信息流、语音、图像模型要跑在GPU集群上，单纯靠调用cuDNN和TensorFlow自带算子，性能天花板很快就撞到了。同一个卷积算子，别人实现跑1.2毫秒，你自己实现只能跑3.8毫秒，整条链路掐下来，线上服务的吞吐就差了好几倍。

所以百度需要的不再是“会用TensorFlow调用GPU的人”，而是“能写出比TensorFlow默认实现更快的自定义算子的人”，这就是“AI异构计算工程师”这个岗位出现的原因。这套笔试题的所有内容，都是在为这个目标选人。

1.2 笔试题里考的不是API，而是“算子的肌肉记忆”

我当时拿到卷子后的第一印象是：怎么全是“场景题”。比如它问你“当你的kernel在GPU上实际运行速度远低于理论峰值时，你第一反应检查哪几个环节”，这种题没有任何标准API可以背，考的是你对访存、计算、调度是否有完整的认知框架。

整套题可以拆成三大块：硬件架构基础、编程题/优化题、系统链路与通信题。我后面会把这三大块分别展开，并把每道题背后真正想验证的能力点说出来。

2. 硬件基础题型拆解：不懂GPU内存层次，写十行代码错八处

2.1 线程结构与任务映射的错误示范

这批题里有一道非常经典的选择题：给你一段CUDA代码，blockDim是(256, 1, 1)，gridDim是(128, 1, 1)，让你算一共有多少个线程、每个线程的全局索引是多少。很多人栽在这道题上，不是因为不会算，是没理解blockIdx和threadIdx的换算关系。

正确公式很简单：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

线程总数就是gridDim.x * blockDim.x，也就是 128 * 256 = 32768。

但笔试题不会只问这个。它会变着法子考：如果数据量是100000，而线程总数只有32768，你怎么办？这时候就需要所谓的“grid-stride loop”：

for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += gridDim.x * blockDim.x) { // 处理第i个元素 }

这个循环的好处是：不管数据量多大，都能用固定数量的线程块把整个数组遍历完，而且线程数可以按设备算力来调，避免线程开太多导致调度开销变大。

注意：笔试里的那道题并不是直接让你写grid-stride loop，而是给了一段错误的代码，让考生判断“为什么运行结果错误”。错误点在于没有判断边界，线程越界访问了未初始化内存。这个细节提醒我们，写kernel时永远要把“边界检查”放在第一位。

2.2 内存体系是送分题也是送命题

CPU侧的程序员刚接触CUDA时，最容易忽略的一点是：GPU里的“内存”不是一整块，而是分层次的。GPU里有全局内存、共享内存、寄存器、L1/L2缓存、常量内存、纹理内存。每层内存的延迟和带宽差出两个数量级。

笔试里有一道填空题直接给了几个数字，让你判断分别对应哪类内存。我回忆一下大概的数字量级：

我当时在卷子上写答案的时候，漏写了L1缓存，只写了L2。后来复盘才意识到，L1缓存在很多GPU架构里和共享内存共享同一个物理存储，这正是一道考察“你是否知道共享内存其实可以是手动管理的L1”的题。

笔试里还考了一个经典陷阱：“为什么GPU上数组的访问要尽量做到连续”。这背后的原理是全局内存的访存是按“事务”来做的，一个事务可以一次取32个连续字节。如果相邻线程访问相邻地址，32个线程的访问可以被合并成一个或少数几个事务，访存效率最高。如果相邻线程访问的是跨度很大的地址，比如每次只取1个字节但间隔很远，那就要发很多次事务，性能直接崩。现场面试官后来说，这个题很多人能答出“合并访问”四个字，但说不清为什么要合并，就是这个原因。

3. 核心编程题实战：从零实现一个可落地的elementwise算子

3.1 题目还原与常规解法

我记得编程题里有一道是“实现一个向量加法”。看起来很简单，给定两个长度为N的float数组a和b，把结果写到c里。常规的C++写法三行就结束了，但CUDA版本要写一堆配置代码。

第一版解法通常长这样：

#include <cuda_runtime.h> __global__ void vector_add(const float* a, const float* b, float* c, int n) { int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < n) { c[tid] = a[tid] + b[tid]; } } int main() { int n = 1 << 20; size_t bytes = n * sizeof(float); float *h_a, *h_b, *h_c; h_a = (float*)malloc(bytes); h_b = (float*)malloc(bytes); h_c = (float*)malloc(bytes); // 初始化数据... float *d_a, *d_b, *d_c; cudaMalloc(&d_a, bytes); cudaMalloc(&d_b, bytes); cudaMalloc(&d_c, bytes); cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); int threads = 256; int blocks = (n + threads - 1) / threads; vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n); cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); free(h_a); free(h_b); free(h_c); return 0; }

这里有一个容易忽略的点：(n + threads - 1) / threads是向上取整的分块方式，保证block数量总是能覆盖所有数据。很多人在笔试手写代码时，会忘了这一步，直接写n / threads，结果N除以256不整除的时候，最后的元素就没人处理了。

另外一个高频陷阱是：代码里虽然malloc了host内存，也初始化了，但在实际工程里，如果数据量很大，一次cudaMemcpy会阻塞很长时间。笔试不考这个，但面试官会追问：“如果a和b是在GPU上已经算好的结果，你还需要拷贝回host吗？”答案是不需要，直接kernel内就能读取。这个追问考察的是“你是否理解数据应该尽量留在设备侧”。

3.2 进阶优化：向量化、网格步长与算子融合

笔试的加分项，大概率在第二问或者选做题里出现。它不会只让你写一个正确版本，还会问“你能做哪些优化”。我复盘下来的标答有三点：向量化访存、网格步长循环、算子融合。

向量化访存的意思是，用CUDA内置的float4类型一次读4个float，这样每个线程一次可以处理4个元素，减少指令总数和访存请求数。改写后的kernel大概是：

__global__ void vector_add_vec4(const float4* a, const float4* b, float4* c, int n) { int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < n / 4) { float4 va = a[tid]; float4 vb = b[tid]; c[tid].x = va.x + vb.x; c[tid].y = va.y + vb.y; c[tid].z = va.z + vb.z; c[tid].w = va.w + vb.w; } }

调用时，n要传实际元素数除以4，并且需要保证原始数组长度为4的倍数，或者处理剩余尾部元素。

网格步长循环我刚才已经给过代码，它的作用是让线程数量不受数据量的限制，同时让每个线程的工作量更均衡。还有一个隐藏好处：如果同一个kernel后面还要继续对结果做别的操作，网格步长循环可以让你更容易控制L2缓存的复用。笔试里你只要能写出来，再简单说一句“可以配合Persistent Threads技术，让线程块常驻，减少block调度开销”，这道题就稳了。

算子融合是这批题里的一个重头戏。它不会直接让你融合两个算子，而是问：“假设你已经写好了向量加法的kernel，现在还要做逐元素乘法和ReLU，你会怎么写？”如果分开写三个kernel，就要读三遍a、b和写三遍c。如果融合成一个kernel，读一遍就能把d = relu(a * b + c)算完。这本质上是用“减少全局内存访问”来换取性能。

笔试里有一道题的解法我到现在还记得：它问的是“为什么很多推理引擎会把卷积层和后面跟着的偏置加和ReLU融合成一个算子”。答案是：卷积的输出如果直接写到全局内存，再读出来做偏置加和激活，中间多了一轮访存。而GPU上全局内存的带宽是稀缺资源，一次额外的全量读写可能吃掉20%-30%的性能。融合之后，中间结果直接留在寄存器或共享内存里，连L2都不用写，省掉的时间肉眼可见。

我当时在答题时补充了一个小点：融合算子还能减少kernel launch的次数。每次启动kernel至少有5-10微秒的固定开销，如果模型有100个可融合的算子，这就是0.5-1毫秒的纯开销。这个角度面试官很喜欢，因为说明你是从系统层面思考的，而不仅仅是把算子写对。

4. 多卡协同与通信开销：校招题里最容易被低估的分

4.1 通信拓扑与带宽计算

2018年那批笔试里，多卡通信的题不算多，但是分值很高。有一道论述题大概是：“在8卡GPU服务器上训练一个模型，loss收敛速度上不去，你觉得瓶颈在哪”。很多人第一反应是“模型并行化没做好”，其实这道题想让你答的是“通信开销可能已经超过了计算收益”。

在单机多卡场景下，卡间通信走PCIe或NVLink，拓扑不同，带宽差异巨大。NVLink单链路的双向带宽在2018年那代卡上大约是25GB/s左右，PCIe 3.0 x16的单向带宽大约16GB/s。如果通信方式是“所有卡都要和卡0通信”，那卡0就是瓶颈；如果是环形通信，每张卡只需要和邻居通信，带宽压力就能分散。

笔试里还出现过一个计算题：假设每张卡每次迭代要发送128MB的梯度数据，8卡之间需要做一次全规约AllReduce，通信带宽是6.25GB/s，问通信时间是多少。如果你按“每张卡都传给卡0”来算，那通信数据量就是7128MB，再除以带宽，结果大约要143毫秒。如果按环形AllReduce优化，每张卡只需要传2(n-1)/n*数据量，总时间就大幅下降。

这个计算过程很有价值。面试官想看到的不是死记硬背公式，而是“你能把通信模型抽象出来，算清楚通信和计算的占比”。因为在实际AI训练中，如果单次迭代的通信耗时超过计算耗时的20%，整个集群的利用率就会很难看，这时候就要考虑梯度压缩、梯度累积，甚至是换更好的通信拓扑。

4.2 一道送命题：如何优化AllReduce

AllReduce是分布式训练里的最核心通信原语。笔试里问它的优化方式，算是既基础又进阶的题。从朴素版到高性能版，大致有这几个级别：

第一级：不做任何优化，所有卡把梯度发给卡0，卡0规约完再广播回去。优点是实现简单，缺点是卡0带宽压力巨大，而且链路利用率极低。

第二级：采用环形AllReduce。卡和卡之间组成一个环，每张卡把自己的一部分数据发给下一张卡，同时接收上一张卡的数据，做局部规约。数据在环上转一圈之后，再把完整结果转一圈分发回去。这样每张卡的发送量和接收量是均等的，没有热点。

第三级：分层AllReduce。在8卡服务器里，同一块CPU下的4卡先做一次局部规约，减少跨CPU的数据量，再在CPU之间做一次全局规约。这种拓扑感知的优化在真实集群里非常常见。

笔试里还有一道延伸题：“为什么梯度通信前的量化压缩可以提升训练性能”。原因很简单，通信时间 = 数据量 / 带宽，如果把梯度从FP32压缩成FP16或者INT8，数据量减半甚至减到1/4，通信时间也近似按比例缩短。但代价是精度损失，所以需要做误差反馈或混合训练技巧。这个题当年我答得比较浅，只回答了“减少数据量能加速”，面试官后来提醒我，“压缩本身也有计算开销，需要判断阈值条件”。所以准备这类题，不能止步于一个答案。

5. 从笔试到实战：我踩过的坑和自查清单

5.1 当年考场上最容易翻车的三类错误

第一类，是分不清“host”和“device”的指针。笔试手写代码是纯文本环境，没有编译器的帮助，经常有人在cudaMemcpy里把host和device参数的顺序搞反。正确顺序是cudaMemcpy(dst, src, count, kind)，第一个参数是目标，第二个参数是源。这个细节不记清楚，笔试里就相当于裸奔。

第二类，是忘记错误检查。真正的CUDA程序里，任何一个CUDA函数都可能返回错误码，比如cudaMalloc失败、cudaMemcpy超时、kernel启动非法配置。正规的工程代码里应该这样调用：

cudaError_t err = cudaMemcpy(d_c, h_c, bytes, cudaMemcpyDeviceToHost); if (err != cudaSuccess) { fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err)); exit(EXIT_FAILURE); }

笔试里虽然不需要你写全套错误处理，但心里要有这个意识，面试官抽查时能答上来，会加分。

第三类，是没有考虑“数据量不是线程块整数倍”的边界情况。我在前面的向量加法里已经强调过，向上取整后面一定接一个if判断，否则多出来的那些块会越界访问。这个问题在实战线上也容易出现，是我自己在训练框架里排查性能Bug时经常第一眼看到的问题。

5.2 给准备异构计算岗的读者一份备考路线

如果你正在准备类似岗位，我建议按照下面的顺序复习，不要一上来就刷LeetCode式的题目。笔试里更看重的是“你对系统每个环节的损耗有多敏感”。

第一步，熟悉CUDA编程模型。至少手写过3个算子：向量加法、矩阵乘法、规约求和。写的时候不要只看功能，看访存模式，看共享内存的使用方式，看线程束发散。矩阵乘法一定要自己实现一个基于shared memory tile的版本，因为很多笔试题目就是从这变形来的。

第二步，弄懂GPU硬件的关键参数。知道“线程束大小是32”“共享内存上限一般几十KB”“L2缓存是最后一级统一缓存”，知道“寄存器溢出到局部内存”是什么概念。不用背全部参数，但从参数里能推导出性能瓶颈的方向。

第三步，去读一遍主流的推理引擎或训练框架的算子融合实现。TensorRT的plugin机制、XLA的fusion pass、OneFlow的算子体系都可以。看它们是怎么做kernel融合、怎么做常量折叠、怎么调度通信的，这些都是面试官眼中“有实战经验”的证明。

第四步，练一遍多卡通信的计算题。不用真的搭集群，但要能写出AllReduce在不同拓扑下的通信量公式，能算出环形优于星形的具体倍数。这类题是笔试里的“压轴拉分题”，会算的人不多。

我个人的体会是，准备异构计算岗，最忌讳的是只背API。API会渐渐过时，但“内存层次决定访存优化空间”“通信开销和计算收益要一起评估”这些底层思维，在哪个芯片上都成立。2018年这套题即使放到现在看，考察的底层逻辑也没有过时，反而因为AI芯片种类变多，异构计算的“异构”二字被赋予了更广的含义。你把这个思维打牢，不管以后面试的是GPU、NPU还是其他AI加速卡岗位，都能站得住。

最后再分享一个实际经验：笔试里的编程题，如果能用到的优化手段，就大胆写上去。我当时在向量加法的题后面额外加了一段说明，描述“可以把cudaMemcpy改成cudaMemcpyAsync”并配合stream来重叠数据传输和kernel计算。这个点本身不是题目要求，但考官在评语里专门提了一句“对异步执行有理解”。所以别小看笔试里的“附加分”，有时候就是你多写的那两行设计思路，让你从一堆同样写对答案的人里被挑出来。