商汤科技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 Computencu看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 / 4K1024时也不超过0.5 FLOP/Byte完全受内存带宽限制。一个经典的优化版本是每个block负责计算一个TILE x TILE的输出子矩阵先把A和B对应tile加载到shared memory再在tile内做乘加。以TILE16为例每个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.xy同理然后马上判断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原因。工具/资料用途说明优先级NVIDIA CUDA C Programming Guide官方编程手册知识树的权威来源高Nsight Computencu分析kernel的瓶颈指标定位stall原因高CUB / Cutlass 源码学习工业级归约和GEMM实现思路中GTC相关演讲视频了解前沿优化技巧和硬件设计思想中自己实现GEMM并对比cuBLAS把优化从0带到80%的最佳练习极高我不能说这些资料“看完就稳了”因为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优化领域是最真实的写照。