我参加过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缓存、常量内存、纹理内存。每层内存的延迟和带宽差出两个数量级。笔试里有一道填空题直接给了几个数字让你判断分别对应哪类内存。我回忆一下大概的数字量级内存类型延迟量级容量量级主要用途寄存器几个周期KB级/线程私有线程内局部计算共享内存20-30个周期几十到上百KB/blockblock内线程间数据共享L2缓存200-400个周期MB级全局内存的缓存层全局内存400-800个周期GB级主数据存储常量内存广播优化后可近似L1延迟64KB所有线程只读的常量我当时在卷子上写答案的时候漏写了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_addblocks, 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 一道送命题如何优化AllReduceAllReduce是分布式训练里的最核心通信原语。笔试里问它的优化方式算是既基础又进阶的题。从朴素版到高性能版大致有这几个级别第一级不做任何优化所有卡把梯度发给卡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计算。这个点本身不是题目要求但考官在评语里专门提了一句“对异步执行有理解”。所以别小看笔试里的“附加分”有时候就是你多写的那两行设计思路让你从一堆同样写对答案的人里被挑出来。
