商汤科技GPU优化工程师笔试复盘：CUDA核心考点与备考路线

2018年秋天，我参加了商汤科技校招的GPU优化工程师第一场笔试。那年头“GPU优化”还不像今天这样被频繁提起，但商汤作为AI视觉领域的第一梯队公司，专门为这个岗位单独出题，本身就释放了一个信号：AI公司开始认真对待底层性能了。深度学习训练和推理都被卡在GPU利用率上，谁能把kernel跑快，谁就能把迭代周期缩短，把推理成本打下来。我当时准备的方向还是传统的高性能计算，做矩阵运算优化，考完之后复盘了一遍，发现这套笔试非常成体系，既考CUDA基础概念，也考手动计算和代码分析，几乎把GPU优化工程师的核心能力圈画清楚了。

这份复盘不是回忆那套题的具体答案（毕竟年代久远，也承诺过不泄题），而是把笔试考察的知识结构、解题思路、备考方向完整拆解出来。无论你是准备面试GPU优化、高性能计算、推理引擎开发，还是单纯想系统入门CUDA优化，我都建议花十分钟把这篇看完。整篇分四个部分：笔试考察逻辑、核心考点拆解、现场答题节奏、复盘后的备考路线。最后还会附上我自己踩过的坑和工具推荐，尽量做到直接可参考。

1. 笔试考察什么：一套岗位能力模型的拆解

1.1 为什么第一题先考CUDA线程模型

笔试第一道大题通常是CUDA线程模型相关的内容，比如线程块怎么划分、线程ID如何计算、为什么blockDim和gridDim要分开设计。这不是随便出的送分题，线程模型是GPU编程的地基。你如果连线程索引转换都搞不清楚，后面写kernel基本就是瞎写，连数据对齐都可能出问题。

我当时拿到这类题，通常先画一张图：grid是外层网格，block是中间结构，thread是最小执行单元。一维情况下threadId = blockIdx.x * blockDim.x + threadIdx.x，这个公式必须形成肌肉记忆。二维甚至三维索引时，很多新手会算错偏移，笔试里经常出现一个二维grid加二维block的题目，让你把线性索引算出来，这是典型的考察点。

为什么公司要考这个？因为GPU优化本质上就是调度逻辑的优化。你把线程排布想清楚了，内存访问模式大概率不会跑偏；线程排布想不清楚，后面tile切分、shared memory复用、bank conflict规避全是空中楼阁。线程模型不只是语法，它直接影响占用率、并行度和访存行为，是后续一切优化的入口。

1.2 出题的三层结构：概念、计算、代码

商汤这套笔试整体上呈现出明显的三层结构。第一层是概念题，考察CUDA 内存体系、线程层次、流与事件、同步机制等基础概念，属于“背过就能答”的层次。第二层是计算题，比如给一个GPU型号和kernel配置，手算占用率、理论带宽、浮点峰值，这个层次就需要真正理解硬件参数的含义。第三层是代码分析题，贴一段有问题的kernel，让你指出性能瓶颈并给出改进方案，这层考察的是实际优化经验。

这种三层结构其实映射了GPU优化工程师的日常工作方式：先理解硬件模型，再估算理论性能上限，最后通过分析工具定位实际瓶颈。笔试用三小时模拟了这个完整链路。所以我建议准备这类笔试时不要只刷概念，要动手写代码，用Nsight Compute（ncu）看Metrics报告，把“为什么快”“为什么慢”练成直觉。纸上谈兵到这里是过不了关的。

1.3 笔试内容背后隐藏的岗位价值观

再往深一层看，这套笔试还隐藏着一个岗位价值观——用数据说话。面试官想知道你是否具备“先计算理论峰值，再对比实测数据，最后定位瓶颈”的思维习惯。比如给你一个数组求和任务，你不仅要知道用归约算法，还要知道在V100上理论带宽是900GB/s，如果你写出来的kernel有效带宽只有200GB/s，说明存在严重的访存问题。

我当时在答题时特意把计算过程写在草稿上，卷面答案只写关键数字和结论。批卷人看的不只是答案，更看你的思路和工程判断。比如“该kernel受带宽限制”和“该kernel受延迟限制”是完全不同的优化方向，你要能说出依据。这种决策能力，才是公司真正考察的核心。

2. 核心考点逐个攻破：线程、内存、计算三个维度

2.1 内存层级与访问模式的经典陷阱题

GPU的内存体系是笔试的重灾区。寄存器（Register）、shared memory、global memory、local memory、constant memory、texture memory，每种内存的速度和用法必须了然于胸。笔试最常见的陷阱是local memory。很多教材说local memory是“线程私有”，新手误以为它像寄存器一样快，实际上local memory物理上位于global memory，只是因为被L1缓存命中，性能才比直接访问global好一些。考这个点就是看你是否真的理解硬件。

