前几天整理资料时翻到当年准备商汤校招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则靠大缓存和分支预测减少延迟。这是两种完全不同的设计哲学。更细一点要理解SMStreaming 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[rowNk]在同一warp内不同col的线程访问时是连续地址合并访问没问题。但B[kNcol]就麻烦了不同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的子块加载到共享内存再用这些数据做计算。全局访存量从每个输出元素读NN个数变成每个block读2×N×TILE_SIZE个数一个block算TILE_SIZE×TILE_SIZE个输出整体访存量降为原来的1/TILE_SIZE。TILE_SIZE16时访存量直接降16倍。注意两个__syncthreads()的位置第一个保证数据全部加载后再计算第二个保证所有线程算完再用下一轮数据覆盖共享内存。笔试如果让你补全代码这两个同步是必考点。还有个细节容易被忽略加载As时A[row*N tile*TILE_SIZE threadIdx.x]同一warp的threadIdx.x连续地址连续合并访问。加载Bs时B[(tile*TILE_SIZEthreadIdx.y)*Ncol]同一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 个人实操后的性能改观数据这里分享一个用分块优化矩阵乘法的实测数据帮助大家建立直觉。环境是个人电脑上的入门级GPUN1024的矩阵乘法朴素版本核函数耗时约12.6ms增加分块共享内存优化后耗时降到约2.1ms再叠加float4向量化后耗时约1.4ms这三步优化从12.6ms降到1.4ms接近9倍提升中间没有用任何高深技巧只是把访存模式理顺了。这就是GPU优化的魔力也是笔试想考察的核心能力不是记住某个奇技淫巧而是掌握系统性的优化思维。如果你在笔试中能把这个优化思路一步步写清楚就已经超过大部分考生了。我个人在准备这场笔试时的一个深刻体会是GPU优化不是一个靠背就能拿高分的科目它需要你在真实硬件上反复试验感受每一次改动带来的性能变化。准备笔试的过程本质上就是一次系统性的并行计算入门训练。即使最后没进面试这个过程中建立的对性能分析的敏感度对后续做任何高性能计算相关的工作都极有帮助。希望这篇拆解能帮你少走一些弯路祝笔试顺利。