看到这个标题我愣了一下2018年商汤科技校招GPU优化工程师第一场笔试在当年的技术圈里几乎成了一个符号——那是少数把“底层性能”作为核心考点的算法类笔试。今天想认真复盘一下这场笔试的考点分布、出题意图和备考路径给准备AI基础设施、高性能计算、算子优化方向笔试面试的同学一份真正能参考的经验。它不是一场考API记忆的考试而是一场考你懂不懂硬件执行真相的考试。1. 这场笔试的底层筛选逻辑招的不是“会调库”的人1.1 2018年为什么GPU优化工程师突然变得抢手2018年恰好是深度学习训练规模暴涨的节点。模型越来越大训练数据越来越多单卡根本扛不住分布式训练成为标配。但算法工程师普遍用Python写模型底层那些真正吃性能的算子比如卷积、矩阵乘、BatchNorm、Reduce全都要靠底层优化工程师来加速。商汤作为以算法见长的AI公司训练卡的数量和负载是很大的算法团队跑一个实验要多久直接影响整个公司的迭代效率。所以GPU优化工程师在这个时间点变得特别重要。这个岗位干的事很清晰把算子的执行效率榨干把训练框架的调度开销打下去把显存带宽用满。这就决定了笔试不会考“你会不会用PyTorch”而是考“你知不知道GPU是怎么把一条指令执行完的”。1.2 “第一场”笔试的定位商汤校招GPU优化方向一般不止一场笔试所以才会有“第一场”的说法。第一场的功能更像海选用它来快速过滤掉不具备硬件基础的人。这说明它的题量不会太小知识点覆盖会比较宽但深度又会控制在“一个计算机基础扎实的人能推出来”的范围内。和常见的LeetCode算法笔试有很大区别算法岗笔试考的是数据结构和逻辑推理而GPU优化岗笔试更看重你对体系结构、并行模型、访存行为这些底层知识的理解。举个很直白的例子——算法笔试让你写一个动态规划而GPU优化笔试可能会让你算一个kernel在特定GPU上跑完要多少微秒。这两者的思维模式完全不同。1.3 考卷的典型结构按当年圈里流传的复盘信息这类笔试大多由三部分组成基础客观题概念选择或填空覆盖CUDA编程模型、GPU架构、内存体系、并行原语。计算与分析题给一个场景手算访存带宽、线程配置、共享内存大小、占用率等。编程与优化题手写CUDA kernel或者分析一段给定kernel的优化空间。整张卷子真正考验的其实就三件事知不知道硬件细节能不能定量估算性能有没有自己的性能优化思路。这三件事做好了哪怕个别题目没完全写对也会在后续面试里被捞回来。2. 架构考点线程组织、warp与内存体系2.1 grid、block、thread的分层逻辑GPU的编程模型是分层的一个kernel启动后对应一个gridgrid由多个block组成block内部再分thread。这个分层不是随意设计的它对应的是硬件执行的真实节奏。block是软件概念但它在SM流式多处理器上执行。一个block里的线程被分到同一个SM上它们在SM内部通过共享内存协作通过__syncthreads()同步。所以block不能设得太大硬件有上限。老一代CUDA版本里一个block最多1024个线程直到现在这个限制还是存在。笔试经常在这埋坑问“一个block最多能开多少个线程”很多人写2048或者4096都是错的正确的写法是“对于绝大多数现代架构是1024”。block也不能设得太小。一个SM里有固定的线程槽位比如Maxwell、Pascal这些架构上一个SM最多能驻留2048个线程。如果你把block设成64个线程又想让SM满负荷跑那就得同时调度32个block但block数量本身也有限制每个SM最多16~32个block视架构而定。所以笔试里经常出现“给你一个GPU的规格参数让你设计合适的block大小”这种题。它考的并不只是记住最大值而是理解block大小如何影响SM内的资源分配和调度效率。2.2 warp永远等于32SIMT的真相GPU的最小调度单位不是thread而是warp。一个warp固定包含32个线程。硬件按warp来发射指令一个warp中的线程在同一时刻执行同一条指令只是作用于不同的数据。这就是SIMT单指令多线程。这里有一个笔试非常爱考的点分支发散divergence。当一个warp里的线程因为if (tid % 2 0)之类的条件走了不同分支时硬件会把执行串行化——先执行真分支的一部分线程再执行假分支的另一部分线程。这意味着这个warp的吞吐直接减半。优化方法通常是把同一个分支的线程尽量凑到同一个warp里或者用数据重组织消除发散。笔试出题一般不会让你写复杂的分支消除代码但会给你一段代码问“这个kernel在多少个warp上发生了发散如果每个warp只有一条分支指令浪费了多少执行周期”这时候只要能抓住“warp是调度单位发散会串行化”这一点就能算明白。2.3 内存层次GPU优化的分水岭GPU优化本质上就是内存优化这句话在笔试里体现得淋漓尽致。GPU有复杂的内存层次每一层的容量、延迟、带宽都不一样优化思路都是围绕“让数据尽量在快层里待着”展开。内存层次作用域典型容量延迟量级特点与使用场景寄存器线程私有每个SM几百KB纳秒级速度最快但容量有限溢出会落到local memory共享内存sharedblock内共享每个SM几十到上百KB20~30个周期需要显式管理block内线程通信的主要手段L2缓存设备全局几MB200~400个周期对global memory访问的透传缓存所有SM共享全局内存global设备全局16~32GB400~800个周期容量最大延迟最高带宽受限于显存规格常量内存/纹理内存特殊用途较小依赖于缓存命中只读场景下可走专用缓存路径笔试最常考的就是寄存器、共享内存、全局内存三者之间的差异。一个很典型的题目是写出GPU端不同类型内存的访问延迟排序。有人把L2缓存放在共享内存前面有人把常量内存当成全局内存的替代都是概念不清。实际上寄存器最快共享内存次之L2再次全局内存最慢。还有一个高频考点是bank conflict。共享内存在物理上被分成32个bank每个bank在同一个周期里只能响应一次访问。如果同一个warp里的32个线程同时访问同一个bank的不同地址就会发生冲突冲突的访问会被串行化。笔试题目经常是这样的一个warp的线程访问shared[tid * 2]问有多少路bank conflict。按32个bank分布tid*2的跨步是2会导致2路冲突该访问需要两个周期而不是一个周期。这种题只要画一下bank分配图就能看懂但很多人在考场上因为没专门练过而卡住。3. 计算题让你当场估算性能的套路3.1 从需求反推block和grid配置笔试里有一类必出现的计算题给你一个数组大小、一个kernel让你算需要多少block、每个block多少线程。这类题看着简单实际很容易丢分。举个例子一个kernel要处理4M个float数据每个block设置256个线程问需要多少个block。[ 4,000,000 / 256 15625 ]这是送分题。但进阶一点会加条件每个线程一次处理4个元素问需要多少block。这时每个线程覆盖4个元素实际需要的是[ 4,000,000 / (256 \times 4) 3906.25 ]注意这里不能直接取3906因为会漏掉最后0.25个block的数据。更好的做法是向上取整到3907个block然后在kernel内部做边界判断防止越界访问。很多人在这一步忽略了“除不尽的尾块”导致一个越界读的bug笔试改卷时这种细节很扎眼。再进一步如果你用grid-stride loop让每个线程在全局维度上以gridDim.x * blockDim.x为步长循环处理数据那block/grid的配置就不需要严格匹配数据量。这种写法在现代优化里非常常见笔试里主动提出来会是很强的加分项因为它说明你不只满足于“能跑”还关心“怎么跑得更优雅”。3.2 访存带宽估算判断计算密集还是访存密集GPU优化决策的第一步永远先问这个任务是被计算卡住还是被访存卡住。笔试喜欢通过给参数让你算“理论耗时”来考察这个思维。假设一个kernel要读取4MB的float数据做一次简单的加法和写入GPU的显存带宽是700GB/s问在理想情况下读写数据的时间是多少。读取4MB[ 4MB / 700GB/s \approx 5.7\mu s ]写入4MB[ 4MB / 700GB/s \approx 5.7\mu s ]总访存时间约 ( 11.4\mu s )。如果这个kernel的计算量非常小比如每个元素只做一个乘加那计算时间可能只要几微秒甚至更短整体耗时大概率会被访存时间盖住。结论就是这个kernel是memory bound优化时要优先减少访存量比如用共享内存缓存重复用到的数据而不是去调整指令流水。这类题目背后的核心概念是Roofline模型把计算量和访存量画在一张图上斜率就是机器的浮点算力与带宽之比。低于斜率的任务是访存瓶颈高于斜率的任务才是计算瓶颈。笔试可能不会直接要求你画Roofline但“通过估算访问用时判断瓶颈”就是Roofline的朴素形式。在答题时把这个逻辑写清楚比生硬地套公式更能体现功底。3.3 占用率、寄存器数与共享内存的博弈占用率occupancy指一个SM上活跃线程占总线程槽位的比例。占用率越高越能隐藏访存和指令延迟。但占用率不是越高越好因为寄存器和共享内存都是有限的。拿Pascal架构举例一个SM最多驻留2048个线程寄存器文件有65536个32位寄存器。如果每个线程用32个寄存器那么SM上最多能同时承载[ 65536 / 32 2048 ]刚好能跑满线程槽位占用率100%。但如果每个线程的寄存器用量涨到64个能承载的线程数就变成1024占用率直接掉到50%。笔试经常会这样问某个kernel每个线程占用了40个寄存器block大小设为512问SM上能并发多少个block。答案不是简单的2048/5124而是要先算寄存器约束[ 65536 / (512 \times 40) 3.2 ]向下取整是3个block。也就是说就算线程槽位允许放4个block寄存器也不够用。这种题考的是“资源是联动的不能单看一个指标”。同理共享内存也会限制并发block数。如果一个block需要20KB共享内存而SM上可用共享内存只有48KB那最多并发2个block。笔试中给出“每个block用X KB共享内存SM共享内存总量Y KB问最多几个block并发”是经典套路本质就是一个整数除法和向下取整。4. 手写kernel从“能跑”到“能打”的reduce路线4.1 为什么用reduce作为编程题一提到手写CUDA笔试最常见的就是大数组归约。这题经典的原因在于它足够简单人人都能写一个能跑的版本它又足够丰富能展示出访存、同步、归约树、甚至warp shuffle等一堆优化点。改卷人一眼就能看出你是背了一个模板还是真正理解每一步在干什么。最简单的版本是每个线程从全局内存取一个元素同一个block内先做一个树形归约最后再对block的部分和做二次归约。这个版本能跑但性能一般因为每个线程只读了一个数据访存完全没有做到一次读取多个连续元素索引计算和内存事务的开销占比很高。4.2 朴素版本骨架下面这个版本适合作为笔试答题的起点先把逻辑写对再谈优化。__global__ void reduce_kernel(const float* in, float* out, int n) { __shared__ float sdata[256]; int tid threadIdx.x; int i blockIdx.x * blockDim.x tid; // grid-stride loop保证一个线程处理多个元素减少block数量 float v 0.0f; while (i n) { v in[i]; i gridDim.x * blockDim.x; } sdata[tid] v; __syncthreads(); // 树形归约 for (int s blockDim.x / 2; s 0; s 1) { if (tid s) { sdata[tid] sdata[tid s]; } __syncthreads(); } if (tid 0) { atomicAdd(out, sdata[0]); } }启动时block数量不需要精确匹配数据量因为grid-stride loop会保证每个block内的线程以固定步长循环读取直到把整个数组读完。atomicAdd负责把每个block的部分和累加到全局结果变量里。这个版本的正确性关键点有三处第一次写入共享内存后的同步、树形归约每一层结束后的同步、以及对n可能不是blockDim整数倍的边界处理。4.3 优化点连续访问、向量化与warp shuffle把这个朴素版本写完真正的分水岭在于能不能继续往下列优化点。首先是连续访问。朴素版本里in[i]的访问模式是每个线程跳着读。如果改成每个线程连续读取一段区间比如先用一个循环让线程读取连续的4个float会明显改善内存事务效率。背后的原因是GPU访存是按128字节为粒度的事务发起的如果同一个warp里的线程访问的是连续地址一次事务就能把数据带回来如果是跨步访问可能需要多个事务。第二个是向量化访存。用float4类型一次读4个float可以把内存事务数量降到原来的四分之一。2018年的笔试里明确提出float4已经是比较进阶的答案但我相信在改卷人眼里这比只会写“每个线程处理多个元素”要强很多。第三个是warp shuffle。Volta和Pascal架构支持__shfl_down_sync之类的指令warp内部的线程可以直接交换寄存器里的数据不需要经过共享内存。这能省掉一部分共享内存的访问和同步。笔试不要求把代码完全写对但如果能写出“在warp内先做shuffle归约再把每个warp的结果写到共享内存”这样的思路明显比只会死记硬背共享内存归约的人高一个层次。4.4 容易翻车的代码细节手写CUDA的翻车点非常固定但每年都有人踩忘记同步第一次把局部值写入共享内存后如果没有__syncthreads()另一个线程可能读到旧值。面试和笔试都特别喜欢盯着这个点问。浮点累加顺序浮点数不满足结合律不同线程的累加顺序会带来微小结果差异。笔试如果问“为什么reduce结果和CPU串行版本不完全一致”答案就是浮点舍入。边界处理block数超过数据量不够除尽时不加判断直接访问会越界。这个错在真实项目里是宕机级别的。atomicAdd参数类型老版本CUDA对double类型的atomicAdd支持不完整笔试题目如果用了double要格外注意。grid-stride loop的步长有人写成i blockDim.x漏了gridDim.x结果每个block都从同样的位置开始算出现重复累加。这是非常隐蔽的逻辑错误笔试改卷时一眼就能看出来。5. 商汤这类AI公司笔试的隐藏分算子优化背后的系统观5.1 深度学习算子和通用并行优化是同构的很多同学不理解为什么AI公司GPU优化笔试不直接考卷积、不考cuDNN而是考reduce、矩阵乘、访存分析。因为卷积在算子层面确实很复杂但把它拆开之后核心组件无非是矩阵乘GEMM、数据重排im2col、归约reduce和向量化访存。能写好通用并行kernel的人经过短期训练就能上手深度学习算子反过来只背过卷积公式但不懂访存行为的人面对真实优化场景会完全无从下手。商汤这类公司内部会有大量自研算子和训练链路优化工作不会单纯依赖cuDNN/cuBLAS。因为框架场景千变万化从模型结构到输入shape都可能定制底层算子也要跟着定制。笔试考通用并行能力本质上是想选拔“遇到新算子能快速分析瓶颈并动手优化”的人而不是“见过多少个现成算子”的人。5.2 瓶颈定位能力profile-first思维我自己的体会是GPU优化工程师最重要的能力不是手写kernel而是定位瓶颈的能力。笔试里那些计算题实际对应到工作中就是你拿到一个算子得先判断它是访存密集还是计算密集再决定优化方向。优化算子的正确路线基本是固定的先用profiler跑一遍看访存吞吐、计算吞吐、占用率、bank conflict、分支发散这些指标然后根据瓶颈做针对性优化优化完再profile验证。这个过程在笔试里没法完全展现但可以通过答题套路暴露出来。比如遇到“给一段代码问哪里慢”的题如果你能先写“首先看访存模式其次看分支和同步最后看占用率”而不是一上来就改算法说明你有profile-first的思维习惯。这种答题习惯是实打实的加分项。备考这类笔试我个人建议把下面这些概念当成“必须能用自己的话讲清楚”的底线Roofline模型、Amdahl定律、访存密集与计算密集、warp发散、bank conflict、占用率、网格跨步循环、向量化访存、浮点精度陷阱、原子操作的性能代价。每个概念背后都对应一个笔试题型也对应一个真实工作中的优化场景。5.3 如果只准备两到四周怎么规划时间紧的情况下与其刷一堆算法题不如直接走“最小可行准备路径”第一周学CUDA编程模型和GPU架构基础把grid、block、warp、共享内存、寄存器这些概念彻底搞明白每天手写一个简单kernel。第二周专攻内存优化。做矩阵转置、矩阵乘、reduce这三个经典例子反复对比朴素版本和优化版本的性能差距用profiler验证。第三周整理性能估算的计算套路把访问带宽、耗时、占用率、bank conflict这些计算题各练十道做到看到参数能马上列公式。第四周复盘笔试风格把每个知识点都能讲出一个“为什么”比如为什么block不能设太大、为什么共享内存要手动管理、为什么需要网格跨步循环。把自己写的kernel放到真实GPU上跑一跑看看profiler输出和你理论估算的差距这一步比多刷十道题都管用。真正面试时能讲清楚“我做了哪些优化、每一步带来了多少提升”的人远比只会背概念的人有竞争力。我自己的体会是这种笔试真正筛选的不是“背过多少CUDA API”而是“对性能有没有直觉”。哪怕你卷子里的代码没有完全跑通只要每道题都体现出了“我知道瓶颈在哪、我知道怎么权衡资源、我知道先估算再动手”的思维方式面试官都会愿意多聊几句。那张卷子之后很多年我判断一个新人写的代码值不值得合入第一反应依然是“它访问了多少次不该访问的全局内存”。这大概就是当年那场笔试留给我的职业病也是我特别想推荐给后来者的核心能力。