1. 从一份笔试样题看英伟达到底想招什么样的人英伟达的暑期实习笔试在圈子里一直是个挺有意思的话题。它不像某些大厂那样上来就是四五道算法题往死里卷也不像纯八股文考试那样背一背就能过。我前后帮几个学弟学妹复盘过这套题也自己动手做过一遍最大的感受是这套卷子的设计逻辑非常“英伟达”——它不指望你什么都会但它非常在意你是否理解硬件和软件之间的那层关系。先把结论放在前面这份笔试样题的核心考察面集中在GPU 并行计算模型、深度学习基础算子、内存层次结构、以及 C/C 底层编程能力这四个方向上。它不会考你 PyTorch 的 API 怎么调但会考你__syncthreads()到底在同步什么它不会让你推导反向传播的完整公式但会问你一个卷积层在特定 stride 和 padding 下的输出尺寸是多少、显存占用大概多少。说白了它要的是能看懂 kernel、能算清楚资源开销、能定位性能瓶颈的人。这篇文章我会把这份样题拆开来讲不只是给答案更重要的是讲清楚每道题背后英伟达想验证的能力点是什么以及你在准备类似笔试时应该往哪个方向使劲。适合正在准备英伟达实习、GPU 相关岗位校招或者单纯想检验自己并行计算和深度学习底层功底的朋友。哪怕你暂时不投英伟达这套题的训练价值也足够高因为它考的东西在实际做 GPU 推理优化、算子开发的时候天天都会碰到。2. 笔试整体结构与考察意图拆解2.1 题型分布与时间压力分析从流传出来的样题结构看英伟达暑期实习笔试通常分为三个板块选择题约 15 到 20 道、填空题约 5 到 8 道、编程题2 到 3 道。总时长一般在 90 到 120 分钟之间。这个时间分配其实挺紧的尤其是编程题部分如果前面选择题卡太久后面基本写不完。选择题覆盖面很广从 CUDA 编程模型、GPU 架构基础、深度学习概念到 C 语法细节、操作系统内存管理甚至偶尔会冒出一两道概率统计题。填空题往往集中在具体数值计算上比如给你一个矩阵乘法的维度让你算 shared memory 需要多少字节或者给你一个卷积层的参数让你算 FLOPs。编程题则偏向CUDA kernel 实现和算法题的组合前者考察你对线程组织、内存访问模式的理解后者考察基础编码能力。我个人的判断是英伟达出这套题的底层逻辑是选择题筛掉基础不牢的填空题筛掉只会调包不会算账的编程题筛掉动手能力不行的。三个板块层层递进任何一块明显短板都会导致挂掉。2.2 为什么英伟达偏爱“计算密集型”题目你如果仔细看这些题会发现一个很明显的特征大量题目需要你动手算。不是那种“以下哪个选项正确”的概念题而是“给定条件 X请计算 Y”的数值题。这跟英伟达的业务性质高度相关。GPU 计算的核心就是资源管理——寄存器、shared memory、global memory、线程块、网格每一种资源都是有限的你写的 kernel 能不能跑起来、跑得快不快取决于你对这些资源的精确计算和控制。英伟达招实习生尤其是做 CUDA 开发、算子优化、深度学习框架底层的人第一要求就是对数字敏感对资源开销有直觉。举个例子一道典型的填空题可能是这样的一个线程块有 256 个线程每个线程需要 32 个寄存器shared memory 每个线程块分配 48KB问在寄存器文件大小为 64K 每 SM、shared memory 为 100KB 每 SM 的架构上每个 SM 最多能同时驻留多少个线程块这道题你需要同时考虑寄存器限制和 shared memory 限制取两者中更严格的那个。这种题没有捷径就是得把 GPU 架构的基本参数记牢然后一步步算。2.3 深度学习部分考什么、不考什么深度学习相关的题目在笔试中占比大概三成左右但英伟达的考法跟一般算法岗笔试完全不同。它不考你模型结构设计不考你调参经验也不考你最新论文的细节。它考的是深度学习算子的底层实现原理和计算量。比如卷积操作它可能会问你输入特征图 224×224×3卷积核 7×7×64stride 为 2padding 为 3输出特征图的尺寸是多少这个用公式一算就出来输出尺寸 floor((输入尺寸 2×padding - 卷积核尺寸) / stride) 1 floor((224 6 - 7) / 2) 1 floor(223/2) 1 111 1 112。答案是 112×112×64。再比如它可能会问你某个操作在 GPU 上的显存占用。一个 batch size 为 32 的 3×224×224 输入经过一个输出通道为 64 的 7×7 卷积层中间激活值需要多少显存你需要算输出特征图的大小32×64×112×112然后乘以每个元素的字节数通常是 4 字节 float32得出大约 32×64×112×112×4 ≈ 102MB。这种计算在模型部署和显存优化时是家常便饭。所以准备深度学习部分重点不是刷 LeetCode 那种题而是把 CNN 的基本组件卷积、池化、批归一化、激活函数的输入输出关系和计算量公式烂熟于心同时理解它们在 GPU 上执行时的主要开销来源。3. CUDA 编程模型核心考点与实操解析3.1 线程层次结构从 grid 到 thread 的映射关系CUDA 的线程组织是笔试必考的内容没有之一。你需要非常清楚这几个概念之间的关系grid 包含多个 blockblock 包含多个 thread。在 kernel 内部你可以通过blockIdx、threadIdx、blockDim、gridDim这几个内置变量来定位当前线程。一个典型的题目会这样出给你一个长度为 N 的一维数组你启动了一个 grid每个 block 有 256 个线程问需要多少个 block 才能覆盖整个数组答案是(N 255) / 256这是向上取整的标准写法。然后它可能会追问在 kernel 内部第 i 个线程处理的数组下标怎么算答案是int idx blockIdx.x * blockDim.x threadIdx.x;然后加一个边界检查if (idx N)。这些看起来简单但笔试中经常会在细节上设坑。比如把blockDim.x写成blockDim在 CUDA 里blockDim是一个 dim3 结构体必须加.x才能取到 x 维度的值或者忘记边界检查导致数组越界。我在帮人复盘的时候发现很多人概念都懂但一写代码就漏掉边界判断这个在笔试编程题里是直接扣分的。还有一个进阶考点是多维线程块。比如你处理一个二维矩阵可能会用dim3 block(16, 16)来组织线程这样每个 block 有 256 个线程threadIdx.x和threadIdx.y分别对应列和行。计算全局索引的时候就是int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x;。这个映射关系一定要练到条件反射的程度。3.2 内存层次与访问延迟为什么 shared memory 是性能关键GPU 的内存层次是另一个高频考点。你需要记住这个层级关系寄存器 → shared memory → L1/L2 cache → global memory访问延迟依次递增。寄存器最快基本没有延迟shared memory 次之大约几十个时钟周期global memory 最慢可能要几百个时钟周期。笔试中常见的考法是给你一段 kernel 代码让你判断它的性能瓶颈在哪里。比如一个矩阵乘法的朴素实现每个线程直接从 global memory 读取 A 和 B 的元素进行计算那么它的瓶颈就是 global memory 的访问带宽。优化方法就是使用 shared memory 做分块tiling把数据先加载到 shared memory 中然后在 shared memory 上做计算减少 global memory 的访问次数。这里有一个经典的计算题两个 N×N 的矩阵相乘朴素实现中每个线程需要读取 2N 个 global memory 元素总访问次数是 2N³。如果使用 T×T 的分块每个线程块需要加载 2T×T 个元素到 shared memory然后每个线程从 shared memory 读取 2T 个元素global memory 的访问次数降低到 2N³/T。这个 T 就是分块大小通常取 16 或 32。注意shared memory 的使用不是越多越好。每个 SM 的 shared memory 总量是有限的比如 100KB 左右具体取决于架构如果你每个 block 分配太多 shared memory会导致 SM 上能同时驻留的 block 数量减少反而降低 occupancy。这个权衡在笔试中经常考。3.3 同步与竞态__syncthreads() 的正确使用姿势__syncthreads()是 block 级别的同步屏障它的作用是让 block 内所有线程都执行到这个点之后再继续往下走。在 shared memory 分块矩阵乘法中你需要在加载完数据之后、开始计算之前调用一次__syncthreads()确保所有线程都完成了数据加载。计算完成后如果还要进行下一轮加载还需要再加一次同步。笔试中常见的坑是在分支语句中使用__syncthreads()。比如if (threadIdx.x 128) { __syncthreads(); }这种写法是未定义行为因为 block 内不同线程可能走不同的分支导致部分线程永远等不到同步点程序会挂死。正确的做法是把__syncthreads()放在所有线程都会执行到的位置。另一个考点是__syncthreads()和__syncwarp()的区别。前者同步整个 block后者只同步一个 warp32 个线程。在 Volta 架构之后warp 内的线程可以独立调度所以如果你只需要 warp 级别的同步用__syncwarp()更轻量。这个知识点在较新的笔试中出现的频率越来越高。3.4 一个完整的 CUDA 编程题实战演示假设笔试中出现了这样一道题实现一个 CUDA kernel计算两个向量的点积dot product。输入是两个长度为 N 的 float 数组 a 和 b输出是一个标量。要求使用 shared memory 做归约reduction。这道题的解题思路分三步第一步每个线程计算一对元素的乘积然后通过 shared memory 做 block 内的归约第二步每个 block 得到一个部分和写入 global memory第三步用一个单独的 kernel 或者 CPU 代码把所有的部分和加起来。核心代码大概长这样__global__ void dotProductKernel(float* a, float* b, float* result, int N) { __shared__ float cache[256]; int tid threadIdx.x; int idx blockIdx.x * blockDim.x threadIdx.x; float temp 0.0f; if (idx N) { temp a[idx] * b[idx]; } cache[tid] temp; __syncthreads(); // 归约 for (int stride blockDim.x / 2; stride 0; stride 1) { if (tid stride) { cache[tid] cache[tid stride]; } __syncthreads(); } if (tid 0) { result[blockIdx.x] cache[0]; } }这道题有几个容易出错的地方第一cache数组的大小必须和 block 的线程数一致这里假设是 256第二归约循环中每次迭代后都需要__syncthreads()否则会出现竞态第三边界检查if (idx N)不能漏否则数组越界。实操心得在笔试中写 CUDA 代码如果时间紧张先把框架搭出来确保线程索引计算、边界检查、同步这三个关键点不出错再去优化细节。我见过太多人因为漏了一个__syncthreads()导致整个 kernel 逻辑错误得不偿失。4. 深度学习算子与 GPU 计算量分析4.1 卷积操作的输出尺寸与计算量公式卷积是深度学习中最核心也最耗算力的操作英伟达笔试中几乎每年都会考。你需要记住两个公式输出尺寸公式和FLOPs 计算公式。输出尺寸公式前面已经提过Output floor((Input 2×Padding - Kernel) / Stride) 1。这个公式对宽和高分别适用。如果是池化操作公式类似只是没有 padding 的常见情况。FLOPs 的计算稍微复杂一点。一个卷积层的浮点运算次数大约是2 × Output_H × Output_W × Output_C × Kernel_H × Kernel_W × Input_C。这里的 2 是因为每次乘加算两次运算。举个例子输入 224×224×3卷积核 3×3输出通道 64stride 1padding 1那么输出尺寸是 224×224×64FLOPs 2 × 224 × 224 × 64 × 3 × 3 × 3 ≈ 2 × 224 × 224 × 64 × 27 ≈ 173 million FLOPs。这个计算在模型分析中非常常用。笔试中可能会给你一个完整的网络结构让你算总 FLOPs 或者总参数量。这时候你需要逐层计算然后累加。参数量相对简单卷积层的参数量就是Kernel_H × Kernel_W × Input_C × Output_C Output_C加上偏置。全连接层的参数量是Input_Dim × Output_Dim Output_Dim。4.2 矩阵乘法在 GPU 上的分块优化原理矩阵乘法是深度学习的另一个核心操作也是 GPU 性能优化的经典案例。笔试中可能会让你分析不同实现方式的性能差异或者让你计算分块大小对 occupancy 的影响。朴素矩阵乘法的问题是 global memory 访问太频繁。对于 C A × B每个输出元素 C[i][j] 需要读取 A 的第 i 行和 B 的第 j 列总共 2N 次 global memory 访问。N×N 个输出元素就是 2N³ 次访问。而 GPU 的 global memory 带宽是有限的这就成了瓶颈。分块优化的思路是把 A 和 B 分成 T×T 的小块每个线程块负责计算 C 的一个 T×T 子块。线程块先把 A 和 B 对应的两个 T×T 块加载到 shared memory然后每个线程从 shared memory 读取数据计算。这样 global memory 的访问次数降低到 2N³/Tshared memory 的访问虽然多了但 shared memory 的带宽远高于 global memory所以整体性能大幅提升。T 的选择需要权衡。T 太小global memory 访问减少不明显T 太大shared memory 不够用occupancy 下降。通常 T 取 16 或 32 是比较好的平衡点。在笔试中如果让你选 T你可以从 shared memory 容量反推每个线程块需要 2×T×T×4 字节的 shared memoryfloat32如果 SM 的 shared memory 是 100KB那么 T 最大可以取到 sqrt(100×1024 / 8) ≈ 113但实际还要考虑寄存器和其他开销所以 32 是比较稳妥的选择。4.3 批归一化与激活函数的 GPU 实现要点批归一化Batch Normalization和激活函数ReLU、Sigmoid 等虽然计算量不大但在笔试中经常作为小题出现。你需要知道它们在 GPU 上的实现要点。批归一化的前向过程是对每个通道计算 batch 内的均值和方差然后做归一化最后乘以缩放因子 gamma 并加上偏移 beta。在 GPU 上实现时计算均值和方差需要做归约操作这跟前面点积的归约类似。笔试可能会问你如果一个 batch 有 32 个样本每个样本有 64 个通道每个通道的特征图是 112×112那么计算均值时需要归约多少个元素答案是 32×112×112 401408 个元素。这个归约操作通常用两个 kernel 完成第一个 kernel 计算每个通道的部分和第二个 kernel 汇总。ReLU 就简单多了就是max(0, x)逐元素操作GPU 上直接一个线程处理一个元素就行。但笔试可能会考你 ReLU 的导数x 0 时为 1x ≤ 0 时为 0。这个在反向传播中会用到。常见问题很多人搞不清楚批归一化在训练和推理时的区别。训练时用当前 batch 的均值和方差推理时用训练阶段累积的移动平均均值和方差。这个知识点在笔试选择题中出现过多次。4.4 显存占用估算一个完整的计算示例显存估算是我认为最实用的考点之一因为在实际工作中你经常需要判断一个模型能不能在给定显卡上跑起来。笔试中通常会给你一个网络结构和 batch size让你估算训练时的显存占用。显存占用主要包括三部分模型参数、梯度、优化器状态、中间激活值。模型参数和梯度的大小就是参数量乘以 4 字节float32。优化器状态如果是 Adam每个参数需要额外存储一阶矩和二阶矩所以是参数量的 2 倍再乘以 4 字节。中间激活值取决于网络结构和 batch size需要逐层计算。举个例子一个简单的两层全连接网络输入 784 维隐藏层 256 维输出 10 维batch size 为 64。参数量 784×256 256 256×10 10 ≈ 200K 2.5K ≈ 203K。参数显存 203K × 4 ≈ 812KB。梯度同样约 812KB。Adam 优化器状态约 1.6MB。中间激活值第一层输出 64×256×4 ≈ 64KB第二层输出 64×10×4 ≈ 2.5KB。总计大约 3.3MB。这个规模很小任何显卡都能跑。但如果换成 ResNet-50参数量约 25M参数显存约 100MB梯度 100MBAdam 状态 200MB中间激活值在 batch size 为 32 时可能达到几个 GB。这时候就需要考虑用混合精度训练、梯度累积、或者模型并行来降低显存占用。5. 编程题解题策略与代码实现细节5.1 算法题部分难度定位与时间分配英伟达笔试的算法题难度大概在 LeetCode Medium 水平偶尔会有一道 Hard。题目类型偏向数组操作、字符串处理、动态规划、以及一些跟 GPU 计算相关的模拟题。比如有一道题是模拟 GPU 的线程调度给你一组任务和线程块大小让你计算需要多少个时钟周期完成所有任务。这种题本质上是一个模拟题需要你仔细读题把逻辑理清楚再写代码。时间分配上我建议选择题控制在 30 分钟以内填空题 20 分钟剩下的时间全部留给编程题。如果编程题有两道每道至少留 25 分钟。如果遇到卡住的题先跳过把能拿的分拿到手。英伟达的笔试通常是按通过测试用例的数量给分所以即使不能完全通过写出一个能过部分用例的暴力解法也比空着强。5.2 CUDA 编程题的常见陷阱与调试技巧CUDA 编程题是英伟达笔试的特色也是最容易拉开差距的地方。除了前面提到的线程索引计算、边界检查、同步问题之外还有几个常见的陷阱。第一个是共享内存的 bank conflict。shared memory 被分成 32 个 bank如果同一个 warp 内的多个线程访问同一个 bank 的不同地址就会发生 bank conflict导致访问串行化。比如cache[tid]这种访问模式如果 tid 连续那么每个线程访问不同的 bank没有冲突。但如果 stride 是 32 的倍数就会全部落到同一个 bank 上。笔试中可能会给你一段代码让你判断有没有 bank conflict或者让你改写代码消除冲突。第二个是warp divergence。如果一个 warp 内的线程走了不同的分支GPU 会串行执行所有分支导致性能下降。比如if (threadIdx.x % 2 0) { ... } else { ... }同一个 warp 内一半线程走 if一半走 else就会产生 divergence。优化方法是尽量让同一个 warp 内的线程走相同的分支或者用一些技巧把分支消除掉。第三个是寄存器溢出。如果 kernel 中使用了太多局部变量编译器可能会把一些变量放到 local memory实际上是 global memory导致性能急剧下降。笔试中可能会给你一个 kernel让你分析它的寄存器使用情况或者让你优化代码减少寄存器压力。实操心得在笔试中写 CUDA 代码如果拿不准某个优化是否有效优先保证正确性。先把朴素版本写对再考虑优化。我见过有人为了追求性能写了一堆复杂的优化结果连正确性都没保证得不偿失。5.3 代码风格与注释阅卷人想看到什么虽然是机考但代码风格和注释仍然很重要尤其是编程题部分。英伟达的阅卷系统可能会人工复核部分代码清晰的代码结构和适当的注释能帮你加分。具体来说函数命名要有意义变量名不要用 a、b、c 这种除非是数学公式中的标准符号。关键步骤要加注释比如// 计算全局线程索引、// 边界检查、// 同步确保数据加载完成。如果用了比较复杂的优化技巧比如 shared memory 分块最好在注释里简单说明一下思路。另外代码的鲁棒性也很重要。比如输入数组可能为空或者 N 可能为 0这些边界情况要考虑进去。虽然笔试的测试用例可能不会覆盖这些情况但写了总比没写好。6. 常见问题与排查技巧实录6.1 笔试中容易犯的低级错误速查表错误类型具体表现后果预防方法线程索引计算错误忘记加 blockIdx.x * blockDim.x部分数据未处理或重复处理写完索引公式后代入几个值验证边界检查遗漏直接访问 a[idx] 而不判断 idx N数组越界程序崩溃所有 global memory 访问前都加检查同步缺失shared memory 写入后未 __syncthreads()数据竞态结果不确定每次 shared memory 写后读前加同步共享内存大小超限声明了过大的 shared 数组kernel 启动失败提前计算 shared memory 需求寄存器溢出局部变量过多性能急剧下降减少局部变量复用寄存器浮点数精度问题累加顺序导致结果偏差测试用例不通过使用 Kahan 求和或调整归约顺序这张表里的每一条我都在实际笔试或帮人复盘时见过。尤其是边界检查和同步缺失几乎是最高频的错误。建议在写完代码后花两分钟逐行检查这几个点。6.2 时间不够用时的取舍策略笔试时间紧张是常态关键是要有取舍策略。我的建议是先扫一遍所有题目把最有把握的题目标记出来优先做这些。选择题中概念题和简单计算题先做复杂的计算题如果 30 秒内没思路就跳过。填空题中如果公式记不清先跳过后面有时间再回来推。编程题中先写暴力解法保证能过部分用例再考虑优化。还有一个技巧是利用选择题的选项反推答案。有些计算题你不需要完整算出结果只需要估算一个范围然后看哪个选项落在范围内。比如算 FLOPs你只需要知道数量级是百万还是十亿就能排除大部分错误选项。6.3 考前一周的冲刺准备清单如果你还有一周就要考了我建议按这个清单来准备第一天到第二天把 CUDA 编程模型的核心概念过一遍重点看线程层次、内存层次、同步机制。手写几个经典的 kernel向量加法、矩阵乘法、归约。第三天到第四天复习深度学习算子的计算公式包括卷积输出尺寸、FLOPs、参数量、显存占用。找几个经典网络LeNet、AlexNet、ResNet练手算。第五天刷一遍 C 的基础知识重点是内存管理、指针、引用、虚函数、模板。英伟达的编程题用 C 写语法不熟会很吃亏。第六天做一套完整的模拟题严格计时模拟真实考试环境。第七天复盘错题把容易忘的公式和概念整理成一页纸考前快速过一遍。注意不要花太多时间在偏题怪题上。英伟达的笔试整体风格是稳扎稳打基础题占大多数。把基础打牢比押题有用得多。6.4 面试环节可能追问的延伸问题笔试通过之后面试环节可能会针对你的笔试答案追问。比如你在编程题中用了 shared memory 分块面试官可能会问你为什么选择这个分块大小有没有考虑过 bank conflict如果矩阵维度不是分块大小的整数倍怎么办这些问题需要你对自己的代码有深入的理解不能只是“背”了一个模板。还有一个常见的追问是如果让你优化这个 kernel 的性能你会从哪些方面入手这时候你可以从减少 global memory 访问、提高 occupancy、消除 bank conflict、减少 warp divergence这几个角度来回答。每个角度都能展开讲很多细节面试官主要看你的思路是否清晰、是否有实际优化经验。7. 从笔试到实战这些考点在实际工作中的映射7.1 算子开发岗位的日常与笔试的关联如果你将来真的进了英伟达做算子开发或者 GPU 优化你会发现笔试考的东西几乎每天都在用。比如你写一个卷积算子第一步就是算输出尺寸和显存占用然后设计线程块和网格接着考虑用不用 shared memory、怎么分块、怎么归约。这些跟笔试编程题的流程一模一样。区别在于实际工作中你还需要考虑更多因素不同 GPU 架构的差异比如 Ampere 和 Hopper 的 shared memory 大小不同、混合精度FP16、BF16、TF32、Tensor Core 的使用、以及跟上层框架的对接。但底层的那套思维模式——算清楚资源开销、设计高效的线程组织、避免常见的性能陷阱——是笔试和实战共通的核心能力。7.2 深度学习框架底层开发的技能树如果你对深度学习框架底层开发感兴趣笔试中的深度学习算子考点就是你的入门基础。你需要进一步学习的是自动微分的基本原理、计算图的构建和优化、算子融合、内存池管理、以及多流并行。这些内容在笔试中不会直接考但它们是笔试考点的自然延伸。举个例子笔试考了卷积的输出尺寸计算实际工作中你可能需要实现一个卷积算子的反向传播这就需要你推导卷积的转置transposed convolution的输入输出关系。笔试考了 shared memory 分块实际工作中你可能需要实现一个高效的矩阵乘法库这就需要你深入理解不同分块策略对 cache 命中率的影响。7.3 如何把笔试准备转化为长期竞争力最后说点实在的。准备英伟达笔试的过程本质上是在补计算机体系结构和并行计算的基础课。这些东西不是考完就忘的它们会一直影响你写代码的方式。即使你最后没去英伟达去了其他做 AI 芯片、自动驾驶、高性能计算的公司这些知识同样适用。我的建议是不要把笔试准备当成一个短期的应试任务而是把它当成一次系统学习的机会。把 CUDA 编程模型、GPU 架构、深度学习算子这三块内容真正搞懂你收获的不仅是一份实习 offer更是一套可以长期使用的技术栈。我在实际工作中遇到性能问题的时候脑子里第一时间浮现的往往就是当年准备笔试时反复练习的那些计算和优化思路。
