商汤GPU优化笔试全解析:从CUDA编程到性能调优
2026/8/29 16:34:07 网站建设 项目流程

商汤科技2018校招的GPU优化工程师笔试,这几年在不少技术论坛和校招群里都被反复翻出来讨论过。原因很简单:这个岗位的笔试题目,基本代表了国内AI公司对"GPU优化"这个方向的真实要求,不是什么概念题背诵,而是直接看你对硬件、对CUDA、对性能分析的底层理解。我入行这些年,带过一些刚毕业的校招生,也帮圈里的朋友做过这类笔试辅导,今天把第二场笔试背后真正想考的东西拆开聊聊。这篇文章不追逐具体的真题答案(网上也传得七零八落),而是把"你以为在考什么"和"实际在考什么"这两层掰开,给准备投GPU优化方向的同学一条完整的复习主线。

凡是投过这类岗位的人应该都有同感:GPU优化工程师这个title听起来很硬核,但很多人在笔试前根本不知道该怎么准备。刷题?不像算法岗有LeetCode。背八股?网上关于CUDA的资料大多是"CUDA编程指南"那种翻译腔,跟实际考察方向对不上。实际上,商汤这场笔试的考察逻辑非常清晰——它不是在考你会不会写某个kernel,而是在考你懂不懂GPU为什么会慢、慢在哪里、怎么定位、怎么解决。这是一套完整的思维链,我今天就沿着这条思维链走一遍。

1. 先说清楚:GPU优化工程师在AI公司里到底干什么

1.1 为什么AI公司要专门养一波人做"优化"

很多人会有个误区,觉得AI公司里GPU优化工程师干的活儿就是"写CUDA"。真进去之后你会发现,这个岗位解决的是个非常现实的问题:一颗几万块的GPU,你到底用出了多少性能?

深度学习训练和推理的计算量不是闹着玩的。拿一个普通的ResNet-50模型来说,单张图片推理就要做几十亿次浮点运算,训练一个商用级模型动辄几百上千GPU小时。GPU本身就是公司最贵的硬件资产之一,如果因为kernel写得烂、显存布局不合理、数据搬运频繁,导致GPU利用率只有30%,那浪费的就是真金白银。优化工程师的存在,就是让同样预算买来的硬件跑出尽可能多的有效计算。

商汤这类公司更特殊,他们的业务里既有大模型的训练任务,又有大量端侧和云端推理任务。训练侧要跟NVIDIA的cuDNN、cuBLAS这些官方库抢性能,推理侧要针对自己的模型结构做算子融合和剪枝后的加速。这些活儿,不会写CUDA能干,但能不能写得比官方库还快、能不能针对自家模型做定制优化,就是优化工程师的看家本事。

很多同学问过我:"现在TensorRT、XLA这些工具不是越来越智能了吗?还需要人肉优化吗?"答案是:需要,而且缺口不小。工具能cover的是通用场景,真实业务里的模型结构千奇百怪,算子组合五花八门,总有些性能瓶颈需要懂底层的人手工分析和优化。这也是为什么这类笔试从来不问"你会不会用TensorRT",而是问底层原理。

1.2 这类笔试真正想筛选什么样的人

既然岗位需求明确了,笔试的筛选逻辑也就清晰了。商汤2018年那场第二场笔试,据我了解的情况和网上流传的讨论来看,考察方向大概分为三个层次。

第一层是"懂不懂GPU硬件"。GPU不是一块可以无限加速的万能芯片,它有自己的一套运作规律:成千上万个线程怎么调度,显存带宽有多大,计算单元和访存单元怎么配合。不理解硬件,写出来的kernel就是"能在GPU上跑"而已,离"跑得快"差了十万八千里。

第二层是"有没有真正的并行编程功力"。这个不是看你知不知道CUDA的API,而是看你能不能写出有合理线程组织、合理数据复用、合理访存模式的kernel。笔试里的编程题可能不大,但考察的都是基本功的细节。

第三层是"有没有性能分析和调优的思维"。给你一段程序,你能不能判断它是计算密集还是访存密集?瓶颈在寄存器还是在shared memory?该用哪种优化手段?这种分析能力,比单纯会写代码更值钱,也恰恰是最难通过临时刷题补上来的。

