又到CUDA练习系列更新的时候了。这次是第6套我特意把出题重心从“API记得熟不熟”挪到了“执行模型理不理解”上。理由很简单这几年帮别人review代码我发现十个报错里六七个都不是语法问题而是线程没对上、同步放错位置、内存访问没合并这些“看起来没毛病、跑起来出怪事”的基础问题。基础不牢后面写归约、写矩阵分块、写多流优化全是踩着坑往前走。这套练习题适合几类人刚学完CUDA基础、想检验自己掌握程度的初学者准备面试、需要快速自查知识盲区的求职者以及正在项目里被莫名其妙bug折磨、想系统排查基本功的开发者。所有题目我都按“题目—答案—解析”的方式放在一起建议你先把答案挡住做一遍再对解析。题目不难但坑不少尤其第6题和第7题是这五套里我见过翻车率最高的两个考点。1. 这套题的出题思路为什么要做“检验型”练习1.1 系列第6套重心从API调用转向执行模型前几套题我分别覆盖了内存布局、线程同步、矩阵运算、API细节等主题。这第6套我刻意调整了比例概念记忆类题目只保留一小部分更多篇幅给了“给一段代码让你找问题”和“给你一个场景让你算配置”这两种题型。原因是我发现一个普遍现象很多人能背出blockIdx.x * blockDim.x threadIdx.x这个公式但给他一个二维grid、二维block让他算某个线程的全局ID他就开始乱套很多人知道__syncthreads()是同步但不知道它为什么不能放在if分支里很多人听说过bank冲突但真给他一段共享内存访问代码他看不出来冲突了几路。这套题就是在这些“知道但不会用”的灰色地带做文章。CUDA真正难的不是语法而是你写的kernel在硬件上到底怎么被执行。理解执行模型比多记十个API函数有用得多。1.2 题目分布与自测方法做题之前先看这张表清楚每道题在考什么方便最后对照自测结果。题号考察点题型难度第1题线程层级与硬件限制判断题基础第2题多维grid/block索引计算选择题基础第3题启动配置与边界检查代码阅读基础第4题内存层次模型选择题基础第5题warp划分与分支开销计算题中等第6题共享内存bank冲突计算题较难第7题__syncthreads正确用法改错题较难第8题两阶段归约设计设计题中等第9题流并发与异步拷贝判断题中等第10题CUDA错误捕获改错题基础第11题访存带宽估算计算题中等建议给自己定个时间前8道题控制在30分钟内后3道题控制在15分钟内。不要一上来就看解析先凭直觉写答案。做错没关系关键是搞清楚错在哪一层——是概念没记住还是理解偏了。1.3 参考答案的正确打开方式很多人做练习题喜欢“先看答案再做题”这等于白做。我提供参考答案的目的是让你做完之后有一个对照而不是让你抄。正确流程是先做题→对答案→看解析→回到代码里验证。验证这个动作特别重要尤其是索引计算和bank冲突这两类题你用printf在kernel里打印几个值或者用性能分析器看一眼共享内存报告比解析文字讲十遍都管用。这系列练习有个特点参考答案不只是给一个结果还会告诉你这个结论在什么条件下成立。CUDA的很多规格限制在不同硬件架构上不一样所以解析里我会特别标注哪些是“通用规则”哪些是“旧限制别记一辈子”。2. 线程层级与索引计算三道送分又容易翻车的题2.1 第1题线程层级与硬件限制题目判断题在CUDA编程模型中一个grid最多只能有65535个block而一个block最多可以包含1024个线程。这道题十个有八个会答“对”因为很多中文教程里写着“grid的block数上限是65535”。这个数字是早期架构的限制拿到现在看已经过时了。答案应该是后半句对前半句错。block最多1024个线程这个限制在现役主流架构上依然成立部分新架构支持更多但要查具体规格。但grid的block数上限并不是65535。在当前的CUDA编程模型里gridDim.x的上限是2^31-1gridDim.y和gridDim.z的上限是65535。也就是说x方向的block数量上限已经远远超过65535这个旧值。这里想提醒大家的是CUDA的规格限制是跟着硬件架构走的不要背死数字。遇到不确定的限制去查该架构对应的规格表格比凭记忆写代码靠谱得多。顺带说一句block内线程数上限受制于寄存器、共享内存等资源不是你想开1024就一定开得满的。2.2 第2题多维grid/block的全局索引计算题目给定gridDim (8, 8, 1)blockDim (16, 16, 1)。请问第(bx, by)个block内第(tx, ty)个线程的全局线性ID计算公式是哪个选项A.(by * gridDim.x bx) * blockDim.x ty * blockDim.x txB.(by * gridDim.x bx) * blockDim.x * blockDim.y ty * blockDim.x txC.(bx * gridDim.y by) * blockDim.x * blockDim.y tx * blockDim.x tyD.bx * blockDim.x * blockDim.y tx by * gridDim.y * blockDim.x * blockDim.y ty正确答案是B。这个公式的核心是“先算block在grid里的线性编号再算线程在block里的线性编号最后乘起来”。二维grid按行主序排布block的线性编号是by * gridDim.x bx也就是先数完x方向所有的block再往上数y方向。block内部同理线程的线性编号是ty * blockDim.x tx。把两者拼起来全局ID (by * gridDim.x bx) * (blockDim.x * blockDim.y) ty * blockDim.x txA错在block内部只乘了blockDim.x漏掉了y方向的线程数。C把grid和block的x/y顺序全搞反了。D的语义混乱把线程维度和grid维度混着乘。我实际带人写代码时发现二维索引算错是最常见的隐性bug。建议你写个小kernel把blockIdx、threadIdx和算出来的全局ID直接printf出来跑一遍、对一遍数自然就记住了。2.3 第3题向量加法的启动配置与边界检查题目现有内核vectorAdd它按int i blockIdx.x * blockDim.x threadIdx.x的方式处理数组。如果N 100000每block开256线程正确的启动配置是什么内核里为什么必须写边界检查正确启动行是vectorAdd(N 255) / 256, 256(d_a, d_b, d_c, N);(N 255) / 256就是向上取整。100000除以256约等于390.625向上取整是391个block。391乘以256等于100096比100000多出96个线程。这96个线程不能访问数组末尾之外的内存所以kernel内部必须有if (i N) { d_c[i] d_a[i] d_b[i]; }没有边界检查多出来的线程会直接访问d_a[100000]及之后的内存。后果分两种运气好那块内存是可读的程序不报错但结果莫名其妙运气差直接触发非法访问程序崩溃或者计算设备被重置。更要命的是这种错误经常是异步的不在kernel这一行报出来而是拖到后面某个cudaMemcpy才炸。很多教程里喜欢写gridDim.x N / blockDim.x这对N恰好是blockDim.x整数倍的情况没问题但实际数据量哪有那么巧。统一用向上取整加边界检查才是工程上能过审的写法。3. 内存体系与Warp执行模型概念题这样考3.1 第4题全局内存、共享内存与寄存器的区别题目下列哪个说法是正确的A. 共享内存位于GPU片外同一grid内的所有block共享同一块共享内存空间。B. 全局内存对所有线程可见访问延迟高通常需要合并访问才能获得较好带宽。C. 寄存器变量在kernel执行结束后仍然保留供下一次kernel继续使用。D. 常量内存和全局内存一样可以在kernel内部写入只是写入速度慢。正确答案是B。逐个拆A错在“片外”和“所有block共享”。共享内存在片上而且是per-block私有的block A看不到block B的共享内存。C错在寄存器是线程私有资源kernel一结束寄存器内容就没了。D错在常量内存不允许在kernel里写它只能通过host端API在kernel启动前写入kernel里只能读。这道题真正的考点是不同内存的可见性和生命周期决定了你该用什么内存存什么数据。全局内存慢但大、谁都能访问共享内存快但小、一个block内可见寄存器最快但线程私有、容量极小。写kernel之前先想清楚数据的流向比动手后再调优省事得多。3.2 第5题warp大小、分支与非整数倍block题目一个block包含127个线程。硬件会把它划分成多少个warp最后一个warp有多少个活跃lane如果warp内的线程发生分支发散最坏情况下性能损失大约是多少倍先算warp数127 / 32向上取整等于4。前3个warp各32个线程第4个warp只剩127 - 3×32 31个活跃lane第32个lane是空闲的。这个知识点很容易被忽略block的线程数不是32的整数倍时硬件依然按warp分配调度资源最后一个warp会用不满。如果你开100个线程的block实际硬件要调度4个warp有28个lane空转。关于分支发散当一个warp内不同线程走了不同分支时硬件会把每个分支路径串行执行一遍。最坏情况下32个线程各自走一条路径性能损失接近32倍。这就是为什么GPU编程里常说“让同一个warp内的线程尽量走同一个分支”。实际经验是block大小尽量取32的整数倍比如128、256、512。非整数倍虽然能跑但白白浪费了warp槽位。3.3 第6题共享内存Bank冲突实战计算题目一个block有32个线程共享内存里有一个float s[128]。如果线程threadIdx.x访问s[threadIdx.x * 4]会发生几路bank冲突如果改成访问s[threadIdx.x]呢共享内存按bank组织32个bank对应32个连续的4字节字。线程访问word_index时bank word_index % 32。硬件一次可以同时给不同bank的线程服务但如果有多个线程访问同一个bank就必须串行。先看s[threadIdx.x * 4]。线程0访问word 0bank 0线程1访问word 4bank 4线程8访问word 32bank 0。因为4 × 8 32线程0、8、16、24都落在bank 0上。同样的分析可知每个bank被4个线程命中所以是4路bank冲突访问速度约为无冲突时的4倍。再看s[threadIdx.x]即线程0访问word0、线程1访问word1……每个线程落在不同bank零冲突。补充一个经典场景二维数组按列访问。假设float tile[32][32]线程访问tile[threadIdx.x][0]也就是取第0列的所有元素。每个元素地址是row * 32 0模32之后全是bank 032个线程挤在一个bank上32路冲突。解决办法是给每行加一个padding比如float tile[32][33]地址变成row * 33 col模32后是(row col) % 32不同row落到不同bank冲突就消除了。bank冲突不报错只是慢。你跑性能分析器才能看到共享内存报告里那一排红色数字。第6题这类计算建议你在心里把word_index % 32当成乘法表一样背熟写访问共享内存的代码前先心算一遍。4. 归约、同步与原子操作从死锁到设计4.1 第7题__syncthreads放错位置的经典死锁题目下面这段归约代码里有隐藏的同步错误请指出问题__shared__ float s[256]; int tid threadIdx.x; s[tid] input[tid]; for (int stride 128; stride 0; stride 1) { if (tid stride) { s[tid] s[tid stride]; __syncthreads(); } }问题在最后一轮。当stride 1时只有线程0进入if然后它停在了__syncthreads()等待其他线程。可是其他线程根本没进这个if不会执行同步指令。结果就是线程0永远等下去整个block死锁。正确写法是把__syncthreads()挪到if外面for (int stride 128; stride 0; stride 1) { if (tid stride) { s[tid] s[tid stride]; } __syncthreads(); }__syncthreads()的使用条件是“同一个block内的所有线程都要执行到这一行”。所以只要代码路径上有某个线程可能跳过它就是定时炸弹。同理如果kernel里有人提前return而其他线程继续走到同步点也会等死。这也是为什么很多代码规范里要求kernel中不要随意提前return宁可包一层if (...) { }也别让部分线程退出。4.2 第8题百万级数组归约的两阶段设计题目给定一个长度为1,000,000的float数组用block256的配置做归约求和。第一层启动多少个block第一层归约完成后还剩下多少个部分和下一步怎么处理block数(1000000 255) / 256向上取整等于3907。每个block内部完成一次共享内存归约输出一个部分和写回全局内存的partial[blockIdx.x]。第一层结束后全局还剩3907个部分和。这时不能再让3907个block每个再归约一个数玩“套娃”了数据量已经很小。标准做法是开第二个小kernel用一个block256线程把这3907个数继续归约成1个最终和或者把partial拷贝回host用CPU做最后几千个数的求和代价也很低。那能不能在第一层直接用atomicAdd把每个block的部分和加到全局变量上可以但要注意3907个block同时抢一个全局地址做原子操作会严重串行化。第一层归约的并行度白白被原子操作拖垮。更合理的方式是第一层老老实实算出3907个部分和第二层再做最终归约。如果只有几十个block原子操作确实简单可行但扩展到几千个block原子操作的代价就上来了。做设计题的时候把“并行度”和“串行瓶颈”这两件事想清楚比背代码重要。5. 流并发与错误捕获项目里最容易出bug的两块5.1 第9题流并发不等于自动重叠题目判断题只要代码里使用了多个stream并把cudaMemcpyAsync放到每个stream里数据传输和kernel计算就一定能重叠执行。答案是不一定这个说法错。多流只是必要条件远不是充分条件。想真正重叠至少满足三条第一host端内存必须是页锁定内存固定内存。如果你用普通的malloc分配内存cudaMemcpyAsync实际上会退化成同步拷贝因为驱动没法在不锁定页面的情况下发起DMA操作。页锁定内存需要用专门的内存分配API在host端分配。第二不同stream上的操作不能有数据依赖。你要是stream 1里拷贝完整数据、stream 2里就靠这份数据算kernel那stream 2必然要等stream 1重叠无从谈起。正确的切分方式是把数组切成几段每个stream负责自己那一段的“拷贝→计算→拷回”。第三别用默认流。默认流有隐式同步特性它会堵住所有其他流导致你的多流编排全部失效。工程上通常都改用非默认流来组织并发。这三条我在实际项目里都踩过第一条没注意写了半天多流发现耗时和串行一模一样第三条更隐蔽代码看起来是两流并行实际被默认流卡成了排队执行。5.2 第10题写CUDA不查错误码等于裸奔题目看下面这段代码找出至少三个问题cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice); kernelgrid, block(d_a); cudaMemcpy(h_a, d_a, size, cudaMemcpyDeviceToHost);第一个问题三个CUDA API调用都没有检查返回值。CUDA几乎每个API函数都返回cudaError_t出错时会返回非cudaSuccess值同时通过cudaGetErrorString可以拿到错误描述。不检查等于程序裸奔。第二个问题kernel是异步启动的它不会在原地等待kernel执行完。如果kernel启动参数有误比如block数超限这个错误不会立刻抛出而是留在错误队列里直到后续某个同步API调用时才被捕获。所以kernel启动后应该立刻调用cudaGetLastError()把所有“异步排队”的启动错误抓出来。第三个问题kernel执行期的错误需要等它真正跑完才能发现。最稳妥的做法是在把数据拷回host前调用cudaDeviceSynchronize()这样一个同步点既能暴露kernel的执行错误又能保证后续cudaMemcpy读到的数据是kernel完成后写的。我自己的代码习惯是封装一个错误检查宏#define CHECK_CUDA(call) do { \ cudaError_t err (call); \ if (err ! cudaSuccess) { \ fprintf(stderr, CUDA error: %s at %s:%d\n, \ cudaGetErrorString(err), __FILE__, __LINE__); \ exit(EXIT_FAILURE); \ } \ } while (0)然后所有调用都写成CHECK_CUDA(cudaMemcpy(...))。虽然代码看起来啰嗦但在项目里能省下几天的排查时间。5.3 第11题一条带宽计算题带你看清优化空间题目某个kernel要读入4,000,000个float计算后写回4,000,000个float实测耗时0.5ms。请问有效访存带宽是多少如果某款GPU的理论带宽约在200GB/s量级这个kernel还有多少优化空间先算访存量。读入4,000,000 × 4字节 16MB。写回同样是16MB。合计32MB。耗时0.5ms也就是0.0005秒。带宽 32MB / 0.5ms 64,000MB/s 64GB/s。对比理论带宽200GB/s量级当前只跑到三分之一左右说明还有明显优化空间。最常见的原因有两个一是内存访问不合并。如果线程访问的地址不是连续的带宽利用率会大幅下降。合并访问要求同一个warp里的32个线程访问的地址尽量落在一个连续的内存段里。二是没有用向量化加载。如果改成每个线程一次处理4个连续的float类似float4可以减少访问事务数量让带宽利用率上一个台阶。把数组看成float4组成的数组线程ID映射到float4索引上一个线程一次读写16字节效果立竿见影。这道题的价值在于让你养成“估算带宽”的敏感度。拿到一个kernel先算它理论最少访存量再对比耗时心里就有个底现在是瓶颈在访存还是计算还是同步开销。有了底优化方向就不会跑偏。6. 自测结果怎么解读附上后续学习路径6.1 按考察点对症下药做完这11道题你可以按下面的表对照看自己哪些地方需要补。题号考察点如果错了建议补什么第1题线程层级限制查官方文档的硬件规格表格只看当前架构那部分第2题多维索引计算写kernel用printf打印blockIdx/threadIdx手工验算几个值第3题启动配置与越界练习(N blockDim - 1) / blockDim向上取整写法所有kernel加边界检查第4题内存层次画一张各内存的容量/延迟/可见性对照表贴在显示器旁边第5题warp与分支重点理解SIMT执行模型拿127线程的block做一次实际运行第6题bank冲突用性能分析器的共享内存报告验证自己的冲突判断反复练到心算准确第7题__syncthreads记住“所有线程必须到达同一点”遇到死锁先查同步位置第8题归约设计手写完整的两阶段归约跑通后再尝试用原子操作对比性能第9题流并发用页锁定内存非默认流写一个分段拷贝/计算重叠的小demo第10题错误捕获把CHECK宏加进自己的所有CUDA项目形成习惯第11题带宽估算学会算访存量尝试用float4优化一个简单kernel并实测提速比如第6题做错说明你对共享内存的bank映射没有形成直觉。这类问题不是多看两篇文章就能解决的必须手动算几组word_index % 32再配合性能分析器验证才能形成条件反射。第7题做错说明你对__syncthreads的本质还没有理解透建议去搜一下官方文档里关于同步的约束原话把“block内所有线程必须执行到同一行”这句话抄三遍。6.2 我给“刚入门开发者”的最后建议最后分享一点我做这套练习的体会。我让几位朋友做过前五套题发现一个有意思的规律第3题这种“背过就会”的正确率特别高而第6题、第7题这种“需要想明白硬件在干什么”的错误率特别高。这不是智力问题是学习顺序问题——很多人学CUDA是背着一个一个API函数学的没有建立起“代码运行在硬件上”的画面感。如果你也卡在这两道题上我的建议很朴素别急着看更花哨的优化技巧先回去把线程层次、内存层次、warp执行这三个最基本的东西搞成一幅能在脑子里跑起来的图。什么时候你看到blockDim.x会条件反射想到这对应一个block里x方向的线程数看到共享内存访问能下意识算出bank看到if里的__syncthreads()会心里发毛那你写kernel的基本功就真的过关了。这套题的目的就是帮你找到那个该补的缺口。