前几天整理资料时翻到当年准备商汤校招GPU优化工程师笔试的笔记,恰好最近又有人问到AI公司基础架构岗位该怎么准备。这个岗位的笔试很典型:看似只考CUDA,实际上一张卷子把并行计算、体系结构、算法复杂度、甚至Linux系统基础全串起来了。我按第一场的题型分布和考点逻辑,把整个拆解过程整理出来,不讲虚的,全是备考时反复琢磨过的内容。
1. 笔试到底在筛什么样的人:从岗位职责倒推考察范围
我准备笔试的习惯是先看岗位JD,再倒推可能考什么。商汤这类AI公司的GPU优化工程师,核心任务是把算法工程师写的模型算子跑得更快,或者把新的算子在GPU上高效实现出来。这意味着你既要懂上层算法在算什么,又要懂底层硬件是怎么执行的,中间全靠CUDA这座桥连着。
1.1 岗位职责决定了知识边界
GPU优化工程师日常做的事围绕三个方向:算子优化(把某个卷积、矩阵乘从慢版本改成快版本)、性能分析(用profiler找出瓶颈在访存还是计算)、框架集成(把优化好的算子接入PyTorch/TensorFlow使用)。这三件事对应到笔试里,正好是三类题目:手写或分析CUDA kernel、对给定代码做性能瓶颈分析、以及对现有算子的优化方案设计。
所以笔试不会只考语法层面的“CUDA怎么写”,而是考“为什么这样写性能更好”。比如同样是一个向量加法,问你用grid-stride loop还是直接映射,绝不只是写法偏好,背后是对启动开销、负载均衡、可扩展性的理解。这类题目没有标准操作步骤,但考察的知识边界非常清晰:CUDA编程模型是基础,GPU架构是底层支撑,访存优化是核心技法,算法基础是隐含前提。
1.2 笔试的整体结构与题型分布
从岗位特点倒推,这场笔试大致分四个部分:选择题/填空题考察基础概念,简答题考察原理理解,编程题考察代码落地能力,压轴的设计题考察整体优化思维。时间一般120分钟左右,题量不算大,但每道题都需要仔细推导,尤其是设计题,给一个算子让你描述优化思路,写清楚每一步为什么这么做,比写完整代码更能拉开差距。
我复盘时的一个直观感受:题目难度梯度拉得很开。前面概念题几乎是送分,但越往后越需要真功夫。如果前面基础题做得犹豫,后面大题基本没时间深入思考。所以策略上要先把送分题稳稳拿到,再集中火力攻设计题。
2. 核心知识点逐个击破:每个考点背后的原理逻辑
这一节把笔试涉及的几个重点模块展开讲,每个点都按“是什么、为什么考、怎么答”三层拆开。这样复习的时候不是死记硬背结论,而是理解推导过程,考场上不管题目怎么变都能应对。
2.1 CUDA编程模型:笔试的绝对核心
CUDA编程模型是整场笔试的主线。从grid、block、thread的三层结构,到kernel启动方式,再到同步机制,每一个都是考点。你需要清楚地知道一个kernel启动后,线程是怎么组织成grid和block的,block内的线程如何协作,block之间又怎么通信。这些概念不仅是语法,更是在描述一个并行任务如何被切分、调度和执行。
考察时常见问法有:一个grid最多能开多少个block,一个block最多能有多少线程,这些硬件限制数字是多少。很多人觉得这类题是死记硬背,其实不是。限制数字背后是硬件的设计取舍。比如block线程数上限是1024,是因为调度器一次管理的线程粒度、寄存器分配的粒度,都围绕这个数量级设计。理解了这层逻辑,数字自然记得住。
另一个高频考点是同步。block内部用__syncthreads()做屏障同步,保证共享内存读写一致。笔试喜欢问:如果block内线程分支发散(if-else),__syncthreads()会不会出问题。答案是会,因为__syncthreads()要求所有线程都到达才能继续,如果某些线程走了分支提前return,就会死锁。这类题考察的是对线程执行模型的理解深度,不是API记忆。
2.2 GPU硬件架构:问的是“为什么快”
如果CUDA是语言,GPU架构就是运行环境。笔试不会直接问某个参数,而是通过性能题间接考察。比如经典问题:为什么GPU适合并行计算而CPU不适合。答题不能只说“GPU核心多”,要落到SIMT执行模型上:GPU通过大量线程的并行隐藏访存延迟,CPU则靠大缓存和分支预测减少延迟。这是两种完全不同的设计哲学。
更细一点,要理解SM(Streaming Multiprocessor)内部的结构。一个SM里有多个SP(流处理器),以warp(通常32线程)为调度单位。warp是GPU调度的最小单位,同一warp的线程执行同一指令,如果分支发散,warp会串行执行每个分支路径,这就是性能损耗的来源。答题时能画出这个执行流程,再解释为什么发散分支慢,就能拿全分。
还有一个常考的点是occupancy(占用率),即SM上活跃warp数与最大warp数的比值。占用率越高,越能隐藏延迟,但也不是越高越好。寄存器使用过多会降低占用率,shared memory使用过多也会限制block数量。笔试常给一个配置让你计算占用率,这类题需要熟练每一代架构的硬件参数,至少要知道SM最大线程数、寄存器文件大小、shared memory大小这些常用数值。
2.3 访存优化:高频考点的重中之重
GPU优化圈有句话:优化访存比优化计算更划算。因为GPU的计算吞吐极高,但显存带宽相对有限,很多kernel跑不满计算单元,瓶颈在等数据搬过来。笔试对这块的考察也最细,常考的知识点包括合并访问、共享内存、bank conflict。
合并访问(coalesced access)是显存访问优化的第一原则:同一warp的32个线程访问连续地址时,硬件合并成少数几次内存事务,带宽利用率最高。笔试会给一段访存代码,问你访问是否合并,如果不合并怎么改。比如按行存储的二维数组,行遍历和列遍历的性能差异极大,原因就在于合并访问。这个例子几乎每年都有。
共享内存(shared memory)是GPU上的片上存储,速度远快于全局内存,但容量小、需要手动管理。典型考法是给一个需要数据复用的场景,让你用shared memory优化,并分析优化前后访存量变化。比如矩阵转置这种经典题,本质就是利用shared memory避免非合并访问。这种题不仅考代码,更考访存行为分析。
bank conflict是共享内存优化的进阶考点。共享内存被划分为32个bank,同一warp多个线程同时访问同一bank的不同地址,就会冲突串行化,性能骤降。笔试常考:给定共享内存数组的访问模式,判断是否有bank conflict,并说明如何padding(加一列)避免。这个知识点对没有实际写过优化的人有点抽象,但确实是区分度很高的题。
2.4 算法与数据结构基础:被低估的一环
GPU优化笔试不是纯CUDA考试,算法基础同样占比重。并行归约(parallel reduction)、前缀和(scan)、Histogram这些GPU上的经典算法模式一定要熟。这些算法本身不难,但在GPU上实现涉及线程组织、同步、负载均衡,比CPU版本复杂得多。
以并行归约为例,笔试常考的是如何避免线程发散。最简单的方式是让奇数线程退出,但这样会造成warp内线程不活跃。更优的是stride减半的方式,让连续线程参与操作,保持warp内活跃度。这类细节如果不实际写代码很难体会到,但笔试就是会考。
另一个容易忽略的是时间复杂度分析。GPU上不能只看计算复杂度,还要算访存复杂度。如果一个算法的访存量是计算量的好几倍,那瓶颈就可能在带宽上。笔试如果让你分析两个kernel的优劣,一定要把计算量和访存量分开分析,再判断谁才是性能瓶颈。
3. 实操题复盘:用矩阵乘法拆解优化全流程
矩阵乘法是GPU优化的“hello world”,也是笔试编程题最常考的原型。下面从朴素版本开始,一步步加优化,每步都讲清楚为什么。这套思路同样适用于卷积、全连接等算子的优化分析。
3.1 朴素版本:先把功能跑通
__global__ void matmul_naive(float *A, float *B, float *C, int N) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; float sum = 0.0f; for (int k = 0; k < N; ++k) { sum += A[row * N + k] * B[k * N + col]; } C[row * N + col] = sum; }这个版本每个线程算输出矩阵的一个元素,每次迭代都要访问A的一行和B的一列。A[rowN+k]在同一warp内不同col的线程访问时是连续地址,合并访问没问题。但B[kN+col]就麻烦了:不同col的线程访问的是同一行不同列,地址不连续,warp的32次访问被拆成32个内存事务,带宽利用率极低。
笔试如果让你分析这个版本,这就是核心答分点:B矩阵的访问是非合并的,导致大量访存浪费。这个问题是理解后面所有优化的基础。
3.2 使用共享内存分块:访存量直接降一个量级
优化的基本思路是利用数据的局部性:C[i][j]的计算需要A的第i行和B的第j列,相邻的输出元素会复用同一行、同一列的数据。这部分数据量远小于计算量,可以加载到共享内存里反复使用。
#define TILE_SIZE 16 __global__ void matmul_tiled(float *A, float *B, float *C, int N) { __shared__ float As[TILE_SIZE][TILE_SIZE]; __shared__ float Bs[TILE_SIZE][TILE_SIZE]; int row = blockIdx.y * TILE_SIZE + threadIdx.y; int col = blockIdx.x * TILE_SIZE + threadIdx.x; float sum = 0.0f; for (int tile = 0; tile < N / TILE_SIZE; ++tile) { As[threadIdx.y][threadIdx.x] = A[row * N + tile * TILE_SIZE + threadIdx.x]; Bs[threadIdx.y][threadIdx.x] = B[(tile * TILE_SIZE + threadIdx.y) * N + col]; __syncthreads(); for (int k = 0; k < TILE_SIZE; ++k) { sum += As[threadIdx.y][k] * Bs[k][threadIdx.x]; } __syncthreads(); } C[row * N + col] = sum; }每次迭代先把一个TILE_SIZE×TILE_SIZE的子块加载到共享内存,再用这些数据做计算。全局访存量从每个输出元素读N+N个数,变成每个block读2×N×TILE_SIZE个数(一个block算TILE_SIZE×TILE_SIZE个输出),整体访存量降为原来的1/TILE_SIZE。TILE_SIZE=16时,访存量直接降16倍。
注意两个__syncthreads()的位置:第一个保证数据全部加载后再计算,第二个保证所有线程算完再用下一轮数据覆盖共享内存。笔试如果让你补全代码,这两个同步是必考点。
还有个细节容易被忽略:加载As时A[row*N + tile*TILE_SIZE + threadIdx.x],同一warp的threadIdx.x连续,地址连续,合并访问。加载Bs时B[(tile*TILE_SIZE+threadIdx.y)*N+col],同一warp的threadIdx.x作为col变化,地址也是连续的(因为col连续),这里B按行存储的话,访问其实是跨行的,但每个线程访问不同行的同一偏移,地址不连续。真正正确的写法应该让Bs也按行加载。笔试遇到过这种陷阱,答题时最好把访存图景画清楚。
3.3 进阶优化:向量化与流水线
分块版本已经能应对大部分笔试场景,但想拿高分,还可以补充几个进阶思路。一是float4向量化,每次读写4个float,减少指令数;二是循环展开,减少循环开销和增加指令级并行;三是双缓冲(用两个共享内存数组交替加载,隐藏加载延迟)。这些点不一定要求现场实现,但在设计题里作为优化方案提出来,能明显体现你的实战积累。
向量化的核心是用float4类型一次读写16字节,配合合并访问,带宽利用率能更高。双缓冲则是在计算当前tile的同时,预加载下一个tile的数据,用计算掩盖访存延迟。笔试时间有限,能写清思路和预期收益就足够,不必追求完整代码。
我把这道题的优化路径完整写出来,是想说明一个答题套路:从朴素实现出发,识别访存瓶颈,用共享内存做数据复用,再逐步叠加高级优化。这个“识别瓶颈→定位优化手段→分析收益”的路径,几乎可以套用所有优化设计题。
4. 做题节奏与答卷策略:稳稳拿分比炫技重要
考场上除了知识点,策略也很重要。一场笔试120分钟,既要保证基础题不丢分,又要给大题留足时间。我自己的经验是分三遍做题,每一遍有不同的目标。
4.1 时间分配:三遍做题法
第一遍快速扫完全卷,把有把握的题直接做掉,大概花30分钟。这里说的“有把握”包括概念题、简单填空题、熟悉的代码分析题。第一遍的目的是稳定心态,确保基础分进账。
第二遍集中攻中等难度题,大概花50分钟。这轮遇到的主要是代码填空题、计算题、简单设计题。答题时注意把计算过程写清楚,即便最终结果出错,过程分也能拿一部分。GPU优化笔试的阅卷很看重分析过程,因为优化本身没有唯一答案,只要推理合理就有分。
第三遍剩下的时间全部投入压轴设计题,大约40分钟。这轮不要追求写完整代码,而是把设计思路结构化地写出来:瓶颈分析→优化方案→预期收益。即使代码没写全,思路对了就能拿大半分数。
4.2 不会做的题怎么办
笔试总会遇到卡壳的题,这时候最忌讳死磕。我的原则是:选择题先标记跳过,最后随便选一个也不要空;简答题写你知道的相关知识点,哪怕不完整也能拿部分分数;代码题哪怕写伪代码,也比不写好。GPU优化是个开放性很强的领域,只要分析方向对,改卷老师会给分的。
具体来说,遇到性能分析题不会算具体数字,就把分析思路写完整:从访存模式看、从占用率看、从同步开销看,列出可能的影响因素,再给出大致的优化方向。这种题的核心是考察你是否具备性能分析的思维框架,结果反而是次要的。
4.3 笔试之后该补什么
笔试结束后我习惯把每道错题对应到具体知识点,再补齐那块短板。比如发现自己shared memory的bank conflict分析老出错,就专门找几道相关题目练;发现自己对线程组织不够敏感,就写几个不同的kernel对比性能。这个复盘过程比做新题更有价值,因为能精准暴露弱点。
补短板时最有效的方式是实际跑代码验证。纸上分析是一回事,跑起来看profiler的输出又是另一回事。笔试备考期间如果条件允许,强烈建议把每个经典优化亲手实现一遍,用ncu看看实际指标,这样分析题才能答得言之有物。
5. 常见失分点与备考避坑建议:前人踩过的坑别踩
最后整理几个备考时常见的失分点和对应的避坑建议。这些坑我基本都踩过,复盘时一个个记下来,给后来的人少走些弯路。
5.1 轻视基础概念,把时间全花在刷题上
很多备考者一上来就刷各种复杂kernel,结果笔试前面概念题反而答得稀烂。GPU优化笔试的基础概念题占比不低,而且往往是最容易拿分的地方。比如线程层次、内存层次、warp大小这些概念,一定要做到闭卷能默写、能画图解释。不要在简单题上丢分。
我的建议是备考前两周先把概念体系过一遍,用思维导图把CUDA编程模型、GPU架构、优化方法串起来。每学一个概念,就问自己两个问题:它解决什么问题,不这么设计会怎样。能用一两句话讲清楚,才算真正理解。
5.2 只记结论不理解推导,换个问法就懵
GPU优化笔试很喜欢换着角度考同一个知识点。比如合并访问,可以考你判断某段代码是否合并,也可以考你解释为什么按列遍历慢,还可以反过来问你如何设计数据布局让访问合并。只记住“连续地址就是合并”这个结论,遇到后面的问法就答不出来了。
正确做法是记住底层原理:内存事务以固定大小为单位,warp内访问地址越集中,需要的事务越少,带宽利用率越高。有了这个理解,不管题目怎么变,都能往原理上推导。
5.3 不重视设计题的结构化表达
设计题其实是笔试中最好拿分的大题,因为评分看的是思路完整性,不要求标准答案。但很多人的答案写得很乱:一会儿说这个优化,一会儿跳到另一个,逻辑线不清晰。
我自己的模板是:先明确算子的计算特征和访存特征,分析当前瓶颈在哪个环节,然后按“先访存优化,再计算优化,最后微调”的顺序给出方案。每一步写清楚做了什么、为什么这么做、预期效果如何。这样写出来的答案,改卷老师一眼就能看到分析能力,分数自然高。
5.4 时间分配失衡,最后大题草草收场
前面也说过时间分配的问题。但这里特别强调一点:不要因为某道题有思路就写太多细节,导致后面大题时间不够。判断标准是:如果一道10分题你写了20分钟还没写完,就该停笔去做后面的20分大题了。考试不是写论文,分值和投入要匹配。
我自己考试时会带一只手表,每隔20分钟左右抬头看一眼时间。如果某一部分超时了,立刻调整节奏。这种时间管理能力,平时做模拟卷就可以练起来。
6. 备考工具的选型与使用心得
笔试备考不只是刷题看书,用对工具会让效率提升很多。我备考时主要用三样:NVIDIA官方的文档和编程指南、Nsight Compute性能分析工具、以及一块支持CUDA的GPU。有真实环境跑代码,很多纸上谈兵的问题会瞬间清晰。
6.1 官方文档怎么读才高效
CUDA C++ Programming Guide是权威资料,但全读不现实也没必要。我的方法是按需查阅:遇到不懂的内存模型就翻对应章节,遇到同步问题就看同步章节,不按目录顺序死读。更重要的是把文档内容和笔试考点对应起来,比如文档里讲coalesced access的部分,正好对应笔试访存分析题。
另一个容易被忽视的资料是NVIDIA的GTC演讲和案例分享。里面有不少真实优化案例,能帮你建立“从实际问题到优化方案”的直觉。虽然不直接对应考点,但对设计题很有帮助。
6.2 用Nsight Compute验证心中猜想
备考时做分析题最怕的就是“感觉是这样但不确认”。我强烈建议把经典优化自己实现一遍,用Nsight Compute看kernel的访存吞吐、计算吞吐、占用率、warp状态等指标。你会发现很多纸上分析时会忽略的细节,比如shared memory的bank conflict实际没那么明显,block大小对性能的实际影响也不是线性的。
我备考时有一个小本子,专门记录“纸上分析结果”和“实测结果”的差异。这些差异后来成了我最宝贵的备考资料。笔试遇到性能分析题时,我能写出的细节和深度明显超出只看书的考生。
6.3 个人实操后的性能改观数据
这里分享一个用分块优化矩阵乘法的实测数据,帮助大家建立直觉。环境是个人电脑上的入门级GPU,N=1024的矩阵乘法:
- 朴素版本核函数耗时约12.6ms
- 增加分块共享内存优化后,耗时降到约2.1ms
- 再叠加float4向量化后,耗时约1.4ms
这三步优化从12.6ms降到1.4ms,接近9倍提升,中间没有用任何高深技巧,只是把访存模式理顺了。这就是GPU优化的魔力,也是笔试想考察的核心能力:不是记住某个奇技淫巧,而是掌握系统性的优化思维。如果你在笔试中能把这个优化思路一步步写清楚,就已经超过大部分考生了。
我个人在准备这场笔试时的一个深刻体会是:GPU优化不是一个靠背就能拿高分的科目,它需要你在真实硬件上反复试验,感受每一次改动带来的性能变化。准备笔试的过程,本质上就是一次系统性的并行计算入门训练。即使最后没进面试,这个过程中建立的对性能分析的敏感度,对后续做任何高性能计算相关的工作都极有帮助。希望这篇拆解能帮你少走一些弯路,祝笔试顺利。