企业官网建设流程全解析

这套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 = 3，blockDim.x = 64，threadIdx.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~63，block 1 占 64~127，block 2 占 128~191，block 3 占 192~255，threadIdx.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的执行模型只告诉你这个线程属于哪个block（blockIdx）、在block里的哪个位置（threadIdx），并不会提供一个全grid唯一的编号。要定位“我这个线程负责数组里的哪个元素”，必须手动把blockIdx、threadIdx和blockDim结合起来换算。这种设计表面上是多了一步操作，实际给了开发者极大的灵活性——你可以自由决定一个block包含多少线程、以什么样的形状覆盖数据，而不是被固定编号绑死。

3. 第二模块：内存体系——笔试与面试最爱挖坑的地方

3.1 一张表理顺五种存储的边界

内存体系是整个CUDA学习的重中之重，也是笔试面试里出现频率最高的考点。面试官不会只问你“共享内存快不快”，他们会追问“快在哪里、慢在哪里、生命周期怎样、需不需要手动控制”。背概念是不够的，得把每种存储的边界理清楚。

表格下方的核心认知是：存储类型从声明方式开始就分道扬镳。核函数里的普通局部变量放在寄存器；__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 conflict（bank冲突）。解析：共享内存在硬件上被划分成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的人听，看他能不能听懂。如果讲到一半发现要甩出一堆术语才能圆过去，那说明你自己其实还没吃透。这也是我坚持做这套练习题系列的初衷——很多知识点，不落到纸上、不亲手算一遍，永远不知道哪个环节是虚的。