• 寄存器：每个线程私有，最快，但容量有限，超限后会溢出到local memory。
• shared memory：同一个block内共享，速度接近寄存器，但需要处理bank conflict。
• global memory：所有线程共享，延迟几百周期，必须靠合并访问和缓存提升效率。
• constant memory：适合所有线程读同一个地址的场景，有专用缓存。
• texture memory：适合二维空间局部性强的访问模式，图像处理里常用。

笔试中还经常出现“判断以下访问是否合并”的题目。比如一个float数组，线程t访问arr[tid * 32]，这显然是非合并访问，每个warp的32个线程访问的地址间隔128字节，一个内存事务只能采到1/32的数据，访存效率惨不忍睹。改成arr[tid]就对了，相邻线程访问相邻地址，硬件能将这些请求合并成少量事务。这个原理我在后面分析题里还会用。

2.2 手算占用率与带宽：笔试里的硬功夫

计算题往往是拉开差距的地方。先给出一个典型例子：在一块Tesla V100上，每个SM最多2048个线程，一个block大小设为256线程，那么理论上一个SM最多驻留2048 / 256 = 8个block。但如果kernel每个线程用到了40个寄存器，每个SM的寄存器文件总共65536个，那么每个block需要256 * 40 = 10240个寄存器，65536 / 10240 = 6.4，取整就是6个block。实际占用率是6 * 256 / 2048 = 75%。

另一个高频计算题是带宽使用率。给一个kernel，读入8GB数据，写出8GB结果，在某块GPU上实测耗时30ms，那么有效带宽就是16GB / 0.03s ≈ 533GB/s。如果这块GPU的HBM理论带宽是900GB/s，效率约为59%。这个数字并不理想，需要排查是否发生了非合并访问或者访存指令过少导致的延迟暴露。笔试里列出这种计算步骤，比你直接写结论要有说服力得多。

2.3 算子优化分析：GEMM与归约是必考题

算子优化类题目中最常出现的就是矩阵乘法GEMM和归约Reduction。GEMM几乎是深度学习底层最核心的算子，笔试要么给一段朴素实现让你分析性能问题，要么让你写一个使用shared memory的tile版本。

朴素的GEMM实现通常是三重循环，每个线程算C矩阵的一个元素。问题在于内层循环里，对B矩阵的列访问是非合并的，同时对A矩阵的行访问连续但每个元素只使用一次，数据复用率极低。整个kernel的算术强度约等于2K / 4，K=1024时也不超过0.5 FLOP/Byte，完全受内存带宽限制。

一个经典的优化版本是：每个block负责计算一个TILE x TILE的输出子矩阵，先把A和B对应tile加载到shared memory，再在tile内做乘加。以TILE=16为例，每个block只需要从全局内存加载16 x K + K x 16个元素，而原始的乘积需要16 x 16 x K次全局内存访问，访存量下降约TILE/2 = 8倍。这就是tile切分的核心价值：提高数据复用率，把访存密集型计算转化为计算密集型。

归约（Reduction）又是另一种套路。笔试让写一个数组求和kernel，最容易犯的错误是直接让每个线程对全局内存做原子操作，性能被锁竞争拖垮。正确做法是两个阶段：先让每个block内的线程把数据归约到shared memory，块内用树形归约；再让每个block输出一个部分和，由少量线程做最终归约。注意归约过程中要避免divergence，比如线程ID大于当前步长时直接退出，还有shared memory的bank冲突问题，必要时做padding。

3. 笔试实操与答题节奏：现场场景还原

3.1 拿到卷子先做什么：时间分配策略

笔试三小时，题量中等，但代码分析题很耗时间。我拿到卷子第一件事不是从头开始做，而是花五分钟快速翻完全卷，按题型标注难度预估时间。概念题每题控制在5分钟以内，计算题每题15分钟，手写kernel题每题25分钟，代码分析题每题20分钟。最后至少留20分钟检查边界条件和索引计算错误。

这种时间分配方式来自我平时做性能优化的习惯：先全局审题再动手。如果一开始就卡在某个计算题上，把时间耗掉，后面手写kernel题目必然仓促。我记得那场笔试题里后面有一道多线程优化题，需要写完整kernel并解释优化点，分值最高。我给自己留足了时间，最后还有空把几道概念题的表述改得更严谨。时间管理本来就是工程能力的一部分。

3.2 手写kernel时的规范细节

手写kernel题目的评分不仅看功能对不对，还看代码风格和边界处理。我总结了一套自己的答题模板，建议你也准备一套。首先是索引初始化，如果是二维场景，一定要写清x = blockIdx.x * blockDim.x + threadIdx.x，y同理，然后马上判断x < width && y < height，防止索引越界。

其次是循环设计。不要假设线程数恰好等于数据量，用grid-stride loop更稳妥：for (int i = begin; i < n; i += stride)，其中stride = gridDim.x * blockDim.x。这样即使数据量远大于启动的线程数，也能保证每个线程处理多个元素，而且循环间隔是连续的，访存可以保持合并。

