2018年秋天我投了商汤科技校招的GPU优化工程师岗位简历通过后收到了笔试通知。第二场笔试当天机房一进去就感觉气氛不太一样没有普通算法岗那种刷不完的LeetCode发下来的卷子更像是“CUDA实战性能分析”的混合体题目乍看都能下手细看处处是坑。今天把这场考试的复盘整理出来结合这几年做算子优化和推理引擎踩过的坑把考了什么、为什么这么考、怎么准备说清楚。这篇内容适合正在投AI公司高性能计算方向的应届生也适合刚接触CUDA、想系统建立GPU优化思维的同学。1. 先搞清楚这个岗位到底在考什么1.1 题型分布与考察维度先说结论第二场笔试的核心不是考你会不会调参而是考你有没有一套稳定的GPU性能优化方法论。我当时拿到的卷子大概分三个大板块基础概念、手写代码、策略分析。基础概念题围绕CUDA线程模型、内存布局、同步机制、编译器优化入门。这类题分值不高但答不好很难看因为它决定了你能不能和团队用同一套语言沟通。手写代码与优化题一般是两到三道要求你把一个常见算子从“能跑”改到“尽量快”。常见载体是归约、矩阵乘法、图像处理、逐元素变换这类东西。综合策略题给你一个业务场景问你如何优化某个环节或者让你估算性能瓶颈在哪里。这类题虽然不写代码但最考验对真实系统的理解。从分值占比看概念题大概三成代码与优化题五成策略分析两成。这个比例代表一个信号GPU优化工程师校招的筛选重点不是背了多少理论而是能不能在限时里写出一段可以编译、可以跑、还知道往哪个方向改的代码。阅卷人不会只盯着你的最终答案更看重你的思考路径里有没有“先量级判断、再逐层优化”的职业习惯。这类笔试设计其实映射了真实岗位的内容。商汤的业务大量以计算机视觉为主从图像分类到目标检测、分割模型从训练到上线每个环节都依赖GPU算子。哪怕只是一个图像缩放、一个归一化写不好也要吃掉不少耗时。GPU优化工程师日常就是和这些算子较劲所以笔试自然要把“能不能识别瓶颈、会不会改算子”放在最前面。1.2 这些考点背后的真实业务场景为什么第二场笔试要考矩阵乘法和归约因为它们是深度学习底层算子库的骨架。乘法不仅仅是线性代数教材里的公式它在卷积、全连接、attention这些结构里反复出现归约操作是softmax、batch norm、layer norm等算子最基本的一环。考题以它们为载体实际上是在筛选“能不能理解深度学习推理引擎里到底发生了什么”的人。我记得当时有一道综合题是在线视频流里需要同时对多路视频做缩放、颜色空间转换和归一化再送入检测模型请设计方案。这就比单纯写一个kernel更接近真实业务。你要考虑数据是NCHW还是NHWC要不要用GPU做预处理多路视频怎么调度会不会和推理争抢显存甚至要考虑CPU和GPU之间的拷贝次数。这类题没有唯一标准答案但阅卷人能很明显地看出谁在拿模板套话、谁真的处理过类似系统。这也解释了一个现象很多同学单独背优化点时头头是道一到综合题就抓瞎。因为真实场景里的优化从来不是单点技巧而是多个约束条件之间的权衡。笔试故意把业务背景摆出来就是要筛掉只会背口诀、不会落地的人。2. 高频考点拆解体系结构与并行编程2.1 线程与内存模型笔试里最容易被“会一点”坑死的部分GPU编程里线程模型和内存模型不是两个独立知识点它们是同一件事的两面。笔试里很多题看似考语法实际考你有没有从硬件角度想问题。先说线程模型。CUDA把线程组织成线程组成线程块线程块组成网格。一个线程块内的线程可以同步、可以通过共享内存通信不同线程块之间基本没有有序协作手段。硬件调度时会把线程块切分成warp通常32个线程一起执行。这意味着你的代码里一个分支如果让一个warp内的线程走向不同路径warp就要串行执行所有路径计算能力被白白浪费。所以写kernel时尽量让warp内所有线程行为一致这是笔试里反复出现的隐藏考点。内存模型更需要重视。GPU的内存体系大致是寄存器、共享内存、全局内存、常量内存、纹理内存。越靠近计算单元越贵越快容量越小。下面这个表格是我在多次面试和笔试中反复用到的速查表内存类型位置访问延迟容量生命周期寄存器每个线程私有最低基本零延迟每线程几十到上百个线程内共享内存每个block私有低几十周期早期默认最多48KBblock内全局内存所有线程共享最高数百周期显存容量可到几十GB整个程序常量内存只读广播缓存命中时较低64KB整个程序纹理内存只读有缓存依访问模式而定显存容量整个程序笔试常问为什么共享内存这么重要因为全局内存延迟太高一个线程直接读全局内存可能要等几百个周期。如果同一块数据会被多个线程反复使用把它先搬到共享内存让block内所有线程各自命中就能把远距离读变成近距离读。这也是矩阵分块优化的核心动机。另一个常考点是同步。__syncthreads()的作用是让block内线程在屏障前等待彼此确保共享内存写入后可以被其他线程安全读取。笔试里最常见的错误是一个线程写共享内存另一个线程立刻去读中间没加同步结果数据是脏的。反过来如果所有线程都调用了__syncthreads()但某个分支提前退出也可能卡死这就是死锁风险。2.2 访存合并与bank conflict两道回忆真题2018年那场笔试里有两道题我现在还记得很清楚。第一道是概念题假设kernel的线程ID为tid访问全局数组a请问a[tid]和a[tid * 2]哪个效率更高为什么第二道是代码分析题共享内存数组s[32][32]线程按s[tid][tid]方式访问会发生什么现象这两道题考的本质都是访存模式。先看全局内存。GPU从全局内存读数据时并不是一个线程单独读一个字节而是把整个warp的请求合并成尽量少的内存事务。如果warp内32个线程访问的地址是连续的硬件一次就能把所有数据取回来这叫合并访问。如果像a[tid * 2]这样跳着读地址不连续硬件就要拆分成多个事务带宽利用率明显下降。所以优化第一原则是让同一个warp内的线程访问相邻地址。再看共享内存。共享内存虽然快但被划分成32个bank同一时刻一个bank只能响应一个访问请求。如果warp内多个线程访问不同bank硬件可以并行处理如果多个线程访问同一个bank的不同地址就会发生bank conflict访问被串行化。经典坑位是共享内存数组的“对角线访问”不一定产生冲突但像s[tid][4 * tid]这种访问要提前看列索引是不是落在同一个bank内。笔试遇到这类题建议先画一张32×32的bank映射表别凭感觉猜。类似的题还有分支发散一个warp里有线程走if分支、有线程走else分支怎么办答案不是“尽量避免if”而是“尽量让warp内线程走向一致”必要时可以用条件计算代替分支。这些细节单拎出来都不难但放在笔试里就是用来测试你有没有真正理解GPU的执行模型。3. 手写代码题从能跑通到能拿高分3.1 向量归约一道题串起三个优化层次归约是GPU笔试里出现率最高的题目没有之一。它看起来太简单了给一个数组求出所有元素的和。但正因为简单才能看出一个人有没有优化直觉。先写一个最基础的版本思路__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; sdata[tid] (i n) ? in[i] : 0.0f; __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只处理256个元素如果数组有100万个数就要开几千个block最后每个block都调一次atomicAdd原子操作竞争会很严重。第二归约过程中从s128开始每个步骤只有一半线程在工作利用率越来越低。第三每次循环内部都有一整次__syncthreads()同步次数多调度开销大。改良思路通常分三层。第一层先让每个线程用strided loop处理多个元素降低block数量减小原子竞争。第二层把块内归约改成交错式让线程在归约过程中尽量保持活跃比如让前一半线程连续读取共享内存的前半部分与后半部分相加。第三层用warp shuffle指令让32个线程在warp内部完成归约完全不经过共享内存同步次数大幅减少。这里要注意不同架构对shuffle的支持有差异笔试提到“利用warp内部指令”通常是很加分的回答。很多同学笔试时只写了最基础的版本就觉得完事了。实际上在真实场景里百亿级元素的归约不同版本性能可能差出一个数量级。阅卷人看的就是你在优化链上能走多远能写出正确版本说明基础扎实能说到共享内存说明有优化意识能提到warp shuffle说明你真的研究过GPU指令集。3.2 矩阵乘法典型算子的完整优化链矩阵乘法在GPU优化笔试里几乎算是必考题因为它就是深度学习的骨架算子。很多同学第一反应是写最朴素的版本每个线程算输出矩阵C的一个元素内层循环逐个读A的行和B的列。这个版本逻辑上没错但性能一言难尽——内层循环每算一次都要读全局内存相当于把A和B反复从显存里搬带宽根本扛不住。改进的第一个节点是分块。假设用16×16的线程块每个线程负责一个输出元素块内只需要把A的一个16×16的tile和B的一个16×16的tile搬到共享内存然后重复利用。计算访存比瞬间上升性能通常能提升好几倍。如果笔试只要求写优化思路讲到这一步已经及格。第二个节点是处理bank conflict。共享内存里的B tile如果按列方向读取容易产生冲突常见做法是声明时给行补一个偏移比如用16×17而不是16×16存储列访问时bank错开冲突就消掉了。第三个节点是向量化和展开。让一个线程连续处理4个输出元素用float4一次读4个浮点可以减少指令数、提高内存带宽利用率。再配合循环展开编译器能更好调度寄存器。我当年笔试时写了一个简化版的tile内核核心流程如下__global__ void matmul_tile(const float* A, const float* B, float* C, int M, int N, int K) { __shared__ float As[16][16]; __shared__ float Bs[16][16]; int row blockIdx.y * 16 threadIdx.y; int col blockIdx.x * 16 threadIdx.x; float sum 0.0f; for (int k0 0; k0 K; k0 16) { As[threadIdx.y][threadIdx.x] A[row * K k0 threadIdx.x]; Bs[threadIdx.y][threadIdx.x] B[(k0 threadIdx.y) * N col]; __syncthreads(); #pragma unroll for (int k 0; k 16; k) { sum As[threadIdx.y][k] * Bs[k][threadIdx.x]; } __syncthreads(); } C[row * N col] sum; }这个版本边界没处理笔试里足够展示思路。真正工作中还要处理M、N、K不能被tile整除、使用双缓冲隐藏加载延迟、甚至做寄存器级别的精细调度。多层优化后2048×2048规模的计算性能差异可以非常明显优化阶段主要改动相对性能参考朴素版本直接访问全局内存1倍16×16分块用共享内存缓存tile约3到5倍避免bank conflict共享内存数组补padding额外提升20%到40%向量化加展开单线程多元素、float4读取可再提升1到2倍笔试中如果你能把这一条链路讲清楚基本可以覆盖大多数阅卷老师的预期。4. 性能优化方法论与“描述优化思路”题的答题框架4.1 先用Roofline判断瓶颈再谈优化笔试进入综合题后经常不是让你写代码而是让你对着一个算子或一个业务场景说优化方案。这时候最容易暴露“只会背优化点不知道怎么切入”的问题。我建议所有准备GPU方向的同学都掌握Roofline模型。它的核心就两个量一个是算力也就是GPU每秒能浮点运算多少次另一个是带宽也就是GPU每秒能从全局内存搬多少字节。把算力除以带宽能得到一个临界点叫算术强度。算术强度 总计算量 / 总访存量。如果一个算子的算术强度低于临界点它是访存受限优化重点应该是减少数据搬运、改善访问模式、提高缓存命中而不是拼命加运算如果算术强度高于临界点它是计算受限优化重点才是提高指令级并行、减少计算浪费、用更高效的数学指令。比如一个逐元素加法的kernel每个元素读两次写一次总共24字节访存只做一次浮点加法算术强度很低显然访存受限。这时候写多漂亮的向量化都没大用瓶颈在带宽。矩阵乘法如果分块足够大算术强度能升到几十甚至上百变成计算受限那才值得把精力放在FMA、寄存器复用上。笔试里遇到“如何优化一个算子”的题先不要急着列优化点先给一个量级判断“这个算子访存量很大、计算量不高我认为瓶颈在内存带宽”。这句话能证明你有系统级思维而不仅仅是背了几条套路。4.2 给一个算子按这七步走基本不会翻车我后来带过一些新人发现大家面对“描述优化思路”题时最大的问题不是不知道优化点而是不知道按什么顺序讲。我总结了一个七步框架笔试和面试都能用第一步明确算子输入输出数据的形状、布局、精度。这一步决定后面的所有分析。第二步估算计算量和访存量粗略算一下这个算子要做多少次浮点运算、访问多少字节数据。量级对了就行。第三步判断瓶颈类型用算术强度或经验判断访存受限还是计算受限。第四步保证并行度检查线程数是否足够、线程块大小是否合理、是否存在负载不均衡。第五步优化访存模式让线程访问连续地址、合并访问、能复用则复用减少全局内存往返。第六步优化计算资源利用使用向量化指令、循环展开、共享内存或寄存器缓存、减少分支发散。第七步用profiler验证跑基线看吞吐率、占用率、kernel时间确认瓶颈是否真如预期。举个例子笔试曾出现“把一段C语言写的图像高斯模糊改成GPU kernel”的题。按这个框架我们会先判断高斯模糊是典型的访存受限算子因为每个像素只做固定次数的乘加但需要反复读邻居像素。所以第一件要事不是优化计算而是让采样相邻像素时尽量命中缓存用共享内存缓存图像块再处理边界。每一步都有明确目的而不是“把数组搬到共享内存”这种无脑优化。这套框架还有一个好处当你不确定题目期望时按顺序展开至少能覆盖大部分得分点不会出现“憋了半天想不出优化点”的情况。5. 笔试环境与常见坑位提前踩过才算稳5.1 上机环境的真实状况笔试不一定给你一个配置好的现代集群。2018年我在考场里用的是一台带NVIDIA显卡的普通机器CUDA版本也不见得新没有IDE运行代码基本靠命令行。写代码要注意几点不要用太新的API特性除非你确定编译器支持文件命名、编译命令要注意题目要求提交前一定要验证能否编译。这个提醒放到今天依然适用。很多同学在本地用一个版本开发到笔试环境突然发现编译器版本老、库路径不对、显存被占用代码跑不起来。笔试环境通常不会给你太多排查时间平时练习就要养成“代码写出来顺手编译一遍”的习惯。5.2 常见运行错误与排查优先级上机笔试最容易翻车的不是思路而是一堆低级运行错误。我整理了一份速查表报错或现象常见原因排查优先级illegal memory access数组越界、空指针、索引算错1invalid argument核函数参数错误、grid或block配置非法1结果不正确或随机波动缺少同步、共享内存未初始化、浮点累加顺序变化2程序卡住不退出__syncthreads()使用不当导致死锁、无限循环2shared memory过大编译失败超过默认共享内存上限或未配置动态共享内存3性能比CPU还慢启动开销太大、访存模式太差、并行度太低3排查时不要慌先看索引和边界条件再看同步最后才怀疑硬件。我发现很多同学一旦报错就怀疑显卡出问题其实绝大多数是自己代码写错了。这里有个特别实用的技巧平时调试kernel时可以先跑一遍compute-sanitizer之类的工具能直接定位越界地址。笔试也许不允许用工具但平时训练一定要养成习惯。5.3 没GPU怎么练本地、WSL与云实例怎么选很多同学会问我笔记本没有独显怎么准备这种笔试这个问题其实好几年前就存在越早解决越好。如果你的机器有NVIDIA独立显卡第一选择是本地环境。Windows下可以装WSL在WSL里安装CUDA工具链注意检查驱动和容器权限。常见情况是WSL里装好驱动后运行程序时遇到无法初始化NVML的提示说GPU访问被操作系统拦截。这种问题大多不是代码问题而是WSL对GPU访问的权限或版本匹配问题要先更新Windows端驱动再确认WSL版本和CUDA支持。如果你的机器是纯核显或者根本没有NVIDIA GPU别在CPU模拟上死磕。CPU上虽然可以编译CUDA代码但没法跑也看不到真实性能和访存行为。最靠谱的做法是租一台带GPU的云主机或按小时收费的GPU实例把练习环境搭好。选镜像时注意驱动和CUDA Toolkit版本匹配装好后先跑一个vectorAdd验证环境再开始经典算子练习。顺便说一句训练GPU编程能力不能只看不写。至少要把归约、矩阵乘法、图像处理这类经典算子每个都手写一遍再用profiling工具看一遍报告对比优化前后差异。这种习惯远比考前刷题更能锻炼真正的优化直觉。6. 笔试之后的一些心里话6.1 一次笔试复盘带来的认知转变写到这里我想起那次笔试结束后我和同考场一个同学在走廊里聊了几句。他说题目不难但每一道都问得很细
GPU优化工程师校招笔试复盘：CUDA性能优化核心考点解析
2018年秋天我投了商汤科技校招的GPU优化工程师岗位简历通过后收到了笔试通知。第二场笔试当天机房一进去就感觉气氛不太一样没有普通算法岗那种刷不完的LeetCode发下来的卷子更像是“CUDA实战性能分析”的混合体题目乍看都能下手细看处处是坑。今天把这场考试的复盘整理出来结合这几年做算子优化和推理引擎踩过的坑把考了什么、为什么这么考、怎么准备说清楚。这篇内容适合正在投AI公司高性能计算方向的应届生也适合刚接触CUDA、想系统建立GPU优化思维的同学。1. 先搞清楚这个岗位到底在考什么1.1 题型分布与考察维度先说结论第二场笔试的核心不是考你会不会调参而是考你有没有一套稳定的GPU性能优化方法论。我当时拿到的卷子大概分三个大板块基础概念、手写代码、策略分析。基础概念题围绕CUDA线程模型、内存布局、同步机制、编译器优化入门。这类题分值不高但答不好很难看因为它决定了你能不能和团队用同一套语言沟通。手写代码与优化题一般是两到三道要求你把一个常见算子从“能跑”改到“尽量快”。常见载体是归约、矩阵乘法、图像处理、逐元素变换这类东西。综合策略题给你一个业务场景问你如何优化某个环节或者让你估算性能瓶颈在哪里。这类题虽然不写代码但最考验对真实系统的理解。从分值占比看概念题大概三成代码与优化题五成策略分析两成。这个比例代表一个信号GPU优化工程师校招的筛选重点不是背了多少理论而是能不能在限时里写出一段可以编译、可以跑、还知道往哪个方向改的代码。阅卷人不会只盯着你的最终答案更看重你的思考路径里有没有“先量级判断、再逐层优化”的职业习惯。这类笔试设计其实映射了真实岗位的内容。商汤的业务大量以计算机视觉为主从图像分类到目标检测、分割模型从训练到上线每个环节都依赖GPU算子。哪怕只是一个图像缩放、一个归一化写不好也要吃掉不少耗时。GPU优化工程师日常就是和这些算子较劲所以笔试自然要把“能不能识别瓶颈、会不会改算子”放在最前面。1.2 这些考点背后的真实业务场景为什么第二场笔试要考矩阵乘法和归约因为它们是深度学习底层算子库的骨架。乘法不仅仅是线性代数教材里的公式它在卷积、全连接、attention这些结构里反复出现归约操作是softmax、batch norm、layer norm等算子最基本的一环。考题以它们为载体实际上是在筛选“能不能理解深度学习推理引擎里到底发生了什么”的人。我记得当时有一道综合题是在线视频流里需要同时对多路视频做缩放、颜色空间转换和归一化再送入检测模型请设计方案。这就比单纯写一个kernel更接近真实业务。你要考虑数据是NCHW还是NHWC要不要用GPU做预处理多路视频怎么调度会不会和推理争抢显存甚至要考虑CPU和GPU之间的拷贝次数。这类题没有唯一标准答案但阅卷人能很明显地看出谁在拿模板套话、谁真的处理过类似系统。这也解释了一个现象很多同学单独背优化点时头头是道一到综合题就抓瞎。因为真实场景里的优化从来不是单点技巧而是多个约束条件之间的权衡。笔试故意把业务背景摆出来就是要筛掉只会背口诀、不会落地的人。2. 高频考点拆解体系结构与并行编程2.1 线程与内存模型笔试里最容易被“会一点”坑死的部分GPU编程里线程模型和内存模型不是两个独立知识点它们是同一件事的两面。笔试里很多题看似考语法实际考你有没有从硬件角度想问题。先说线程模型。CUDA把线程组织成线程组成线程块线程块组成网格。一个线程块内的线程可以同步、可以通过共享内存通信不同线程块之间基本没有有序协作手段。硬件调度时会把线程块切分成warp通常32个线程一起执行。这意味着你的代码里一个分支如果让一个warp内的线程走向不同路径warp就要串行执行所有路径计算能力被白白浪费。所以写kernel时尽量让warp内所有线程行为一致这是笔试里反复出现的隐藏考点。内存模型更需要重视。GPU的内存体系大致是寄存器、共享内存、全局内存、常量内存、纹理内存。越靠近计算单元越贵越快容量越小。下面这个表格是我在多次面试和笔试中反复用到的速查表内存类型位置访问延迟容量生命周期寄存器每个线程私有最低基本零延迟每线程几十到上百个线程内共享内存每个block私有低几十周期早期默认最多48KBblock内全局内存所有线程共享最高数百周期显存容量可到几十GB整个程序常量内存只读广播缓存命中时较低64KB整个程序纹理内存只读有缓存依访问模式而定显存容量整个程序笔试常问为什么共享内存这么重要因为全局内存延迟太高一个线程直接读全局内存可能要等几百个周期。如果同一块数据会被多个线程反复使用把它先搬到共享内存让block内所有线程各自命中就能把远距离读变成近距离读。这也是矩阵分块优化的核心动机。另一个常考点是同步。__syncthreads()的作用是让block内线程在屏障前等待彼此确保共享内存写入后可以被其他线程安全读取。笔试里最常见的错误是一个线程写共享内存另一个线程立刻去读中间没加同步结果数据是脏的。反过来如果所有线程都调用了__syncthreads()但某个分支提前退出也可能卡死这就是死锁风险。2.2 访存合并与bank conflict两道回忆真题2018年那场笔试里有两道题我现在还记得很清楚。第一道是概念题假设kernel的线程ID为tid访问全局数组a请问a[tid]和a[tid * 2]哪个效率更高为什么第二道是代码分析题共享内存数组s[32][32]线程按s[tid][tid]方式访问会发生什么现象这两道题考的本质都是访存模式。先看全局内存。GPU从全局内存读数据时并不是一个线程单独读一个字节而是把整个warp的请求合并成尽量少的内存事务。如果warp内32个线程访问的地址是连续的硬件一次就能把所有数据取回来这叫合并访问。如果像a[tid * 2]这样跳着读地址不连续硬件就要拆分成多个事务带宽利用率明显下降。所以优化第一原则是让同一个warp内的线程访问相邻地址。再看共享内存。共享内存虽然快但被划分成32个bank同一时刻一个bank只能响应一个访问请求。如果warp内多个线程访问不同bank硬件可以并行处理如果多个线程访问同一个bank的不同地址就会发生bank conflict访问被串行化。经典坑位是共享内存数组的“对角线访问”不一定产生冲突但像s[tid][4 * tid]这种访问要提前看列索引是不是落在同一个bank内。笔试遇到这类题建议先画一张32×32的bank映射表别凭感觉猜。类似的题还有分支发散一个warp里有线程走if分支、有线程走else分支怎么办答案不是“尽量避免if”而是“尽量让warp内线程走向一致”必要时可以用条件计算代替分支。这些细节单拎出来都不难但放在笔试里就是用来测试你有没有真正理解GPU的执行模型。3. 手写代码题从能跑通到能拿高分3.1 向量归约一道题串起三个优化层次归约是GPU笔试里出现率最高的题目没有之一。它看起来太简单了给一个数组求出所有元素的和。但正因为简单才能看出一个人有没有优化直觉。先写一个最基础的版本思路__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; sdata[tid] (i n) ? in[i] : 0.0f; __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只处理256个元素如果数组有100万个数就要开几千个block最后每个block都调一次atomicAdd原子操作竞争会很严重。第二归约过程中从s128开始每个步骤只有一半线程在工作利用率越来越低。第三每次循环内部都有一整次__syncthreads()同步次数多调度开销大。改良思路通常分三层。第一层先让每个线程用strided loop处理多个元素降低block数量减小原子竞争。第二层把块内归约改成交错式让线程在归约过程中尽量保持活跃比如让前一半线程连续读取共享内存的前半部分与后半部分相加。第三层用warp shuffle指令让32个线程在warp内部完成归约完全不经过共享内存同步次数大幅减少。这里要注意不同架构对shuffle的支持有差异笔试提到“利用warp内部指令”通常是很加分的回答。很多同学笔试时只写了最基础的版本就觉得完事了。实际上在真实场景里百亿级元素的归约不同版本性能可能差出一个数量级。阅卷人看的就是你在优化链上能走多远能写出正确版本说明基础扎实能说到共享内存说明有优化意识能提到warp shuffle说明你真的研究过GPU指令集。3.2 矩阵乘法典型算子的完整优化链矩阵乘法在GPU优化笔试里几乎算是必考题因为它就是深度学习的骨架算子。很多同学第一反应是写最朴素的版本每个线程算输出矩阵C的一个元素内层循环逐个读A的行和B的列。这个版本逻辑上没错但性能一言难尽——内层循环每算一次都要读全局内存相当于把A和B反复从显存里搬带宽根本扛不住。改进的第一个节点是分块。假设用16×16的线程块每个线程负责一个输出元素块内只需要把A的一个16×16的tile和B的一个16×16的tile搬到共享内存然后重复利用。计算访存比瞬间上升性能通常能提升好几倍。如果笔试只要求写优化思路讲到这一步已经及格。第二个节点是处理bank conflict。共享内存里的B tile如果按列方向读取容易产生冲突常见做法是声明时给行补一个偏移比如用16×17而不是16×16存储列访问时bank错开冲突就消掉了。第三个节点是向量化和展开。让一个线程连续处理4个输出元素用float4一次读4个浮点可以减少指令数、提高内存带宽利用率。再配合循环展开编译器能更好调度寄存器。我当年笔试时写了一个简化版的tile内核核心流程如下__global__ void matmul_tile(const float* A, const float* B, float* C, int M, int N, int K) { __shared__ float As[16][16]; __shared__ float Bs[16][16]; int row blockIdx.y * 16 threadIdx.y; int col blockIdx.x * 16 threadIdx.x; float sum 0.0f; for (int k0 0; k0 K; k0 16) { As[threadIdx.y][threadIdx.x] A[row * K k0 threadIdx.x]; Bs[threadIdx.y][threadIdx.x] B[(k0 threadIdx.y) * N col]; __syncthreads(); #pragma unroll for (int k 0; k 16; k) { sum As[threadIdx.y][k] * Bs[k][threadIdx.x]; } __syncthreads(); } C[row * N col] sum; }这个版本边界没处理笔试里足够展示思路。真正工作中还要处理M、N、K不能被tile整除、使用双缓冲隐藏加载延迟、甚至做寄存器级别的精细调度。多层优化后2048×2048规模的计算性能差异可以非常明显优化阶段主要改动相对性能参考朴素版本直接访问全局内存1倍16×16分块用共享内存缓存tile约3到5倍避免bank conflict共享内存数组补padding额外提升20%到40%向量化加展开单线程多元素、float4读取可再提升1到2倍笔试中如果你能把这一条链路讲清楚基本可以覆盖大多数阅卷老师的预期。4. 性能优化方法论与“描述优化思路”题的答题框架4.1 先用Roofline判断瓶颈再谈优化笔试进入综合题后经常不是让你写代码而是让你对着一个算子或一个业务场景说优化方案。这时候最容易暴露“只会背优化点不知道怎么切入”的问题。我建议所有准备GPU方向的同学都掌握Roofline模型。它的核心就两个量一个是算力也就是GPU每秒能浮点运算多少次另一个是带宽也就是GPU每秒能从全局内存搬多少字节。把算力除以带宽能得到一个临界点叫算术强度。算术强度 总计算量 / 总访存量。如果一个算子的算术强度低于临界点它是访存受限优化重点应该是减少数据搬运、改善访问模式、提高缓存命中而不是拼命加运算如果算术强度高于临界点它是计算受限优化重点才是提高指令级并行、减少计算浪费、用更高效的数学指令。比如一个逐元素加法的kernel每个元素读两次写一次总共24字节访存只做一次浮点加法算术强度很低显然访存受限。这时候写多漂亮的向量化都没大用瓶颈在带宽。矩阵乘法如果分块足够大算术强度能升到几十甚至上百变成计算受限那才值得把精力放在FMA、寄存器复用上。笔试里遇到“如何优化一个算子”的题先不要急着列优化点先给一个量级判断“这个算子访存量很大、计算量不高我认为瓶颈在内存带宽”。这句话能证明你有系统级思维而不仅仅是背了几条套路。4.2 给一个算子按这七步走基本不会翻车我后来带过一些新人发现大家面对“描述优化思路”题时最大的问题不是不知道优化点而是不知道按什么顺序讲。我总结了一个七步框架笔试和面试都能用第一步明确算子输入输出数据的形状、布局、精度。这一步决定后面的所有分析。第二步估算计算量和访存量粗略算一下这个算子要做多少次浮点运算、访问多少字节数据。量级对了就行。第三步判断瓶颈类型用算术强度或经验判断访存受限还是计算受限。第四步保证并行度检查线程数是否足够、线程块大小是否合理、是否存在负载不均衡。第五步优化访存模式让线程访问连续地址、合并访问、能复用则复用减少全局内存往返。第六步优化计算资源利用使用向量化指令、循环展开、共享内存或寄存器缓存、减少分支发散。第七步用profiler验证跑基线看吞吐率、占用率、kernel时间确认瓶颈是否真如预期。举个例子笔试曾出现“把一段C语言写的图像高斯模糊改成GPU kernel”的题。按这个框架我们会先判断高斯模糊是典型的访存受限算子因为每个像素只做固定次数的乘加但需要反复读邻居像素。所以第一件要事不是优化计算而是让采样相邻像素时尽量命中缓存用共享内存缓存图像块再处理边界。每一步都有明确目的而不是“把数组搬到共享内存”这种无脑优化。这套框架还有一个好处当你不确定题目期望时按顺序展开至少能覆盖大部分得分点不会出现“憋了半天想不出优化点”的情况。5. 笔试环境与常见坑位提前踩过才算稳5.1 上机环境的真实状况笔试不一定给你一个配置好的现代集群。2018年我在考场里用的是一台带NVIDIA显卡的普通机器CUDA版本也不见得新没有IDE运行代码基本靠命令行。写代码要注意几点不要用太新的API特性除非你确定编译器支持文件命名、编译命令要注意题目要求提交前一定要验证能否编译。这个提醒放到今天依然适用。很多同学在本地用一个版本开发到笔试环境突然发现编译器版本老、库路径不对、显存被占用代码跑不起来。笔试环境通常不会给你太多排查时间平时练习就要养成“代码写出来顺手编译一遍”的习惯。5.2 常见运行错误与排查优先级上机笔试最容易翻车的不是思路而是一堆低级运行错误。我整理了一份速查表报错或现象常见原因排查优先级illegal memory access数组越界、空指针、索引算错1invalid argument核函数参数错误、grid或block配置非法1结果不正确或随机波动缺少同步、共享内存未初始化、浮点累加顺序变化2程序卡住不退出__syncthreads()使用不当导致死锁、无限循环2shared memory过大编译失败超过默认共享内存上限或未配置动态共享内存3性能比CPU还慢启动开销太大、访存模式太差、并行度太低3排查时不要慌先看索引和边界条件再看同步最后才怀疑硬件。我发现很多同学一旦报错就怀疑显卡出问题其实绝大多数是自己代码写错了。这里有个特别实用的技巧平时调试kernel时可以先跑一遍compute-sanitizer之类的工具能直接定位越界地址。笔试也许不允许用工具但平时训练一定要养成习惯。5.3 没GPU怎么练本地、WSL与云实例怎么选很多同学会问我笔记本没有独显怎么准备这种笔试这个问题其实好几年前就存在越早解决越好。如果你的机器有NVIDIA独立显卡第一选择是本地环境。Windows下可以装WSL在WSL里安装CUDA工具链注意检查驱动和容器权限。常见情况是WSL里装好驱动后运行程序时遇到无法初始化NVML的提示说GPU访问被操作系统拦截。这种问题大多不是代码问题而是WSL对GPU访问的权限或版本匹配问题要先更新Windows端驱动再确认WSL版本和CUDA支持。如果你的机器是纯核显或者根本没有NVIDIA GPU别在CPU模拟上死磕。CPU上虽然可以编译CUDA代码但没法跑也看不到真实性能和访存行为。最靠谱的做法是租一台带GPU的云主机或按小时收费的GPU实例把练习环境搭好。选镜像时注意驱动和CUDA Toolkit版本匹配装好后先跑一个vectorAdd验证环境再开始经典算子练习。顺便说一句训练GPU编程能力不能只看不写。至少要把归约、矩阵乘法、图像处理这类经典算子每个都手写一遍再用profiling工具看一遍报告对比优化前后差异。这种习惯远比考前刷题更能锻炼真正的优化直觉。6. 笔试之后的一些心里话6.1 一次笔试复盘带来的认知转变写到这里我想起那次笔试结束后我和同考场一个同学在走廊里聊了几句。他说题目不难但每一道都问得很细
![[光学原理与应用-521]：对光的错误理解与纠偏](http://pic.xiahunao.cn/yaotu/[光学原理与应用-521]：对光的错误理解与纠偏)
