2018年秋天我在考场上拿到百度这套“AI异构计算工程师”笔试题第一批的时候第一反应是这公司是真的打算把做AI的人按硬件工程师的标准来筛。当时这个岗位在市面上非常少见大部分公司招算法工程师笔试考的是机器学习理论和手推公式最多再来一两道LeetCode。可百度的题目完全不是这个路数。整套卷子围绕一条主线展开假如让你把一个深度学习模型跑得飞快你到底懂不懂底层发生了什么。如果你只会写Python、调框架那这份卷子大概率会让你怀疑人生。但如果你本身有CUDA编程经验或者认真读过GPU架构相关的材料很多题其实是送分题。这篇文章我按“考后复盘”的思路把当年印象最深的几类题目、背后考的知识点、以及我当时踩过的坑完整拆一遍。内容不保证和原卷逐字一致但考点分布和解题思路是可以直接参考的对准备AI基础设施、高性能计算、异构计算相关岗位的同学尤其有用。1. 考后复盘这份卷子到底在筛选什么样的人1.1 从岗位JD反推考点分布百度在2018年校招里单独立一个“AI异构计算工程师”的岗位本身就是很明确的信号深度学习已经不只是算法的竞争而是工程效率的竞争。模型得跑在GPU上推理要压低延迟训练要提升吞吐算力集群要降低成本。这些全落到“异构计算”这四个字上。所以笔试题的考点分布其实非常合理我当时整理了一下大概是这么几块考点方向常见出题形式对应能力异构计算基本概念简答/选择是否理解CPUGPUFPGA等协同工作的本质GPU硬件架构选择/填空是否知道SM、Warp、线程层级这些基础概念CUDA编程模型程序分析/改错是否真正写过并调过CUDA代码访存优化计算题/设计题是否理解内存合并、共享内存、Bank冲突AI算子优化推导题是否接触过卷积的im2col、GEMM等实现并行算法手写代码是否具备并行思维能写正确的并发程序这个结构放到今天依然成立。别指望只背几个面试八股就能过它考的是“你有没有亲手做过”的细节。1.2 整体难度节奏与时间压力整套卷子的时间我记得是两小时左右题量不算特别大但计算推导题很占时间。最难受的地方在于前面的简答题你觉得自己写得挺好结果到了后面手写CUDA内核的时候光是索引对不对、同步加没加、边界怎么处理就得反复推敲。我个人做下来的感受是前40分钟做概念题中间60分钟死磕CUDA和算子优化最后20分钟写并行归约时间刚刚好。如果你在概念题上犹豫太久后面代码题基本没救。从难度曲线上看这张卷子不是那种“从易到难”的友好型试卷而是“先易后难最后突然开放”的结构。最后一道题往往没有标准答案考的是你分析问题的思路。这种题千万别空着哪怕只写一个优化方向也能拿到部分分数。2. 硬件基础题把CPU、GPU、存储器层次讲清楚2.1 异构计算概念题怎么答才满分这套卷子开篇第一题基本是概念题大致问法是“什么是异构计算为什么深度学习需要异构计算”这种题看着简单但想拿满分其实有讲究。我当时是这么回答的异构计算指在一个系统里同时使用多种不同类型的处理器来协同完成计算任务。CPU擅长复杂逻辑控制和少量数据的快速处理GPU擅长大规模数据并行计算FPGA擅长低延迟、可定制的流水线计算ASIC则在特定场景下做到极致的能效比。异构计算的关键不是堆硬件而是把任务按照特性拆分到最合适的计算单元上。深度学习需要异构计算的原因可以从训练和推理两个阶段来看。训练阶段核心操作是大量稠密矩阵乘法这正是GPU的强项而数据预处理、分布式通信调度、模型参数的更新逻辑依然由CPU负责。推理阶段更复杂高吞吐场景可以用GPU做批量推理低延迟场景可能需要FPGA定制流水线端侧场景只能依赖NPU或DSP。所以整个AI系统天然就是一个异构系统。如果你只是写“异构计算就是CPU加GPU”那只能拿一半分。加分点在于“为什么适用”和“任务如何切分”。我当时还顺手补了一句异构计算的本质是“用合适的计算单元做合适的事”这句话放在任何场景都能概括核心思想。2.2 线程层级与索引计算细节第二类必考是GPU线程模型。题目会让你简述CUDA的Grid、Block、Thread三级结构并说明它们和硬件SM、Warp的映射关系。标准答法是一个Kernel启动后对应一个GridGrid由多个Block组成Block由多个Thread组成Block会被调度到SM上执行一个SM可以同时驻留多个BlockGPU实际执行指令的最小单位是Warp通常包含32个线程Warp内的32个线程以SIMT单指令多线程方式执行即同一时刻执行同一条指令但各自处理不同的数据这里有一个非常容易错的细节线程的调度顺序。Block内的线程不是一次性全部执行而是按Warp为单位切分。如果一个Block有128个线程它会被分成4个Warp。Warp之间在GPU上轮流执行遇到访存等长延迟操作时会切换其他Warp来隐藏延迟。这个机制叫“线程级并行”配合“延迟隐藏”是GPU和CPU在设计哲学上最大的区别。我当时做这题的时候还补了一张线程索引的推算思路一个Block内第tid个线程的全局ID如果是一维结构就是blockIdx.x * blockDim.x threadIdx.x如果是二维Block则是blockIdx.x * blockDim.x threadIdx.x这是最低频的一层上面还有blockIdx.y * gridDim.x blockIdx.x这样逐层叠加的过程。虽然2018年的卷子没有直接要求写二维索引但清楚这个逻辑对后边的代码题是隐形加分。注意写CUDA代码时索引错误是最隐蔽的bug。程序不会报错只是结果不对而且通常在数据量大到一定程度才暴露。笔试现场切忌只写主体逻辑不验证索引边界建议每一步都推演一两个小数据量样例。3. CUDA编程题矩阵乘法与共享内存优化3.1 第一版朴素写法得基础分矩阵乘法是这套卷子的重头戏。题目一般给一个矩阵乘法的CUDA实现框架要求你补全内核或写出一个可运行版本。2018年的题我记得大致是给定两个1024x1024的float矩阵A和B写一个CUDA Kernel计算C A * B。最朴素的内核是这样的__global__ void matmul_naive(float* A, float* B, float* C, int N) { int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; if (row N col N) { float sum 0.0f; for (int k 0; k N; k) { sum A[row * N k] * B[k * N col]; } C[row * N col] sum; } }这个写法在功能上完全正确理论上也能跑但性能非常差。原因在于内层循环每次都直接访问全局内存而A和B的访存模式完全没有利用局部性。每次循环都要从显存读两个float一个矩阵乘法做下来全局内存访问次数是O(N^3)而理想情况应该是O(N^2)。为什么性能差核心在于GPU的运算单元比访存带宽“快得多”。2018年主流的P100/V100FP32算力都在10 TFLOPS以上但HBM显存带宽只有几百GB/s。这种算力和带宽的差距意味着如果每个计算都需要从全局内存取数计算单元大部分时间都在等数据GPU根本跑不满。要解决这个问题只有一条路利用共享内存和寄存器做数据复用把“重复从全局内存读数据”变成“从片上缓存高速读取”。3.2 分块优化和共享内存tile的写法分块矩阵乘法的思路是把大矩阵切成多个小Tile每个Block负责计算输出矩阵中的一个Tile。计算时先把A和B对应位置的Tile从全局内存拷贝到共享内存然后从共享内存中反复读取数据参与计算。#define TILE_SIZE 32 __global__ void matmul_tiled(float* A, float* B, float* C, int N) { __shared__ float As[TILE_SIZE][TILE_SIZE]; __shared__ float Bs[TILE_SIZE][TILE_SIZE]; int tx threadIdx.x; int ty threadIdx.y; int row blockIdx.y * TILE_SIZE ty; int col blockIdx.x * TILE_SIZE tx; float sum 0.0f; for (int t 0; t N / TILE_SIZE; t) { As[ty][tx] A[row * N t * TILE_SIZE tx]; Bs[ty][tx] B[(t * TILE_SIZE ty) * N col]; __syncthreads(); for (int k 0; k TILE_SIZE; k) { sum As[ty][k] * Bs[k][tx]; } __syncthreads(); } C[row * N col] sum; }这个版本比朴素版快了一个数量级。关键优化点有三个数据复用A矩阵的每个元素会被B矩阵同一Tile内的所有线程复用通过共享内存将全局内存访问次数从O(N^3)降为O(N^2 / TILE_SIZE)内存合并在拷贝As和Bs时相邻线程访问相邻地址满足全局内存合并访问的条件带宽利用率高同步保证__syncthreads()确保共享内存写入完成后再读取防止数据竞争笔试判分的时候面试官主要看你的内核里有没有用__shared__以及有没有正确处理同步。如果只写朴素版基本就是及格分写出分块版直接进入加分区间。3.3 Bank冲突真题计算与手算过程共享内存虽然快但也不是没有限制。它被划分成32个Bank每个Bank的宽度是4字节。当同一个Warp内的线程同时访问同一个Bank的不同地址时硬件会把这次访问串行化导致性能下降这就是Bank冲突。当年卷子里的经典计算题我记忆很深题干大概是这样共享内存中有float数组假设每个float占4字节某Warp内线程threadIdx.x访问第 threadIdx.x * 8 个float元素请计算这些地址分布到哪些Bank判断是否存在Bank冲突冲突几路。手算过程如下Bank编号的计算方式是地址按4字节为单位除以32取余数。第i个线程访问的是数组下标 threadIdx.x * 8 的元素换算成字节地址是 threadIdx.x * 8 * 4 threadIdx.x * 32。再除以4字节得到Bank单元编号threadIdx.x * 8。对32取余线程00 % 32 0线程18 % 32 8线程216 % 32 16线程324 % 32 24线程432 % 32 0线程540 % 32 8…规律逐渐清晰每4个线程映射到同一个Bank32个线程总共只用到8个Bank每个Bank被4个线程击中。这就是典型的4路Bank冲突实际访存耗时是理想情况的4倍。解决Bank冲突的常见手段有两种。一是给数组加Padding例如在第二个维度末位多申请一个float空间让原本对齐到同一Bank的地址错开二是在计算索引时乘以一个和Bank数互质的系数让地址重新散列。笔试时只要写出“加Padding”和“修改索引映射”任意一种再配合上面的手算过程基本就能拿满分。4. AI算子题卷积展开与访存优化4.1 im2col推导与空间代价百度这套卷子既然定位是AI异构计算卷积相关题目是绕不开的。2018年考的卷积题现在看其实是后来所有推理引擎的基础操作。题目问的是简述im2colimage to column的原理并分析这种方法的空间复杂度。im2col是把卷积运算转换成矩阵乘法的高效方法。假设输入是N个通道数为C、高H、宽W的特征图卷积核大小是K输出特征图的高宽设为H_out和W_out。im2col的过程是取出输入上每一个卷积窗口覆盖的区域将其展平成向量拼成一个二维矩阵。原始输入经过im2col后得到一个形状为 (C * K * K) × (N * H_out * W_out) 的大矩阵。这时候卷积核本身也可以展开成 K * K * C 行 × M列M为卷积核数量两者做矩阵乘法结果就是卷积输出。因为GPU上的GEMM通用矩阵乘法已经被优化到了接近硬件峰值所以im2col GEMM的流程在很长一段时间内是CNN推理的主流实现。空间代价是这道题的灵魂。展开矩阵的规模大约是原始输入数据的 K * K 倍。以3x3卷积为例im2col之后的数据量约是原始数据的9倍。对于大特征图这种空间开销是灾难级的。所以后续才有内存优化的变种比如在GEMM执行过程中动态做im2col不显式生成完整矩阵或者改用隐式im2col直接用地址计算代替物理拷贝。理解这个演进过程比单纯背im2col的定义重要得多。4.2 1x1卷积为什么在GPU上如此高效还有一道考1x1卷积的题当年回答的时候我只写了“相当于逐点线性变换”后来复盘发现其实没答到点上。1x1卷积在GPU上效率高的原因是它在做跨通道的信息融合而不改变空间尺寸。把特征图视为 N * H * W 个像素点每个像素点有C个通道值1x1卷积等价于对每个像素点做一次 C→M 的矩阵乘。在NCHW布局下这本身就是一次标准的GEMM而且K维度恰好是输入通道数C属于中小尺寸矩阵乘法GPU的Tensor Core对这种形状有专门优化。更重要的是1x1卷积天然适配内存合并访问。在NCHW布局下同一像素点的C个通道值在内存中不连续需要做转置或使用特殊的访存模式但如果把数据布局改成NHWC同一像素点的通道值是连续的读取时可以一次性载入寄存器再参与矩阵运算。这也是为什么很多推理引擎在GPU上会把数据转换成NHWC核心原因之一就是让1x1卷积和后续的GEMM算子访存更友好。实操心得笔试遇到这种“为什么在GPU上高效”的开放题千万别只答表面原因。尽量往数据布局、访存连续性、硬件指令集这三个方向靠分自然就高了。我当时只答了“矩阵乘”后来反思应该把NHWC布局的访存优势也写进去那才是GPU性能优化的真正命门。5. 并行算法题手写归约求和5.1 线程分歧是隐藏丢分点并行归约Reduction是异构计算笔试里出现概率极高的代码题它考察的是并发编程最核心的思维如何把一个整体任务拆成众多小任务然后高效地合并结果。题目一般是给定一个包含1024个浮点数的数组写出一个CUDA Kernel计算所有元素之和。这题看似简单但陷阱极多。最容易出问题的是线程分歧。如果你这样写__global__ void reduce_bad(float* input, float* output) { __shared__ float sdata[1024]; int tid threadIdx.x; int i blockIdx.x * blockDim.x threadIdx.x; sdata[tid] input[i]; __syncthreads(); for (int stride 512; stride 0; stride 1) { if (tid stride) { sdata[tid] sdata[tid stride]; } __syncthreads(); } if (tid 0) output[blockIdx.x] sdata[0]; }这个写法对吗功能上对。但性能上有很大的问题随着stride减小每轮参与计算的线程越来越少但Warp依然要按照32个线程一组的粒度执行。当stride小到一定程度时大量Warp里的线程只有部分是活跃的其他线程空转造成严重的线程分歧和资源浪费。更好的做法是从一开始就让线程处理多个元素先用循环进行一个block内的部分归约把数据量降到blockDim.x个再进入树状归约阶段。另一个优化点是让stride从blockDim.x/2开始而不是从总长度的一半开始这样能保证每一轮不浪费线程。还有一点是避免分支在关键循环里让所有线程执行相同的指令只是在数据选择上不同可以通过地址计算而不是if判断来处理边界条件减少控制流的开销。5.2 一个更稳的归约写法笔试现场推荐下面这个版本。它结合了“每线程多元素”和“交错归约”两种策略__global__ void reduce_good(float* input, float* output, int n) { __shared__ float sdata[256]; int tid threadIdx.x; int i blockIdx.x * blockDim.x threadIdx.x; float sum 0.0f; // 每个线程处理多个元素减少block数量提高数据复用 int stride gridDim.x * blockDim.x; for (int j i; j n; j stride) { sum input[j]; } sdata[tid] sum; __syncthreads(); // 树状归约 for (int s blockDim.x / 2; s 0; s 1) { if (tid s) { sdata[tid] sdata[tid s]; } __syncthreads(); } if (tid 0) output[blockIdx.x] sdata[0]; }核心改进有两个。第一每个线程先循环累加多个元素避免启动过多Block导致调度开销第二树状归约只对共享内存中的数据操作速度极快。这版代码还有一个隐藏知识点blockDim.x取256而不是取512或1024。原因是共享内存越大一个SM上能同时驻留的Block越少占用率会下降。256线程的block通常能让SM驻留足够多的block来隐藏访存延迟是实践中很均衡的配置。6. 现场调试与高频翻车点6.1 同步漏写为什么会导致随机错误笔试现场还有一类题目是程序改错给一段CUDA代码让你指出错误。当年最经典的陷阱就是漏写__syncthreads()。在矩阵乘法分块版本里如果没有在共享内存写完后加同步就会出现一个经典问题Block内不同的Warp执行进度不一致。Warp 0可能已经把数据写到了共享内存正打算读而Warp 1可能还没开始写。如果Warp 0抢跑去读读到的就是未初始化或上一轮残留的脏数据。更麻烦的是这种错误是随机且不可复现的。同一份代码有的机器上跑一万次没问题换个GPU型号或改用更大规模数据结果就错了。笔试现场没有GPU可跑所以只能靠“代码审阅”发现这类问题。我的经验是看到共享内存写入后一律条件反射地确认后面有没有同步这是保命习惯。另一个高频错误是核函数启动参数里没有检查边界。比如矩阵大小不是blockDim的整数倍时如果每个线程只计算一个元素那么grid外侧的线程就会越界访问。正确做法是在内核开头加上if (row N col N) { // 计算 }这个边界检查虽然会让代码多几行但它是鲁棒性的保障笔试中写出边界检查会明显加分。6.2 分析工具与查错顺序的建议如果要我推荐一套查错路径顺序是这样的先看Kernel有没有越界或同步缺失再看数据布局是否满足合并访问接着看Bank冲突和占用率最后才看指令级优化。网上很多人喜欢一上来就谈NVVP/Nsight的指标但我个人的体会是如果没有确认代码逻辑正确性能分析毫无意义。当年笔试最后有一道开放式题目问“如何定位CUDA程序性能瓶颈”如果你只知道答“用profiler看热点”得分有限。更系统的答法是先通过单元测试验证功能正确性再用有限的实验数据比如矩阵规模从小增大观察耗时增长曲线判断算法复杂度是否符合预期然后利用profiler对比实际访存带宽和理论峰值带宽的差距如果两者接近问题在计算如果差距大问题在访存最后针对访存问题检查缓存命中率、合并访问率和Bank冲突情况。这个排查思路放在今天依然通用而且适用于任何性能调优场景不局限于CUDA。7. 给后来人的复习建议与延伸方向这套笔试题虽然出自2018年但它的考点框架在AI基础设施领域一直有效。后来无论是我去面试其他大厂的AI加速岗还是在开源社区看一些高性能推理引擎的设计方案都能在当年这套卷子里找到原点。如果现在有同学准备类似岗位我给三点建议第一不要只刷题一定要动手写CUDA。哪怕手头没有GPU也可以用在线平台跑小样例。只有亲手写过一遍你拿到题目的时候才知道哪些地方容易出边界错误哪些优化是纸上谈兵。第二重视内存布局和访存模式这比背诵各种并行模式重要得多。绝大多数性能瓶颈不是算力不够而是数据喂不饱计算单元。第三对卷积、矩阵乘、归约、转置这些经典算子至少要能手写一套可运行版本同时能解释清楚优化前后访存次数的变化量级。另外异构计算这些年也在快速演进。从早期单纯CPUGPU到引入Tensor Core、NPU、可重构计算再到如今围绕大模型训推的算力集群设计本质都是在“任务切分”和“数据搬运”之间做权衡。理解了2018年这套笔试题想考的东西后续学任何新硬件、新框架都会觉得有迹可循。最后一个个人经验笔试答开放题时哪怕不会完整推导也要把“我理解的瓶颈是什么”“我打算从哪个方向优化”写清楚。这类题看的不是你记住了多少现成答案而是面对一个陌生性能问题时你有没有一套自己的分析路径。这种能力不是短期突击能练出来的需要平时多读代码、多看硬件白皮书并且在踩坑中不断复盘。
百度AI异构计算工程师笔试题复盘：CUDA与GPU性能优化核心考点
2018年秋天我在考场上拿到百度这套“AI异构计算工程师”笔试题第一批的时候第一反应是这公司是真的打算把做AI的人按硬件工程师的标准来筛。当时这个岗位在市面上非常少见大部分公司招算法工程师笔试考的是机器学习理论和手推公式最多再来一两道LeetCode。可百度的题目完全不是这个路数。整套卷子围绕一条主线展开假如让你把一个深度学习模型跑得飞快你到底懂不懂底层发生了什么。如果你只会写Python、调框架那这份卷子大概率会让你怀疑人生。但如果你本身有CUDA编程经验或者认真读过GPU架构相关的材料很多题其实是送分题。这篇文章我按“考后复盘”的思路把当年印象最深的几类题目、背后考的知识点、以及我当时踩过的坑完整拆一遍。内容不保证和原卷逐字一致但考点分布和解题思路是可以直接参考的对准备AI基础设施、高性能计算、异构计算相关岗位的同学尤其有用。1. 考后复盘这份卷子到底在筛选什么样的人1.1 从岗位JD反推考点分布百度在2018年校招里单独立一个“AI异构计算工程师”的岗位本身就是很明确的信号深度学习已经不只是算法的竞争而是工程效率的竞争。模型得跑在GPU上推理要压低延迟训练要提升吞吐算力集群要降低成本。这些全落到“异构计算”这四个字上。所以笔试题的考点分布其实非常合理我当时整理了一下大概是这么几块考点方向常见出题形式对应能力异构计算基本概念简答/选择是否理解CPUGPUFPGA等协同工作的本质GPU硬件架构选择/填空是否知道SM、Warp、线程层级这些基础概念CUDA编程模型程序分析/改错是否真正写过并调过CUDA代码访存优化计算题/设计题是否理解内存合并、共享内存、Bank冲突AI算子优化推导题是否接触过卷积的im2col、GEMM等实现并行算法手写代码是否具备并行思维能写正确的并发程序这个结构放到今天依然成立。别指望只背几个面试八股就能过它考的是“你有没有亲手做过”的细节。1.2 整体难度节奏与时间压力整套卷子的时间我记得是两小时左右题量不算特别大但计算推导题很占时间。最难受的地方在于前面的简答题你觉得自己写得挺好结果到了后面手写CUDA内核的时候光是索引对不对、同步加没加、边界怎么处理就得反复推敲。我个人做下来的感受是前40分钟做概念题中间60分钟死磕CUDA和算子优化最后20分钟写并行归约时间刚刚好。如果你在概念题上犹豫太久后面代码题基本没救。从难度曲线上看这张卷子不是那种“从易到难”的友好型试卷而是“先易后难最后突然开放”的结构。最后一道题往往没有标准答案考的是你分析问题的思路。这种题千万别空着哪怕只写一个优化方向也能拿到部分分数。2. 硬件基础题把CPU、GPU、存储器层次讲清楚2.1 异构计算概念题怎么答才满分这套卷子开篇第一题基本是概念题大致问法是“什么是异构计算为什么深度学习需要异构计算”这种题看着简单但想拿满分其实有讲究。我当时是这么回答的异构计算指在一个系统里同时使用多种不同类型的处理器来协同完成计算任务。CPU擅长复杂逻辑控制和少量数据的快速处理GPU擅长大规模数据并行计算FPGA擅长低延迟、可定制的流水线计算ASIC则在特定场景下做到极致的能效比。异构计算的关键不是堆硬件而是把任务按照特性拆分到最合适的计算单元上。深度学习需要异构计算的原因可以从训练和推理两个阶段来看。训练阶段核心操作是大量稠密矩阵乘法这正是GPU的强项而数据预处理、分布式通信调度、模型参数的更新逻辑依然由CPU负责。推理阶段更复杂高吞吐场景可以用GPU做批量推理低延迟场景可能需要FPGA定制流水线端侧场景只能依赖NPU或DSP。所以整个AI系统天然就是一个异构系统。如果你只是写“异构计算就是CPU加GPU”那只能拿一半分。加分点在于“为什么适用”和“任务如何切分”。我当时还顺手补了一句异构计算的本质是“用合适的计算单元做合适的事”这句话放在任何场景都能概括核心思想。2.2 线程层级与索引计算细节第二类必考是GPU线程模型。题目会让你简述CUDA的Grid、Block、Thread三级结构并说明它们和硬件SM、Warp的映射关系。标准答法是一个Kernel启动后对应一个GridGrid由多个Block组成Block由多个Thread组成Block会被调度到SM上执行一个SM可以同时驻留多个BlockGPU实际执行指令的最小单位是Warp通常包含32个线程Warp内的32个线程以SIMT单指令多线程方式执行即同一时刻执行同一条指令但各自处理不同的数据这里有一个非常容易错的细节线程的调度顺序。Block内的线程不是一次性全部执行而是按Warp为单位切分。如果一个Block有128个线程它会被分成4个Warp。Warp之间在GPU上轮流执行遇到访存等长延迟操作时会切换其他Warp来隐藏延迟。这个机制叫“线程级并行”配合“延迟隐藏”是GPU和CPU在设计哲学上最大的区别。我当时做这题的时候还补了一张线程索引的推算思路一个Block内第tid个线程的全局ID如果是一维结构就是blockIdx.x * blockDim.x threadIdx.x如果是二维Block则是blockIdx.x * blockDim.x threadIdx.x这是最低频的一层上面还有blockIdx.y * gridDim.x blockIdx.x这样逐层叠加的过程。虽然2018年的卷子没有直接要求写二维索引但清楚这个逻辑对后边的代码题是隐形加分。注意写CUDA代码时索引错误是最隐蔽的bug。程序不会报错只是结果不对而且通常在数据量大到一定程度才暴露。笔试现场切忌只写主体逻辑不验证索引边界建议每一步都推演一两个小数据量样例。3. CUDA编程题矩阵乘法与共享内存优化3.1 第一版朴素写法得基础分矩阵乘法是这套卷子的重头戏。题目一般给一个矩阵乘法的CUDA实现框架要求你补全内核或写出一个可运行版本。2018年的题我记得大致是给定两个1024x1024的float矩阵A和B写一个CUDA Kernel计算C A * B。最朴素的内核是这样的__global__ void matmul_naive(float* A, float* B, float* C, int N) { int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; if (row N col N) { float sum 0.0f; for (int k 0; k N; k) { sum A[row * N k] * B[k * N col]; } C[row * N col] sum; } }这个写法在功能上完全正确理论上也能跑但性能非常差。原因在于内层循环每次都直接访问全局内存而A和B的访存模式完全没有利用局部性。每次循环都要从显存读两个float一个矩阵乘法做下来全局内存访问次数是O(N^3)而理想情况应该是O(N^2)。为什么性能差核心在于GPU的运算单元比访存带宽“快得多”。2018年主流的P100/V100FP32算力都在10 TFLOPS以上但HBM显存带宽只有几百GB/s。这种算力和带宽的差距意味着如果每个计算都需要从全局内存取数计算单元大部分时间都在等数据GPU根本跑不满。要解决这个问题只有一条路利用共享内存和寄存器做数据复用把“重复从全局内存读数据”变成“从片上缓存高速读取”。3.2 分块优化和共享内存tile的写法分块矩阵乘法的思路是把大矩阵切成多个小Tile每个Block负责计算输出矩阵中的一个Tile。计算时先把A和B对应位置的Tile从全局内存拷贝到共享内存然后从共享内存中反复读取数据参与计算。#define TILE_SIZE 32 __global__ void matmul_tiled(float* A, float* B, float* C, int N) { __shared__ float As[TILE_SIZE][TILE_SIZE]; __shared__ float Bs[TILE_SIZE][TILE_SIZE]; int tx threadIdx.x; int ty threadIdx.y; int row blockIdx.y * TILE_SIZE ty; int col blockIdx.x * TILE_SIZE tx; float sum 0.0f; for (int t 0; t N / TILE_SIZE; t) { As[ty][tx] A[row * N t * TILE_SIZE tx]; Bs[ty][tx] B[(t * TILE_SIZE ty) * N col]; __syncthreads(); for (int k 0; k TILE_SIZE; k) { sum As[ty][k] * Bs[k][tx]; } __syncthreads(); } C[row * N col] sum; }这个版本比朴素版快了一个数量级。关键优化点有三个数据复用A矩阵的每个元素会被B矩阵同一Tile内的所有线程复用通过共享内存将全局内存访问次数从O(N^3)降为O(N^2 / TILE_SIZE)内存合并在拷贝As和Bs时相邻线程访问相邻地址满足全局内存合并访问的条件带宽利用率高同步保证__syncthreads()确保共享内存写入完成后再读取防止数据竞争笔试判分的时候面试官主要看你的内核里有没有用__shared__以及有没有正确处理同步。如果只写朴素版基本就是及格分写出分块版直接进入加分区间。3.3 Bank冲突真题计算与手算过程共享内存虽然快但也不是没有限制。它被划分成32个Bank每个Bank的宽度是4字节。当同一个Warp内的线程同时访问同一个Bank的不同地址时硬件会把这次访问串行化导致性能下降这就是Bank冲突。当年卷子里的经典计算题我记忆很深题干大概是这样共享内存中有float数组假设每个float占4字节某Warp内线程threadIdx.x访问第 threadIdx.x * 8 个float元素请计算这些地址分布到哪些Bank判断是否存在Bank冲突冲突几路。手算过程如下Bank编号的计算方式是地址按4字节为单位除以32取余数。第i个线程访问的是数组下标 threadIdx.x * 8 的元素换算成字节地址是 threadIdx.x * 8 * 4 threadIdx.x * 32。再除以4字节得到Bank单元编号threadIdx.x * 8。对32取余线程00 % 32 0线程18 % 32 8线程216 % 32 16线程324 % 32 24线程432 % 32 0线程540 % 32 8…规律逐渐清晰每4个线程映射到同一个Bank32个线程总共只用到8个Bank每个Bank被4个线程击中。这就是典型的4路Bank冲突实际访存耗时是理想情况的4倍。解决Bank冲突的常见手段有两种。一是给数组加Padding例如在第二个维度末位多申请一个float空间让原本对齐到同一Bank的地址错开二是在计算索引时乘以一个和Bank数互质的系数让地址重新散列。笔试时只要写出“加Padding”和“修改索引映射”任意一种再配合上面的手算过程基本就能拿满分。4. AI算子题卷积展开与访存优化4.1 im2col推导与空间代价百度这套卷子既然定位是AI异构计算卷积相关题目是绕不开的。2018年考的卷积题现在看其实是后来所有推理引擎的基础操作。题目问的是简述im2colimage to column的原理并分析这种方法的空间复杂度。im2col是把卷积运算转换成矩阵乘法的高效方法。假设输入是N个通道数为C、高H、宽W的特征图卷积核大小是K输出特征图的高宽设为H_out和W_out。im2col的过程是取出输入上每一个卷积窗口覆盖的区域将其展平成向量拼成一个二维矩阵。原始输入经过im2col后得到一个形状为 (C * K * K) × (N * H_out * W_out) 的大矩阵。这时候卷积核本身也可以展开成 K * K * C 行 × M列M为卷积核数量两者做矩阵乘法结果就是卷积输出。因为GPU上的GEMM通用矩阵乘法已经被优化到了接近硬件峰值所以im2col GEMM的流程在很长一段时间内是CNN推理的主流实现。空间代价是这道题的灵魂。展开矩阵的规模大约是原始输入数据的 K * K 倍。以3x3卷积为例im2col之后的数据量约是原始数据的9倍。对于大特征图这种空间开销是灾难级的。所以后续才有内存优化的变种比如在GEMM执行过程中动态做im2col不显式生成完整矩阵或者改用隐式im2col直接用地址计算代替物理拷贝。理解这个演进过程比单纯背im2col的定义重要得多。4.2 1x1卷积为什么在GPU上如此高效还有一道考1x1卷积的题当年回答的时候我只写了“相当于逐点线性变换”后来复盘发现其实没答到点上。1x1卷积在GPU上效率高的原因是它在做跨通道的信息融合而不改变空间尺寸。把特征图视为 N * H * W 个像素点每个像素点有C个通道值1x1卷积等价于对每个像素点做一次 C→M 的矩阵乘。在NCHW布局下这本身就是一次标准的GEMM而且K维度恰好是输入通道数C属于中小尺寸矩阵乘法GPU的Tensor Core对这种形状有专门优化。更重要的是1x1卷积天然适配内存合并访问。在NCHW布局下同一像素点的C个通道值在内存中不连续需要做转置或使用特殊的访存模式但如果把数据布局改成NHWC同一像素点的通道值是连续的读取时可以一次性载入寄存器再参与矩阵运算。这也是为什么很多推理引擎在GPU上会把数据转换成NHWC核心原因之一就是让1x1卷积和后续的GEMM算子访存更友好。实操心得笔试遇到这种“为什么在GPU上高效”的开放题千万别只答表面原因。尽量往数据布局、访存连续性、硬件指令集这三个方向靠分自然就高了。我当时只答了“矩阵乘”后来反思应该把NHWC布局的访存优势也写进去那才是GPU性能优化的真正命门。5. 并行算法题手写归约求和5.1 线程分歧是隐藏丢分点并行归约Reduction是异构计算笔试里出现概率极高的代码题它考察的是并发编程最核心的思维如何把一个整体任务拆成众多小任务然后高效地合并结果。题目一般是给定一个包含1024个浮点数的数组写出一个CUDA Kernel计算所有元素之和。这题看似简单但陷阱极多。最容易出问题的是线程分歧。如果你这样写__global__ void reduce_bad(float* input, float* output) { __shared__ float sdata[1024]; int tid threadIdx.x; int i blockIdx.x * blockDim.x threadIdx.x; sdata[tid] input[i]; __syncthreads(); for (int stride 512; stride 0; stride 1) { if (tid stride) { sdata[tid] sdata[tid stride]; } __syncthreads(); } if (tid 0) output[blockIdx.x] sdata[0]; }这个写法对吗功能上对。但性能上有很大的问题随着stride减小每轮参与计算的线程越来越少但Warp依然要按照32个线程一组的粒度执行。当stride小到一定程度时大量Warp里的线程只有部分是活跃的其他线程空转造成严重的线程分歧和资源浪费。更好的做法是从一开始就让线程处理多个元素先用循环进行一个block内的部分归约把数据量降到blockDim.x个再进入树状归约阶段。另一个优化点是让stride从blockDim.x/2开始而不是从总长度的一半开始这样能保证每一轮不浪费线程。还有一点是避免分支在关键循环里让所有线程执行相同的指令只是在数据选择上不同可以通过地址计算而不是if判断来处理边界条件减少控制流的开销。5.2 一个更稳的归约写法笔试现场推荐下面这个版本。它结合了“每线程多元素”和“交错归约”两种策略__global__ void reduce_good(float* input, float* output, int n) { __shared__ float sdata[256]; int tid threadIdx.x; int i blockIdx.x * blockDim.x threadIdx.x; float sum 0.0f; // 每个线程处理多个元素减少block数量提高数据复用 int stride gridDim.x * blockDim.x; for (int j i; j n; j stride) { sum input[j]; } sdata[tid] sum; __syncthreads(); // 树状归约 for (int s blockDim.x / 2; s 0; s 1) { if (tid s) { sdata[tid] sdata[tid s]; } __syncthreads(); } if (tid 0) output[blockIdx.x] sdata[0]; }核心改进有两个。第一每个线程先循环累加多个元素避免启动过多Block导致调度开销第二树状归约只对共享内存中的数据操作速度极快。这版代码还有一个隐藏知识点blockDim.x取256而不是取512或1024。原因是共享内存越大一个SM上能同时驻留的Block越少占用率会下降。256线程的block通常能让SM驻留足够多的block来隐藏访存延迟是实践中很均衡的配置。6. 现场调试与高频翻车点6.1 同步漏写为什么会导致随机错误笔试现场还有一类题目是程序改错给一段CUDA代码让你指出错误。当年最经典的陷阱就是漏写__syncthreads()。在矩阵乘法分块版本里如果没有在共享内存写完后加同步就会出现一个经典问题Block内不同的Warp执行进度不一致。Warp 0可能已经把数据写到了共享内存正打算读而Warp 1可能还没开始写。如果Warp 0抢跑去读读到的就是未初始化或上一轮残留的脏数据。更麻烦的是这种错误是随机且不可复现的。同一份代码有的机器上跑一万次没问题换个GPU型号或改用更大规模数据结果就错了。笔试现场没有GPU可跑所以只能靠“代码审阅”发现这类问题。我的经验是看到共享内存写入后一律条件反射地确认后面有没有同步这是保命习惯。另一个高频错误是核函数启动参数里没有检查边界。比如矩阵大小不是blockDim的整数倍时如果每个线程只计算一个元素那么grid外侧的线程就会越界访问。正确做法是在内核开头加上if (row N col N) { // 计算 }这个边界检查虽然会让代码多几行但它是鲁棒性的保障笔试中写出边界检查会明显加分。6.2 分析工具与查错顺序的建议如果要我推荐一套查错路径顺序是这样的先看Kernel有没有越界或同步缺失再看数据布局是否满足合并访问接着看Bank冲突和占用率最后才看指令级优化。网上很多人喜欢一上来就谈NVVP/Nsight的指标但我个人的体会是如果没有确认代码逻辑正确性能分析毫无意义。当年笔试最后有一道开放式题目问“如何定位CUDA程序性能瓶颈”如果你只知道答“用profiler看热点”得分有限。更系统的答法是先通过单元测试验证功能正确性再用有限的实验数据比如矩阵规模从小增大观察耗时增长曲线判断算法复杂度是否符合预期然后利用profiler对比实际访存带宽和理论峰值带宽的差距如果两者接近问题在计算如果差距大问题在访存最后针对访存问题检查缓存命中率、合并访问率和Bank冲突情况。这个排查思路放在今天依然通用而且适用于任何性能调优场景不局限于CUDA。7. 给后来人的复习建议与延伸方向这套笔试题虽然出自2018年但它的考点框架在AI基础设施领域一直有效。后来无论是我去面试其他大厂的AI加速岗还是在开源社区看一些高性能推理引擎的设计方案都能在当年这套卷子里找到原点。如果现在有同学准备类似岗位我给三点建议第一不要只刷题一定要动手写CUDA。哪怕手头没有GPU也可以用在线平台跑小样例。只有亲手写过一遍你拿到题目的时候才知道哪些地方容易出边界错误哪些优化是纸上谈兵。第二重视内存布局和访存模式这比背诵各种并行模式重要得多。绝大多数性能瓶颈不是算力不够而是数据喂不饱计算单元。第三对卷积、矩阵乘、归约、转置这些经典算子至少要能手写一套可运行版本同时能解释清楚优化前后访存次数的变化量级。另外异构计算这些年也在快速演进。从早期单纯CPUGPU到引入Tensor Core、NPU、可重构计算再到如今围绕大模型训推的算力集群设计本质都是在“任务切分”和“数据搬运”之间做权衡。理解了2018年这套笔试题想考的东西后续学任何新硬件、新框架都会觉得有迹可循。最后一个个人经验笔试答开放题时哪怕不会完整推导也要把“我理解的瓶颈是什么”“我打算从哪个方向优化”写清楚。这类题看的不是你记住了多少现成答案而是面对一个陌生性能问题时你有没有一套自己的分析路径。这种能力不是短期突击能练出来的需要平时多读代码、多看硬件白皮书并且在踩坑中不断复盘。