最后是同步问题。多个线程同时写同一个全局地址时，要么分阶段归约，要么用atomic操作。特别要注意的是，shared memory写完之后必须调用__syncthreads()，否则另一个block内的线程读到的是脏数据。笔试中写出同步缺失的代码，面试官一眼就能看出来你踩过坑。

3.3 一道典型分析题的完整推导过程

有一道题让我印象很深，针对一段矩阵转置kernel要求指出瓶颈并优化。代码逻辑大概是：每个线程读A[row][col]，然后写到B[col][row]。直观看起来功能正确，但性能极差。原因在于写B时，同一warp内线程的col是相邻的，写入地址B[col][row]的步长是row stride，也就是矩阵宽度，这完全破坏了合并访问。读A是合并的，但写B不合并，整个kernel的有效带宽被拉低到原来的几十分之一。

我的优化方案是使用shared memory分块转置。每个block加载一个TILE x TILE的A子矩阵到shared memory，加载时保持合并访问；然后同步；再从shared memory读取并转置写入B，写入时同样保持合并访问。这样全局内存的读写都是合并的，shared memory里可以任意转置访问，代价只是增加一次同步和shared memory的bank conflict处理。

这一类分析题的解题逻辑是固定的：先看访存模式是否合并，再看数据复用率，再看是否发生bank conflict，最后看有无同步错误。把这条检查链背下来，放在答题里，比凭感觉写“性能不好”要专业得多。

4. 复盘后的备考行动路线

4.1 从笔试反推需要的知识树

笔试之后我在笔记本上画了一张知识树，把它作为后续几个月的备考索引。最底层是硬件架构，包括SM结构、内存层级、warp调度机制；第二层是CUDA编程模型，包括线程组织、同步、流与事件；第三层是性能分析方法，包括ncu指标阅读、带宽与延迟区分、roofline模型；第四层是算子库和框架源码，包括CUB、Cutlass、cuBLAS的常用优化技巧。

这棵树的好处是帮你快速定位知识空白。比如发现自己shared memory只懂概念不会计算bank conflict，就去找资料专门练；发现自己对缓存命中率没有直觉，就用ncu在真实kernel上反复测量。商业公司招优化的工程师，最看重的就是你能不能把理论和实测对应起来，知识树只是一个框架，真正的延伸需要靠项目实践来填。

4.2 工具链与资料清单

备考必须上手工具。我自己的环境是一张普通消费级卡（显卡算力不算高，但足够跑通绝大多数优化实验），系统装的Ubuntu双系统，CUDA Toolkit版本建议用和实际部署环境接近的版本。我的实测经验是，nvidia-smi只是最基础的监控，真正要熟练使用的是Nsight Compute，它能给你完整的性能分析报告，包括内存吞吐、计算吞吐、stall原因。

我不能说这些资料“看完就稳了”，因为GPU优化本身就是一门实践学科。我的建议是每个资料都要配一个实验：读完某个优化技巧，就把它用到自己写的kernel里，用ncu记录优化前后的指标差异。比如读完shared memory tile优化，就动手写一个16x16或32x32的GEMM版本，看看有效带宽从多少提升到多少。这样学到的知识才是自己的。

4.3 备考路上踩过的坑和提醒

最后分享几个我实际踩过的坑。第一个坑是只关注理论复杂度，忽略访存模式。我刚开始优化一个卷积算子，把计算量从O(n^2)降到O(n)，但实际性能反而下降，因为我用了大量非连续访存，反而让内存子系统成了瓶颈。优化之前一定要先判断算子是计算密集还是访存密集，用roofline模型定位瓶颈面。

第二个坑是过度依赖atomic操作。在一个归约任务里，我图省事让每个线程直接atomicAdd到全局变量，结果因为竞争严重，20个线程版本比1个线程还慢。分布式归约思想在GPU上同样适用，必须分block归约再合并。笔试里这个点也是高频，答错的人很多。

第三个坑是忽略边界和整除条件。手写kernel时如果没有grid-stride loop或者边界判断，数据量不是线程数整倍数时，线程会越界访问，在线上就是随机崩溃。我建议在本地用AddressSanitizer或者cuda-memcheck跑一遍，再提交。

还有一个容易被忽视的点是版本问题。2018年的V100和今天的H100在架构细节上有很大差异，但内存模型、线程模型的基本原理没有变。所以当你翻开官方文档时，注意区分哪些是架构相关参数（比如SM数量、共享内存大小），哪些是通用编程模型原理。笔试和面试中你可以在答案里标注“根据硬件的具体参数调整”，这体现的是工程思考方式。

这次备战让我养成了一个习惯：每次写完kernel，都要问自己“理论峰值是多少，实测达到多少，差距在哪里”。这个问题听起来简单，但能连续问下去的人不多。GPU优化没有银弹，无非是不断逼近硬件极限的过程。如果你也在准备类似的岗位笔试，建议从今天开始就打开Nsight，找一个最简单的kernel，跑一遍，看看那些Metrics到底在说什么。纸上得来终觉浅，这句老话在GPU优化领域是最真实的写照。