那场笔试网上最常被提到的,除了几道CUDA编程题,还有不少概念题和计算题。概念题逃不开内存模型、线程层次、同步机制这些基础;计算题常常会给你一个具体配置的GPU,让你算理论峰值、算访存带宽需求,然后问你一段代码是compute-bound还是memory-bound。这些题目看起来是知识点,实际上考的都是硬件理解。

2. 笔试的知识地图:必须吃透的核心考点

2.1 CUDA编程模型:所有的优化都从这里出发

说CUDA是GPU优化的基石,一点都不过分。笔试里不管你优化的是算子还是端到端模型,落到代码层面,都是写CUDA kernel。我对准备笔试的同学反复强调:CUDA编程模型里最重要的不是API列表,而是线程的层次组织和内存的层次结构这两个概念。

线程层次上,CUDA把线程组织成grid、block、thread三级。一个grid里有很多block,一个block里有很多thread,block里的线程可以通过共享内存和同步机制协作。为什么这个层次结构重要?因为它直接决定了你写kernel时的并行粒度。你在设计一个kernel时,首先要回答的问题就是:用多大的block?每个线程处理多少数据?block内的线程怎么协作?这些选择直接影响性能,笔试里遇到编程题,第一步就是做这种设计。

内存层次更是笔试的重灾区。GPU里寄存器最快但容量最小,每个线程私有;shared memory是block内共享的高速存储,速度也很快;全局内存(global memory)容量大但速度慢,是显存的主战场;还有只读的constant memory和texture memory。优化的核心思路之一,就是让数据尽量在高速的存储层次上复用,减少对慢速全局内存的访问。

举一个最典型的例子:向量加法。朴素写法是每个线程直接读全局内存里的两个数,加完写回全局内存。这个kernel每个线程只需要两次读一次写,性能基本被访存带宽锁死。但如果要做更复杂的操作,比如矩阵乘法,数据会被反复读取,这时候就需要把数据先搬到shared memory里,让block内的线程反复从shared memory读,而不是每次都要访问全局内存。这就是tiling分块的核心思想,也是笔试编程题最常见的考察点。

2.2 GPU硬件架构:懂硬件才能谈优化

很多同学复习CUDA时只看软件层面的API,结果笔试遇到"为什么这段代码跑得慢"就懵了。原因是:不了解GPU硬件,你无法判断一段代码的瓶颈在哪。

笔试里最常涉及的硬件概念就这么几个:SM(Streaming Multiprocessor)、CUDA core(也叫SP)、warp和warp scheduler。一个GPU里有多个SM,每个SM里有若干计算核心和调度器。线程执行时,block会被分配到某个SM上,SM里的线程实际不是一条一条执行的,而是按warp为单位调度,一个warp通常32个线程。这意味着同一个warp里的32个线程,在同一个时刻执行的是同一条指令,如果它们走入了不同的分支,就得串行执行,这就是"分支发散"(warp divergence)的由来,也是笔试里常用来考察概念理解的点。

还有一个绕不开的概念是memory coalescing,可以翻译成"访存合并"。全局内存的访问效率,取决于一个warp里的32个线程访存的地址是否连续。如果连续,硬件可以把这32次访问合并成少数几次内存事务,效率最高;如果离散,每次访问都可能触发一次独立的事务,延迟成倍增加。我在帮人看代码时,见过太多性能瓶颈都在这里——逻辑没问题,数据也对了,就是地址不连续,跑得奇慢无比。

笔试怎么考这些?通常不会让你去画芯片架构图,而是给一段代码,问你能优化哪里。你要能看出:这段代码的访存模式是不是连续的?有没有共享内存的使用?有没有不必要的全局同步?有没有bank conflict?这些都是硬件层面才能看出来的问题。

2.3 访存优化与计算优化:优化的两大主线

把所有GPU优化手段归纳起来,其实就两条主线:访存优化和计算优化。笔试和面试题,基本都围着这两条线转。

