这套CUDA基础练习题已经做到第6套了。前5套发出去之后我收到最多的反馈不是“答案不对”而是“我以为自己会一做题才发现很多概念只是眼熟”。这恰恰是我坚持做这个系列的原因——CUDA入门最不缺教程缺的是能逼你动手推算的检验题。这一套我刻意加大了手算题和易错判断题的比重覆盖线程索引换算、内存分层、同步与原子操作、性能调优基础、错误调试五个模块共23题。全部做完大约需要40分钟建议先别急着翻解析动笔算一算再对照参考答案这样暴露出来的问题才是真问题。1. 这套题怎么用先限时自测再对照解析1.1 出题范围与设计思路CUDA的基础知识看似就那几块实际学起来特别容易“眼高手低”。概念课上听懂了一到写代码就不知道索引该乘还是该加面试被问到共享内存和全局内存的区别能背出几条但一问“为什么共享内存快”就卡壳。所以第6套题我没有追求偏题怪题而是把所有内容收敛到五个最核心的模块线程组织与索引计算、内存体系、同步机制与原子操作、性能调优基础、错误处理与调试。设计题目时我刻意做了三件事第一增加手算题逼你实打实算几个索引而不是看一眼公式觉得会了第二增加判断题很多错误认知在选择题里往往能被运气掩盖判断题更考验概念边界第三参考答案全部带解析每道题都会说清楚“考的是什么”和“容易错在哪”这样才能起到巩固知识的作用。1.2 建议的做题流程我的建议是拿一张纸一支笔给自己限时40分钟闭卷作答。不会的题果断跳过并做好标记不要边做边翻资料那样检验效果会大打折扣。做完之后先别急着对答案把不确定的题回去再想一遍然后才看后面的参考答案。对答案的时候重点看两样东西一是解析里提到的“常见错误”看看自己是不是刚好掉进同一个坑二是把错题对应的知识点写在一张卡片上比如“blockIdx和threadIdx的换算”“__syncthreads的同步边界”哪怕只是几个关键词也比反复读教程有用。如果某个模块的错题超过一半说明这一块需要重回基础而不是继续刷题。2. 第一模块线程组织与索引计算的手算能力2.1 为什么索引换算是第一道坎线程组织是CUDA编程模型的地基而索引换算是地基里的钢筋。一个核函数启动之后你面对的是成千上万个线程每个线程需要知道自己负责哪部分数据。CUDA不会像单线程程序那样给你一个从0到N-1的全局编号它只提供blockIdx、threadIdx、blockDim三个量你必须自己拼出最终的全局索引。这个模块的5道题前4道是纯手算题第5道是简答题。我遇到过不少同学写代码时习惯把“blockIdx.x * blockDim.x threadIdx.x”当成万能公式粘贴一旦需要二维或三维线程块立刻出错。手算的过程能逼你重新理解这个公式到底在算什么。2.2 题目1-5第1题填空题一个grid包含4个block每个block包含128个线程这个grid一共会执行____个线程实例。第2题计算题一维grid中blockIdx.x 3blockDim.x 64threadIdx.x 17求该线程在grid中的全局线性索引是多少第3题判断题threadIdx.x的取值范围是[0, blockDim.x - 1]。第4题计算题某个核函数启动时的blockDim (32, 8, 4)某线程的threadIdx (5, 3, 2)。按x维度变化最快的线性化方式该线程在block内部的线性ID是多少第5题简答题为什么CUDA核函数里经常要自己计算全局线程ID直接给每个线程一个“现成的全局编号”不行吗2.3 参考答案与易错点第1题参考答案512。解析4 × 128 512。这个题目看着简单真正想考的是你有没有养成“先看维度再相乘”的习惯。如果题目把block写成blockDim (128, 1, 1)那每个block的线程数就是128如果写成blockDim (16, 8, 1)那每个block的线程数就是16 × 8 128而不是16或8。网格里的总线程数永远要按维度乘积来算不能只看某一个值。第2题参考答案209。解析全局线性索引 blockIdx.x × blockDim.x threadIdx.x 3 × 64 17 209。很多人在这个公式上出错不是不知道乘blockDim.x而是把式子写反成“threadIdx.x × blockDim.x blockIdx.x”。可以自己画一条数轴block 0 占 0~63block 1 占 64~127block 2 占 128~191block 3 占 192~255threadIdx.x 17 自然落在192 17 209。只要画过一遍这个公式就不容易再记混。第3题参考答案对。解析CUDA的线程索引从0开始和C语言的数组下标一致。blockDim.x表示这一维度有多少个线程那么threadIdx.x的合法范围就是[0, blockDim.x - 1]。这个看起来是常识却有不少人在手写越界检查时写成了“threadIdx.x blockDim.x”导致多检查了一个根本不会出现的值虽然不影响正确性但暴露了对索引边界的模糊。第4题参考答案613。解析block内线性ID threadIdx.x blockDim.x × (threadIdx.y blockDim.y × threadIdx.z) 5 32 × (3 8 × 2) 5 32 × 19 613。这里的关键是“x最快变化”也就是把三维索引按x、y、z的顺序摊平到一维类似CPU端多维数组的按行优先存储。常见错误有两个一是把式子写成threadIdx.x threadIdx.y × blockDim.y threadIdx.z × blockDim.z维度配错了二是把z当成最内层写成z最快变化。建议把公式固定在脑子里最内层永远是x其次是y最外层是z。第5题参考答案要点CUDA的执行模型只告诉你这个线程属于哪个blockblockIdx、在block里的哪个位置threadIdx并不会提供一个全grid唯一的编号。要定位“我这个线程负责数组里的哪个元素”必须手动把blockIdx、threadIdx和blockDim结合起来换算。这种设计表面上是多了一步操作实际给了开发者极大的灵活性——你可以自由决定一个block包含多少线程、以什么样的形状覆盖数据而不是被固定编号绑死。3. 第二模块内存体系——笔试与面试最爱挖坑的地方3.1 一张表理顺五种存储的边界内存体系是整个CUDA学习的重中之重也是笔试面试里出现频率最高的考点。面试官不会只问你“共享内存快不快”他们会追问“快在哪里、慢在哪里、生命周期怎样、需不需要手动控制”。背概念是不够的得把每种存储的边界理清楚。存储类型作用域生命周期物理位置访问特点全局内存grid内所有线程从cudaMalloc到cudaFree设备显存容量大访问延迟高通常数百个周期共享内存block内所有线程block生命周期片上存储延迟远低于全局内存需手动管理和同步寄存器单个线程线程生命周期片上访问速度极快数量有限常量内存grid内所有线程同全局内存显存带缓存只读全线程读同一地址时效率高容量仅64KB本地内存单个线程线程生命周期实际位于显存寄存器溢出时使用访问延迟高表格下方的核心认知是存储类型从声明方式开始就分道扬镳。核函数里的普通局部变量放在寄存器__shared__修饰的变量放在共享内存通过指针访问的、由cudaMalloc分配的内存是全局内存。很多人把“本地内存”理解成片上存储其实是个大坑——本地内存在物理上是全局内存的一部分只是它只能被单个线程访问。3.2 题目6-10第6题选择题下列哪种内存在核函数内所有线程都可见且生命周期一直持续到显式释放为止 A. 共享内存 B. 全局内存 C. 寄存器 D. 常量内存第7题填空题在核函数中使用__shared__修饰的数组数据存放在____中。它比全局内存访问快的主要原因是____。第8题判断题常量内存的容量限制是64KB所有线程同时读同一个常量地址时访问效率会比较高。第9题简答题为什么说共享内存不是通常意义上的“高速缓存”它和CPU里的L2 Cache最本质的区别是什么第10题判断题寄存器是线程私有的生命周期覆盖整个线程所以寄存器溢出对性能没有任何影响。3.3 参考答案与解析藏在代码里的坑第6题参考答案B。解析全局内存的作用域是整个grid所有block里的线程都能访问生命周期从cudaMalloc开始到cudaFree结束由程序员显式管理。共享内存的生命周期只有block存活期间寄存器根本不可能跨线程共享常量内存虽然所有线程可读但它的生命周期实际上依附于全局内存分配方式不能单独“释放”。第7题参考答案共享内存片上存储、访问延迟低、带宽高。解析这是最基础的一组概念但有两个常见的延伸误区。第一共享内存的容量很小一个block能用的量通常只有几十KB级别写代码时必须精打细算第二共享内存没有自动初始化使用时需要手动填充而且它的访问还有个隐藏陷阱——bank conflict后面的模块会专门考。第8题参考答案对。解析常量内存的标称大小就是64KB这是硬件层面的限制。因为它有缓存机制当整个warp内的线程都读同一个地址时一次读取就能满足所有线程也就是所谓的广播读取效率很高。不过一旦warp内的线程读的是不同的常量地址就可能退化成多次串行访问性能反而不如预期。所以常量内存适合存放“全体线程都会用到同一份数据”的场景例如数值计算里的系数数组。第9题参考答案要点CPU里的Cache对程序员是透明的硬件自动维护数据一致性代码不需要干预。而CUDA的共享内存是显式管理的什么时候往里面填数据、什么时候用__syncthreads()做同步、block结束后数据作废这些都要程序员自己控制。与其说它是“缓存”不如说它是一块可控的高性能暂存区。这道题考的是对“显式管理”和“自动透明”的理解很多背书选手会在这一点上翻车。第10题参考答案错。解析寄存器确实在片上、速度极快但数量非常有限。如果一个线程需要的寄存器数量超过硬件能分配的上限寄存器就会溢出到本地内存。本地内存在物理上位于设备显存访问延迟比寄存器高几个数量级。我在实际项目里遇到过类似情况一个内核因为占用寄存器太多性能莫名其妙掉了一大截后来用编译选项限制寄存器数量并调整代码才把性能救回来。寄存器溢出这个问题笔试里常以判断题出现实际开发里则经常以“隐性性能杀手”的方式出现。4. 第三模块同步机制与原子操作的真实边界4.1 同步工具能做什么、管不了什么同步是并发编程里最需要敬畏心的话题。CUDA提供了多种同步手段但它们各有边界__syncthreads()只能同步一个block内的线程跨block的全局同步必须借助kernel结束或者协作组等高级机制原子操作能解决竞态但解决得并不“便宜”锁虽然能用但在GPU上代价很高一旦有线程持锁等待很容易演变成死锁或者性能灾难。很多初学者把同步当成“保险措施”哪里不对就加一个__syncthreads()结果要么卡死要么性能变差。这一模块的题目就是帮大家画出同步工具的边界。4.2 题目11-15第11题判断题__syncthreads()可以同步同一个grid内不同block之间的线程。第12题选择题一个warp内部分线程走if分支部分线程走else分支会发生什么 A. 两个分支并行同时执行 B. 先串行执行一个分支再串行执行另一个分支 C. 只有走到if分支的线程被执行else分支被忽略 D. 编译器直接报错第13题填空题多个线程需要对同一个全局变量做累加如果希望最终结果正确且不想使用锁可以使用CUDA内置的____操作。第14题判断题一个block内的线程先把数据写入共享内存数组紧接着另一个线程读取这份数据不需要任何同步因为共享内存访问速度足够快写入会立刻被读到。第15题简答题为什么某些原子操作在高竞争场景下性能下降得非常厉害4.3 参考答案与解析死锁风险和竞争代价第11题参考答案错。解析__syncthreads()是block内的栅栏同步它只保证“当前block里的所有线程都到达这个点之后大家才继续往下走”。不同block之间没有任何依赖关系硬件也不保证它们谁先谁后。如果真想实现跨block同步一个kernel结束前后的内存可见性、或者网格同步等机制才派得上用场但这些机制都更重不是简单一个函数调用能解决的。第12题参考答案B。解析GPU以warp为基本调度单位一个warp里的32个线程执行同一条指令。遇到if-else分支时硬件没办法让一部分线程执行if、另一部分同时执行else只能先让走if的线程执行此时走else的线程被掩蔽再反过来执行else路径。这个现象叫warp divergence分支分歧会导致两个分支串行执行是性能优化里必须警惕的问题。判断一个代码快不快不能只看有没有分支还要看分支是否发生在warp内部。第13题参考答案atomicAdd原子加。解析原子操作是硬件提供的read-modify-write原语保证多个线程同时操作同一个地址时不会互相覆盖。除了atomicAdd还有atomicSub、atomicMin、atomicMax、atomicCAS等。算力比较新的设备对这些操作的支持已经很强但使用时要清楚一点原子操作保护的是“单次读改写”的原子性不是“一堆操作整体不被打断”的互斥区后者仍然需要锁或者更重的机制。第14题参考答案错。解析共享内存访问延迟再低也不能保证内存一致性。block里的线程执行进度各不相同可能线程0刚把数据写到共享内存的某个位置线程1已经去读了读到的自然是旧值。必须用__syncthreads()把整个block的线程拉到同一个执行点才能安全地进行共享内存的生产者-消费者合作。我在写归约内核时经常强调这个习惯写完共享内存立刻同步读完共享内存如果要重新写再同步一次。多一次同步不算丑漏一次同步就是未定义行为。第15题参考答案要点原子操作必须在硬件层面保证“读-改-写”三个动作不可分割。大量线程同时争抢同一个地址时硬件只能让其中一个成功其余线程要么不断重试要么排队等待竞争越激烈无效开销越大。直观类比就是只有一个出入口的停车场车越多每辆车进出的平均时间越长。实际优化思路常见两种一是减少对同一个地址的争抢比如先用共享内存做block内的局部累加再写回全局二是用原子操作把结果分散到多个中间变量最后再做一次合并类似分桶直方图的思路。5. 第四模块性能调优基础题的思考路径5.1 性能题怎么答才不空泛性能优化是CUDA学习的终极目标之一但基础阶段的性能题不会问你某个内核怎么调优到极致而是考察你有没有建立“访问模式、资源占用、并行度”这几个基本视角。答这一类题时最忌讳上来就背“用共享内存、避免分支、合并访问”却说不清背后的机制。这4道题覆盖了三个最常见的入门性能指标合并访问、bank conflict、占用率。这三个概念看似独立实际都指向同一个核心怎么让硬件的工作效率变高而不是让线程数量看起来很多。5.2 题目16-19第16题选择题下列哪种全局内存访问模式对内存带宽最友好 A. warp内相邻线程访问相邻地址 B. warp内线程以较大步长跳跃访问 C. 单个线程反复访问相距很远的多个地址 D. 所有线程访问同一个地址第17题填空题共享内存访问中一个warp内多个线程同时访问同一个bank的不同地址时会产生____现象。第18题判断题核函数使用的寄存器数量越多通常越不利于提高占用率。第19题简答题什么是占用率occupancy为什么它重要但它不等于性能一定最优5.3 参考答案与解析从访存、占用率到Bank Conflict第16题参考答案A。解析GPU的全局内存以缓存行或事务为单位访问warp内相邻线程访问相邻地址时硬件可以把这些访问合并成一次或少数几次内存事务这是所谓的coalesced access合并访问 | 合并访存。选项B和C都会让事务数量明显增加带宽利用率下降。选项D虽然在某些场景下有广播机制但它是特殊优化手段不是通用情况下的“最友好”答案。判断一个内核访问内存是否高效第一反应就该是“warp内这32个线程访问的地址跨度有多大”。第17题参考答案bank conflictbank冲突。解析共享内存在硬件上被划分成32个bank每个bank在同一时钟周期只能服务一个地址。如果warp内多个线程同时访问不同地址但这些地址恰好落在同一个bank上访问就会被拆成多次串行操作这就是bank conflict。典型例子32个线程同时访问地址0、32、64……它们对32取模后都落在bank 0于是产生31次冲突如果访问同一个地址则会被广播合并不算冲突。解决思路通常是调整数据布局或填充让相邻线程访问到不同的bank。第18题参考答案对一般趋势。解析每个SM的寄存器文件总量是固定的如果每个线程占用寄存器太多能同时驻留的线程数就会下降进而影响占用率和延迟隐藏能力。但这里要强调“一般趋势”——有些内核通过较多寄存器换取了更少的指令数整体反而更快。所以这道判断题的考点是趋势判断不是绝对的优化准则。实际调优时我会先用编译选项报告查看寄存器占用再结合内核的实际瓶颈决定是否限制寄存器数量。第19题参考答案要点占用率是SM上活跃warp数与硬件最大可支持warp数的比值。高占用率意味着有更多warp可以轮流执行当一个warp在等待访存结果时调度器可以切换到其他warp从而隐藏延迟。但占用率高不直接等于性能最高原因很简单如果内核已经被内存带宽或指令吞吐量卡死再多warp也无法突破瓶颈甚至占用率太高导致共享内存或缓存被过度挤占还会起反作用。所以性能优化要把占用率当成参考指标而不是目标本身。6. 第五模块错误处理与调试基本功6.1 为什么错误经常“延迟暴露”CUDA新手调试时最容易遇到的一种诡异情况内核启动那一行的cudaError_t明明返回cudaSuccess程序却在后面某个cudaMemcpy或者cudaDeviceSynchronize时报错。根本原因在于内核启动是异步的核函数提交给GPU后CPU不会阻塞等待它执行完。内核里的非法内存访问错误是在执行过程中发生的你无法在启动那一刻立刻拿到结果只能等GPU执行到某个同步点或者主动调用同步API错误才会被抓到。这个机制本身是GPU高并发的基石但它也意味着错误处理的习惯必须跟着变。6.2 题目20-23第20题填空题CUDA内核启动是____的同步/异步调用____可以等待设备端所有任务完成。第21题选择题cudaGetLastError()与cudaPeekAtLastError()最主要的区别是什么 A. 返回值类型不同 B.cudaGetLastError()会清除错误状态cudaPeekAtLastError()不会 C. 两者功能完全相同 D.cudaPeekAtLastError()会阻塞CPU线程第22题判断题核函数内部访问了非法地址这个错误会在内核启动的那条代码之后立刻体现在返回值中。第23题简答题除了在代码里用printf输出中间值还有哪些常用的CUDA调试或检查手段6.3 参考答案与解析诊断手段与断言习惯第20题参考答案异步cudaDeviceSynchronize()。解析内核启动是异步的这是理解CUDA执行模型的关键点之一。cudaDeviceSynchronize()会阻塞CPU直到设备端所有之前提交的任务完成。实际开发中性能分析阶段不要在每行代码后都加这个函数否则等于把并行执行变成串行等待但在调试阶段它是最直接的问题定位手段。第21题参考答案B。解析两者都返回当前的错误码但cudaGetLastError()会同时把错误状态清零而cudaPeekAtLastError()只读不清除。调试时这个区别很实用如果你想检查内核启动后到底有没有潜在错误先用cudaPeekAtLastError()看一眼决定是否处理不要一上来就cudaGetLastError()否则错误状态被清除后后续逻辑可能再也查不到真正的问题。第22题参考答案错。解析正如前面所说内核启动是异步的错误往往要等到同步点才会暴露。这也是为什么很多团队会在关键API之后统一做一个错误检查宏。这里我贴一段自己常用的简版宏它本质上是在每个可能有问题的调用后面加一道保险#define CHECK(call) \ do { \ cudaError_t err (call); \ if (err ! cudaSuccess) { \ fprintf(stderr, CUDA error at %s:%d: %s\n, \ __FILE__, __LINE__, cudaGetErrorString(err)); \ exit(1); \ } \ } while (0)这个宏的好处是出错时能直接打印出文件位置和行号不用满屏找printf。正式项目里可以把exit(1)换成更友好的错误恢复或者日志记录逻辑但调试阶段就靠它快速揪出问题。第23题参考答案要点常用手段大致分四类。一是专门的内存检查工具比如compute-sanitizer能检测越界访问、未初始化内存、漏同步等运行时问题比肉眼盯着代码有效得多二是调试器比如cuda-gdb可以在设备端代码上打断点、查变量三是性能分析工具比如Nsight Systems看时间线、Nsight Compute看内核内部指标这类工具主要用于定位性能瓶颈而非普通bug四是写扎实的错误检查宏在所有关键API后统一检查返回码这是最基础也最容易被忽略的防线。7. 做完第6套我想多说几句第6套题到这里就结束了。按我的经验做题只是第一步合上参考答案之后的复盘才是真正拉开差距的地方。如果你在这套题里有任何一道错题我建议先把对应知识点用自己的话写一遍再拿一道同类的编程题或者配套练习验证一下。比如索引换算错了就再手算一个三维block的线程映射共享内存和全局内存分不清就写一个用共享内存做block内归约的小程序看效果。我在带人入门CUDA的过程中发现大多数人对内存层次能背出概念但没法解释“为什么共享内存快、为什么寄存器溢出是灾难、为什么原子操作高竞争会慢”。这些问题的答案都指向硬件结构而不是API本身。所以复盘时尽量往深处问一层逼自己去理解硬件的工作方式基础才能扎实。最后分享一个小习惯想确认自己是不是真的掌握了一个知识点就试着把它讲给一个完全不懂GPU的人听看他能不能听懂。如果讲到一半发现要甩出一堆术语才能圆过去那说明你自己其实还没吃透。这也是我坚持做这套练习题系列的初衷——很多知识点不落到纸上、不亲手算一遍永远不知道哪个环节是虚的。