手写 CUDA vector add kernel 时，为什么需要边界判断，为什么通常不用 shared memory，block size 和 warp 有什么关系？｜百度 C/C++面经解析

60 秒回答模板

我会先写出 vector add 的基本模型：每个线程负责一个或多个元素，global id = blockIdx.x * blockDim.x + threadIdx.x，如果 id < n，就执行 c[id] = a[id] + b[id]。边界判断必须存在，因为 grid size 通常按 ceil(n / blockDim.x) 向上取整；最后一个 block 里的部分线程可能对应 id >= n，如果不判断就会越界读写。这个 kernel 通常不用 shared memory，因为每个输入元素只读一次、输出只写一次，线程之间没有数据复用或协作归约；把数据先搬进 shared memory 只会增加一次读写和同步，反而可能变慢。性能上它主要是 memory bound，要保证相邻线程访问相邻地址，使 global memory coalescing 好。block=256 表示每个 block 有 256 个线程；如果 warp size 是 32，那么一个 block 有 8 个 warp。选择 256 是常见经验值，通常能提供足够并行度并兼顾调度开销、occupancy 和寄存器资源，但不是唯一答案；要结合数据类型、寄存器、访存模式和 profiler 验证。

考点 索引公式

难度 真实面经题

回答目标 让候选人能用一个简单 vector add 讲清 CUDA 线程模型、越界安全、访存效率、shared memory 适用条件和 block/warp 基础。

深入解析

先说清线程到元素的映射

最基础的一维 vector add 会把数组下标映射到 CUDA 线程。线程的全局下标通常是 blockIdx.x * blockDim.x + threadIdx.x，每个线程处理一个元素；更通用的版本可以用 grid-stride loop，让一个线程按 stride = blockDim.x * gridDim.x 处理多个元素。这个映射是解释边界判断和访存模式的前提。

边界判断保护最后一个 block

数组长度 n 往往不是 blockDim.x 的整数倍。为了覆盖全部元素，host 端会把 gridDim.x 设成向上取整，这会让最后一个 block 中有一部分线程的 id 超过 n - 1。if (id < n) 的作用是避免这些线程访问 a[id]、b[id] 或 c[id] 时越界，保证正确性和内存安全。

shared memory 要看复用收益

Shared memory 适合一个 block 内多个线程反复使用同一批数据，例如 tile GEMM、卷积或归约。Vector add 中每个 a[i] 和 b[i] 通常只被一个线程读一次，没有跨线程复用；搬到 shared memory 需要额外 global load、shared store、shared load，可能还要同步，收益不足。

性能重点是全局内存访问

Vector add 每个元素只做一次加法，却要读两个输入、写一个输出，算术强度很低，常见瓶颈是显存带宽。优化重点不是增加计算技巧，而是让相邻线程访问连续、对齐的地址，形成良好的 coalesced global memory access，并避免不必要的中间写回。

block size 和 warp 的关系

CUDA 线程以 warp 为基本调度单位，常见 warp size 是 32。block=256 意味着一个 block 包含 256 个线程，也就是 8 个 warp。block size 通常选为 warp size 的整数倍，避免大量空 lane；但最终选择还要考虑每个 SM 可同时驻留多少 block、寄存器使用、shared memory 使用和并行度。

验证要覆盖正确性和性能

正确性上要测 n 小于 blockDim、n 等于整倍数、n 非整倍数、n 为 0 或很大数组等情况。性能上要观察带宽、global load/store efficiency、occupancy、kernel launch overhead 和不同 block size 的耗时。对于这样简单的 kernel，代码写对只是第一步，能解释为什么这么写更重要。

易错点

只写 c[i] = a[i] + b[i]，没有解释 grid 向上取整后为什么最后一个 block 会越界。
认为 shared memory 总比 global memory 快，所以任何 CUDA kernel 都应该先搬 shared memory。
把 block=256 理解成 256 个 warp，而不是 256 个线程、通常 8 个 warp。
忽略 coalesced global memory access，线程映射导致相邻线程访问不连续地址。
只追求更大的 block size，没有考虑 occupancy、寄存器、调度和硬件上限。
把基础手写题答成背 API，不会用边界、访存和 warp 调度解释代码设计。

面试官追问

如果去掉 id < n 会发生什么？

当 n 不是 blockDim 的整数倍时，最后一个 block 的多余线程会访问数组边界外的地址，可能导致错误结果、非法内存访问或难以复现的内存破坏。

什么时候 vector add 可能需要 grid-stride loop？

当数组很大或希望限制 gridDim 时，可以让每个线程按总线程数为 stride 处理多个元素。这样 kernel 对任意 n 都可扩展，也便于复用同一个 launch 配置。

block size 选 128、256、512 哪个一定最好？

没有固定最优。256 是常见起点，但最佳值取决于硬件、寄存器、occupancy、访存模式和数组规模。简单 vector add 通常要通过 profiler 或 microbenchmark 比较。

coalescing 在 vector add 里如何体现？

如果连续 thread 访问连续的 a[id]、b[id]、c[id]，一个 warp 的访问可以合并成较少的内存事务，带宽利用率更高。若访问 stride 很大或离散，带宽会下降。

为什么说 vector add 是 memory bound？

每个元素只做一次加法，却至少读两个数、写一个数，计算量相对访存量太低。即使计算单元很强，速度也常被显存读写带宽限制。