访存优化的核心目标是减少慢速访存。常用手段包括:调整数据布局让访存连续(coalescing)、用shared memory做tiling减少重复访问全局内存、避免 bank conflict、合理使用只读缓存等。这类优化的判断依据很简单:如果一段代码的实际耗时远高于计算耗时,大概率瓶颈在访存,优化的重点就是减少内存访问次数或者提高访存效率。

计算优化的核心目标是提高计算单元利用率。常用手段包括:减少分支发散、使用向量化指令(如float4)、提高指令级并行度(ILP)、使用FFMA等融合计算指令、在权衡之后合理设置block大小以提高occupancy等。

笔试里一张经典题目就是矩阵乘法。朴素的三重循环,每个线程算一个输出元素,每个输出元素要读一整行和一整列,数据完全没有复用,访存效率极低。优化版就是分块,让一个block算一个输出tile,先把需要的子矩阵加载到shared memory,然后反复使用。再进一步,还可以在寄存器层面做微调,让每条数据加载被更多输出元素复用。这个演进过程本身就是一道很好的笔试论述题。

3. 典型题目实战拆解:从朴素实现到优化实现

3.1 矩阵乘法:GPU优化的"Hello World"

如果要选一道必然值得提前演练的题,我首推矩阵乘法。它不是是笔试的原题,但它的优化过程几乎涵盖了GPU优化所有基础知识点。从朴素到最优,我认为可以理出四条台阶。

第一台阶:朴素实现。每个线程算一个输出元素,假设矩阵是MNK的C=A*B,每个线程用两层循环累加K次。这段代码写得出来很容易,但性能极差,原因有两个:一是每个线程要读A的一行和B的一列,数据完全没复用;二是所有访存都是随机离散的,完全不连续。

第二台阶:全局内存连续化。调整线程到输出的映射方式,让同一个warp的线程访问的地址尽量连续。比如让线程按列方向映射到C矩阵时,把线程循环顺序调整一下,保证相邻线程处理相邻的列,这样读A时连续。这一步能提升一些性能,但数据复用问题没有解决。

第三台阶:shared memory tiling。这是真正的分水岭。把A和B分别切成大小为BLOCK_SIZE的tile,每个block负责计算C上一个BLOCK_SIZE×BLOCK_SIZE的输出tile。计算前,先把A和B对应的子块加载到shared memory,然后block内所有线程从这个快得多的共享内存里读取数据进行计算,整个kernel结束前不再访问全局内存。这一步之后,对全局内存的访问量从O(N^3)降到了O(N^3/BLOCK_SIZE),性能提升非常明显。

第四台阶:寄存器tiling和向量化。在shared memory的基础上,让每个线程一次计算多个输出元素(比如4×4的tile),这样可以进一步增加数据复用,减少shared memory的访问次数,同时可以用float4这类向量化加载指令来提高访存效率。还能配合double buffering等技术让数据加载和计算重叠。

笔试里如果考矩阵乘法,通常考到第三台阶就能拿不错的分数,但能说清楚第四台阶思路的人,会让面试官明显多看你一眼。关键是每一层优化,你都得能说清楚:这步优化解决了什么问题,代价是什么。比如shared memory tiling的代价是增加了__syncthreads()同步,block尺寸受shared memory容量限制。

3.2 Occupancy计算题:不说具体数字也能答好

笔试里还有一种我很喜欢看到、但很多考生头疼的题——给你一张GPU的参数表,让你计算某个kernel能达到的最大occupancy是多少。这类题真的考计算能力吗?其实考的是你对资源约束的理解。

Occupancy的概念,简单说就是当前GPU上活跃的warp数量与硬件最大支持warp数量的比值。约束条件主要有四个:每个线程用多少寄存器、每个block用多少shared memory、每个SM最多支持多少线程/block/warp、每个block最多多少线程。

举个例子,假设一个GPU的SM最大支持2048个线程、32个block、64K shared memory。你的kernel每个block用了256个线程,每个线程用了40个寄存器,每个block用了16KB shared memory。那么:按线程数算,一个SM最多放2048/256=8个block;按寄存器算,每个SM寄存器总量假设65536个,每个block需要256*40=10240个寄存器,65536/10240=6.4,只能放6个block;按shared memory算,64K/16K=4个block。三者取最小,最多放4个block,也就是1024个线程,occupancy只有50%。

这种题的精髓是"木桶效应"——取所有约束的最小值。有的同学背了公式,但忘了shared memory也是约束,或者忘了寄存器数量也会限制,结果算错。我建议准备笔试时,不要只背公式,要真正理解每一个资源是怎么被占用的。理解了,题目怎么变都不怕。

还有一点值得提醒:occupancy不是越高越好,这是个很经典的"看似有道理但实际不对"的论点。occupancy高意味着调度的余地大,有利于隐藏访存延迟,但高occupancy通常伴随每个线程可用资源减少(比如寄存器变少),可能导致寄存器溢出(spill)到local memory,反而性能下降。笔试里如果遇到"是不是occupancy越高越好"这类开放题,能说出这个"资源权衡"视角,很加分。

3.3 复杂度估算题:怎么判断算得完还是算不够

另一类高频计算题是:给你一个任务,问你它的理论下限时间是多少,或者判断它是compute-bound还是memory-bound。

这类题的思路其实很固定。第一步,算出这个任务需要多少次浮点运算(FLOPs)。第二步,根据GPU的算力(比如FP32峰值10 TFLOPS)算出纯计算需要的时间。第三步,估算数据量,根据显存带宽(比如900 GB/s)算出数据搬移需要的时间。第四步,比较这两个时间,谁大谁是瓶颈,瓶颈的时间就是理论下限。

比如一个任务需要1G FLOPs的浮点运算,数据量是2GB。按10 TFLOPS算力、900GB/s带宽算,计算时间是0.1秒,访存时间是2.2毫秒,显然这是一个compute-bound的任务,理论下限约0.1秒。反过来,如果任务只有10M FLOPs但要搬运2GB数据,那访存至少2.2毫秒,计算只需1毫秒,这时就是memory-bound,再怎么优化计算也没用,重心应该放到减少数据访问上。

这类题看起来是数学题,实际考察的是你在真实优化中的第一反应。拿到一个性能问题,先定位是哪一类,才能选择正确的优化方向。很多实际优化做不下去,就是因为方向错了——该减少数据搬运的时候在那儿死抠指令数,完全是白费劲。

4. 笔试中的常见陷阱与解题策略

4.1 最容易翻车的几个点

根据我看到的笔试复盘和给候选人做模拟题的经验,有几类错误几乎是批量出现的,分享出来帮大家避坑。

第一个坑是只讲API不讲原理。面试官问"怎么用shared memory",有同学答"用__shared__关键字声明就行"。这个回答看似没错,但没有说到点上。更好的回答应该包括:为什么需要shared memory(因为它比全局内存快很多)、什么场景适合用(数据会被block内多个线程重复访问)、使用时注意什么(同步、bank conflict、容量限制)。笔试不是让你背字典,而是要展示你理解了这个机制存在的意义。

第二个坑是忽略精度问题。GPU优化题里,很多同学把所有注意力放在速度上,却忽略了精度。实际工程师如果只优化速度而不管精度,可能在业务上线时直接出事。笔试中但凡涉及代码或者优化方案的题,主动提一句"需要注意浮点累加的顺序会影响精度,可以用Kahan求和等方式降误差",都会显得你想得很全面。

第三个坑是不会用反例思考。给你一道优化题,你把自己能想到的优化手段全部堆上去,最后代码又复杂又难维护,但你没有任何一个地方验证过到底是哪一步起的作用。专业的优化工程师做事,是"每做一步优化,就测一次性能,确认提升来自这一步"。笔试时就算不需要真跑代码,也要有这个思路,方案里体现出"先定位瓶颈,再针对性优化,再验证"的流程,而不是一股脑堆砌手段。

4.2 性能分析工具链的准备

笔试不考工具使用,但准备工作你一定绕不开工具,因为面试环节的追问往往从这里展开。我建议把NVIDIA的性能分析工具至少用过一遍,不用精通,但要清楚它们各自能干什么,以及它们的定位差异。

nvprof是早期的命令行性能分析工具,现在已经逐步被Nsight系列取代了,但很多网上教程还是用它,最好了解。Nsight Systems(nsys)主要做时间线的宏观分析,看kernel、显存拷贝、CuDNN调用这些事件在时间线上的分布,适合发现整体流程层面的浪费,比如host和device频繁同步、kernel之间有空隙。Nsight Compute(ncu)则是kernel级别的微观分析工具,可以看occupancy、访存吞吐、指令混合、bank conflict这些硬指标,是最能回答"我的kernel到底慢在哪"的工具。

还有一个最朴素却最常用的工具,nvidia-smi。它提供GPU的实时占用率、显存占用、温度、功耗这些信息,虽然颗粒度比较粗,但在判断程序是否真的跑在GPU上、显存是否爆了这些问题上,它是第一排查工具。

准备笔试的同学可以花一个下午把一个简单的kernel用ncu分析一遍,看看里面的指标怎么读。我不需要你记住每个指标的缩写,而是通过这个过程理解一个kernel的完整性能画像到底由哪些信息构成。面试时你能说出"我用Ncu看过我的kernel,发现它的SM busy低,是因为访存等待过高",这句话立刻跟只会写代码的人拉开差距。

4.3 项目经验该怎么准备

笔试里经常会有论述题或者方案设计题,这时候如果你有实打实的项目经历支撑,答题会顺手很多。但很多应届生的问题是:我没什么GPU优化方向的正式实习,怎么办?

我的建议是,个人小项目也能证明能力,关键在于选题要有代表性。典型的可做项目包括:自己用CUDA实现一个常用算子,比如卷积、矩阵乘、LayerNorm,然后优化到接近官方库的性能;用Nsight对现成的一个模型推理过程做性能分析,找出瓶颈并尝试优化;优化一个访存密集的简单算子,比如transpose、reduce,从带宽角度去逼近理论极限。

这些项目放在简历上,写清楚几件事:你做了哪个算子,基线性能是多少,优化之后是多少,怎么找到瓶颈的,用了什么手段。这里面最有说服力的是"优化的量化过程"——你给出从baseline到最终版本的性能曲线,每一步对应什么优化手段。我在面试时最怕听到的就是"我做过优化",但问怎么做的、数据是多少,就支支吾吾。反过来说,能清楚讲出"我把矩阵乘从40GFLOPS优化到120GFLOPS,瓶颈第一是访存不连续,我改了数据布局;第二是共享内存bank conflict,我加了padding;第三是寄存器溢出,我调整了block大小",这种回答,没有正式实习经历也完全能打动面试官。

5. 从笔试看这个岗位的长线价值

复盘下来,商汤2018这次GPU优化工程师笔试,其实代表的是一类岗位的考察逻辑:不追求你会多少花哨的框架,而是看你有没有扎扎实实理解清楚"硬件—数据—算法"这三层之间是怎么耦合的。

我给准备这个方向的同学一个复习顺序建议:先把CUDA编程模型彻底吃透,然后找一块真实的GPU(哪怕是租的云服务器),把矩阵乘法从朴素到tiling的优化路径完整走一遍,每做一步都用Nsight测一下,记录数据;然后去研究一下卷积层是怎么用GPU实现的,im2col和Winograd的思维可以了解一下;最后把访存优化、计算优化、Occupancy计算、分支发散这几个主题整理成自己的笔记。这套流程走完,应付笔试和面试都是够用的。

这些年我感受到一个变化:GPU优化的入门门槛在降低,工具越来越智能,框架帮你封装的东西越来越多,但这恰恰意味着真正能理解底层的人变得更稀缺了。当工具本身不能区分高下的时候,理解原理的人才是在效率、成本和稳定性上有决定性优势的那一批。如果你正在准备这个方向的校招,不要怕笔试题难,它其实是在帮你筛选一个值得投入的赛道。把基础打牢,把实践做扎实,这个岗位能带给你的成长空间,远不止一份offer而已。

需要专业的网站建设服务?

联系我们获取免费的网站建设咨询和方案报价,让我们帮助您实现业务目标

立即咨询