想象你是一家“GPU 大排档”的主厨。周日晚上突然涌进 10 万道菜，后厨里站着一千名帮厨。你没法挨个叮嘱每人做什么，只能写一张能贴在墙上让所有人执行的菜单，然后喊一声“开工”。CUDA 里，这张菜单就是 kernel 核函数，那些帮厨就是 GPU 上成千上万的线程。核函数写得好不好，决定这顿饭是十分钟上齐还是两小时崩盘。本文带你走一遍 kernel 编写最重要的七个环节：怎么发菜单、怎么给帮厨编号、怎么从仓库搬货、怎么用小组工作台、怎么处理货不够分的边界、怎么让 32 人小队不窝工，以及怎么把后厨从“能出菜”迭代成“翻台王”。全程用生活直觉打底，再落到代码和公式。

岗位关联

会写 kernel，你就掌握了亲手调教 GPU 的能力：从只会调框架 API，到能写出高效并行算子，AI 工程师的含金量直接上一个台阶。

目录
  1. 可视化演示
  2. 声明与传参：菜单怎么发，人手怎么配
  3. 线程索引计算：帮厨怎么认自己的工号
  4. 合并访问优化：仓库搬货要整整齐齐排队
  5. 共享内存分块：小组工作台让协作更快
  6. 边界条件防护：货不够分时别让帮厨空手乱抓
  7. Warp 发散规避：32 人小队要齐步走
  8. 迭代性能调优：从能出菜到闪电出菜
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

展示CUDA核函数核心优化：全局内存合并访问、shared memory分块、warp发散问题与渐进式优化迭代。支持四种模式步进切换。

本节配有交互动效（进入视口后自动加载，约 3KB）。

'+captions.coalesce+'

声明与传参：菜单怎么发，人手怎么配

后厨接到 10 万道菜，主厨不可能挨个告诉每名帮厨做什么。你只需要写一张通用菜单贴在墙上：所有人都按这份菜单执行，每个人通过拿到的工作编号负责自己的那一份。在 CUDA 里，这张菜单就是 __global__ 修饰的 kernel 函数，它不在主机端直接运行，而是被 GPU 上大量线程同时执行。

调用 kernel 时，我们还要告诉后厨开几组灶台、每组站几人，这就是 kernel<<<grid, block>>>(args) 中的执行配置。grid 决定网格里有多少个线程块，可以理解为“开几组灶台”；block 决定每个线程块里有多少个线程，可以理解为“每组灶台配几个帮厨”。参数列表则相当于菜单上的原料清单，每个线程都能读到这些参数。

要确定 grid 开多大，最常用的方法是按总工作量除以每组人数再向上取整：

菜单发出去了，灶台也排好了，但你发现一个更重要的问题：每个帮厨都站在灶台前，却不知道自己到底该负责哪一份菜。别急，这正是下一步要解决的。

gridDim.x=⌈NblockDim.x⌉\displaystyle\text{gridDim}.x = \left\lceil \frac{N}{\text{blockDim}.x} \right\rceil
通俗说明：如果每组最多 8 人，却有 100 道菜，就需要 13 组灶台，最后一组只有 4 人干活。这个公式就是“总共需要开几组灶台”。
  1. N总工作量，可以理解为需要处理的数组元素个数或菜的份数。
    N\displaystyle N
  2. blockDim.x每个线程块包含的线程数量，也就是一组灶台能容纳的帮厨人数。
    blockDim.x\displaystyle \text{blockDim}.x
  3. 向上取整符号\lceil \cdot \rceil 表示即使最后一组只有几个人，也必须单独开一组，确保所有工作都被覆盖。
    ⌈NblockDim.x⌉\displaystyle \left\lceil \frac{N}{\text{blockDim}.x} \right\rceil
  4. gridDim.x网格在 x 方向上的线程块数量，也就是菜单最终需要的灶台组数。
    gridDim.x\displaystyle \text{gridDim}.x

线程索引计算：帮厨怎么认自己的工号

每个线程执行的代码完全相同，但它们必须知道“我是谁、我在哪”，才能分配到属于自己的数据。这就像一家大后厨给所有帮厨发工号牌：工号由“第几区、第几组、组内第几个”三部分组成。

CUDA 为此提供了三个内置变量：blockIdx.x 表示当前线程块在整个网格中的编号，blockDim.x 表示每个线程块有多少个线程，threadIdx.x 表示线程在块内的编号。通过这三个变量，我们可以算出全局唯一的线程号：

有了工号，每个线程就能从数组里取走属于自己的数据。不过取数据的方式也有讲究：如果 32 个并排的帮厨去仓库搬货时东一个西一个，整个后厨都会被拖慢。

global_tid=blockIdx.x×blockDim.x+threadIdx.x\displaystyle\text{global\_tid} = \text{blockIdx}.x \times \text{blockDim}.x + \text{threadIdx}.x
通俗说明：全局工号 = 前面所有组的工人总数 + 自己在组内排第几。四组每组 8 人，第三组第 2 个工人的工号就是 2×8+1=17。
  1. blockIdx.x当前线程块在网格中的编号，从 0 开始，相当于“第几组灶台”。
    blockIdx.x\displaystyle \text{blockIdx}.x
  2. blockDim.x每个线程块包含的线程数量，相当于“一组灶台的标准人数”。
    blockDim.x\displaystyle \text{blockDim}.x
  3. 乘法用块编号乘以每块线程数，得到“当前块之前已经排了多少个工人”。
    blockIdx.x×blockDim.x\displaystyle \text{blockIdx}.x \times \text{blockDim}.x
  4. 加法再加上块内编号 threadIdx.x，就得到整个网格中的全局唯一线程号。
    global_tid=blockIdx.x×blockDim.x+threadIdx.x\displaystyle \text{global\_tid} = \text{blockIdx}.x \times \text{blockDim}.x + \text{threadIdx}.x

合并访问优化：仓库搬货要整整齐齐排队

全局内存像一个大仓库，离灶台很远，访问一次代价高昂。当 GPU 上的线程需要读数据时，硬件会按 warp（通常 32 个线程）为一组批量搬运。如果这 32 个线程要的数据在货架上紧挨着，一辆货车就能一趟拉完；如果每个人隔着老远，货车就要反复绕路跑很多趟。

在 CUDA 中，这种“一趟拉完”的模式叫合并访问。最典型的写法就是让第 threadIdx.x 个线程去读数组第 threadIdx.x 个元素。比如处理 float 数组时，线程访问的地址可以写成：

如果写成了 data[threadIdx.x * 8]，相当于每名帮厨都隔 7 个货位去拿货，一趟车只能拉一件，性能会断崖式下跌。解决思路是让数据访问尽量顺序化、连续化，规划好仓库路线，才能让货车一趟接一趟满载而归。

address=base+threadIdx.x×4\displaystyle\text{address} = \text{base} + \text{threadIdx}.x \times 4
通俗说明：float 占 4 字节，所以每往后一个线程，地址就往后跳 4 字节。线程编号连续，地址就连续，硬件就能用少量货车把所有数据一次搬回。
  1. base数组的首地址，相当于仓库货架起点的位置。
    base\displaystyle \text{base}
  2. threadIdx.x线程在块内的编号，同一个 warp 里这 32 个编号基本连续。
    threadIdx.x\displaystyle \text{threadIdx}.x
  3. 乘以 4一个 float 占 4 字节，所以相邻线程访问的地址正好相邻 4 字节。
    threadIdx.x×4\displaystyle \text{threadIdx}.x \times 4
  4. 整体公式当地址随线程编号连续递增时，32 个线程的访存请求会被硬件合并成少量事务，这就是合并访问。
    base+threadIdx.x×4\displaystyle \text{base} + \text{threadIdx}.x \times 4

共享内存分块：小组工作台让协作更快

仓库虽然能通过合并访问高效搬货，但延迟依然很高。于是后厨给每个小组配了一张工作台：组长先派人把今天要用的那批货从仓库搬到台面上，之后组员直接在工作台上取用，不用一次次跑仓库。这个工作台就是 CUDA 里的共享内存。

共享内存用 __shared__ 声明，位于每个线程块内部，块内所有线程都能访问。经典的分块做法是：每个线程块协作把矩阵中的一小块数据搬进共享内存，之后反复使用。假设我们用二维线程块处理矩阵，每个线程搬运一个元素，那么它对应的全局坐标是：

搬完货后，所有线程必须等待同伴把货放好，才能从工作台上取用，这时候还要调用 __syncthreads() 同步屏障。有了工作台，同一小组的成员不必反复跑仓库，后厨的效率自然大大提升。

tile_row=blockIdx.y×TILE+threadIdx.y,tile_col=blockIdx.x×TILE+threadIdx.x\displaystyle\text{tile\_row} = \text{blockIdx}.y \times \text{TILE} + \text{threadIdx}.y, \quad \text{tile\_col} = \text{blockIdx}.x \times \text{TILE} + \text{threadIdx}.x
通俗说明：TILE 是工作台的边长；tile_row 和 tile_col 表示当前小组负责的那块数据在整个矩阵中的起始位置。
  1. blockIdx.y 与 blockIdx.x当前线程块在二维网格中的行、列编号，相当于小组在厨房里的位置。
    blockIdx.y,blockIdx.x\displaystyle \text{blockIdx}.y, \quad \text{blockIdx}.x
  2. TILE分块边长，比如 16×16 的小方块，就是工作台的大小。
    TILE\displaystyle \text{TILE}
  3. 乘法blockIdx.y × TILE 跳过上方所有块覆盖的行，定位本块在矩阵中的起点。
    blockIdx.y×TILE,blockIdx.x×TILE\displaystyle \text{blockIdx}.y \times \text{TILE}, \quad \text{blockIdx}.x \times \text{TILE}
  4. 加法加上线程在块内的编号 threadIdx.y/x，得到该线程要搬运的具体矩阵坐标。
    tile_row=blockIdx.y×TILE+threadIdx.y,tile_col=blockIdx.x×TILE+threadIdx.x\displaystyle \text{tile\_row} = \text{blockIdx}.y \times \text{TILE} + \text{threadIdx}.y, \quad \text{tile\_col} = \text{blockIdx}.x \times \text{TILE} + \text{threadIdx}.x

边界条件防护：货不够分时别让帮厨空手乱抓

真实数据量往往不能被线程块大小整除。比如最后一批货只剩 5 份，可这一组灶台却站着 8 名帮厨。如果没拿到货的人硬冲，轻则拿到隔壁桌的脏盘子，重则直接把货架撞塌。边界检查就是给每名帮厨发一块“有没有货”的提示牌。

在 kernel 里，最常见的做法是先用线程索引换算出行列坐标，再判断坐标是否落在数据范围内。设数据大小为 m 行 n 列，当前线程要处理的行列为 row、col，合法条件可以写成：

更稳妥的做法是启动 kernel 时按 ceil 配置 grid，保证线程数覆盖全部数据，再在 kernel 内加合法性判断；当多个线程可能同时更新同一个位置时，还可以用原子操作兜底。边界问题就像后厨的安全网，多一道检查，就少一次崩溃。

valid(row,col)={1,row<m and col<n0,otherwise\displaystyle\text{valid}(\text{row}, \text{col}) = \begin{cases} 1, & \text{row} < m \ \text{and} \ \text{col} < n \\ 0, & \text{otherwise} \end{cases}
通俗说明：如果 row 和 col 都没越界，线程才允许干活；否则立刻收手，不让任何人访问非法内存。
  1. row 和 col由线程索引换算出来的数据行、列坐标，相当于帮厨拿到的取货单。
    row,col\displaystyle \text{row}, \quad \text{col}
  2. m 和 n矩阵的实际行数和列数，也就是货架真正放得下的货物量。
    m,n\displaystyle m, \quad n
  3. 同时满足row < m 保证行不越界，col < n 保证列不越界，两个条件缺一不可。
    row<m and col<n\displaystyle \text{row} < m \ \text{and} \ \text{col} < n
  4. 输出 1/0条件成立才继续执行后面的内存读写，不成立就跳过，避免访问越界。
    valid(row,col)\displaystyle \text{valid}(\text{row}, \text{col})

Warp 发散规避：32 人小队要齐步走

GPU 实际执行指令时，不是一个个线程单独走，而是把 32 个线程编成一个 warp 集体行动。这就像 32 名帮厨并排站在一条流水线上，听同一个口令齐步走。如果队伍中有人要向左、有人要向右，整条流水线就只能先带一批人走完，再回来带另一批人。

比如代码里写 if (threadIdx.x < 16) doA; else doB;，同一个 warp 内一半人做 A、一半人做 B。硬件无法同时执行两条指令，只能先把做 A 的人调度完，再让做 B 的人执行，白白浪费一半时间。这种现场就叫 warp 发散。我们先看 warp 编号怎么算：

避免 warp 发散的方法不只是“少写 if”，而是把条件从 threadIdx.x 上挪开，或者把相同条件的数据凑到同一个 warp。让 32 人小队始终齐步走，GPU 的流水线才不会因为等人而空转。

warp_id=⌊threadIdx.x32⌋,divergence  ⟺  branch condition varies inside a single warp\displaystyle\text{warp\_id} = \left\lfloor \frac{\text{threadIdx}.x}{32} \right\rfloor, \qquad \text{divergence} \iff \text{branch condition varies inside a single warp}
通俗说明：一个 warp 是 32 个并排站的帮厨；如果分支条件在这 32 个人内部不一致，机器只能分成两批干活，就像查身份证时队伍里一半有证一半没证。
  1. threadIdx.x线程在块内的编号，同一个 warp 的 32 个线程编号连续。
    threadIdx.x\displaystyle \text{threadIdx}.x
  2. 除以 32 向下取整把线程按每 32 个一组分组，得到的商就是 warp 编号。
    ⌊threadIdx.x32⌋\displaystyle \left\lfloor \frac{\text{threadIdx}.x}{32} \right\rfloor
  3. 发散判定只要同一个 warp 内的线程在某个分支上的去留不一致，就会产生发散。
    divergence  ⟺  condition varies within warp\displaystyle \text{divergence} \iff \text{condition varies within warp}
  4. 规避方式让分支条件基于 warp 编号而不是单个线程编号，或者重排数据，让同一 warp 的线程尽量走同一条路。
    condition=f(warp_id)\displaystyle \text{condition} = f(\text{warp\_id})

迭代性能调优：从能出菜到闪电出菜

后厨不能一上来就追求完美，必须先保证菜能端出去，再逐步改进流程。Kernel 优化也一样：先写一个朴素版本，确认结果正确，再用性能分析工具找出瓶颈，然后针对性地优化，最后测量收益。

典型迭代路径是：朴素版本 → 优化全局内存合并访问 → 加共享内存分块 → 消除 warp 发散 → 减少多余同步或使用只读缓存。每一步都要单独验证，避免“优化 A 破坏了优化 B”。性能提升的幅度用加速比来衡量：

最终，你的 GPU 大排档从“勉强开张”变成“翻台率王”。但请记住，优化要靠 profiler 数据说话，不能靠直觉拍脑袋。菜单、工号、搬货、分块、边界、齐步走，每一个环节都打磨到位，kernel 才算真正写成高手水平。

Speedup=TbaselineToptimized\displaystyle\text{Speedup} = \frac{T_{\text{baseline}}}{T_{\text{optimized}}}
通俗说明：原来 100 秒上齐的菜，优化后 20 秒上齐，加速比就是 5 倍。比值大于 1 说明优化有效，小于 1 就说明改坏了。
  1. baseline第一版能正确运行但性能平平的 kernel，也就是后厨最初的原始菜单流程。
    Tbaseline\displaystyle T_{\text{baseline}}
  2. optimized经过合并访问、共享内存分块等改造后的 kernel 耗时。
    Toptimized\displaystyle T_{\text{optimized}}
  3. 比值同一份数据上，旧耗时除以新耗时，得到几倍收益。
    TbaselineToptimized\displaystyle \frac{T_{\text{baseline}}}{T_{\text{optimized}}}
  4. 加速比使用每次改动只验证一个优化点，记录前后耗时，确保收益真实可叠加。
    Speedup=TbaselineToptimized\displaystyle \text{Speedup} = \frac{T_{\text{baseline}}}{T_{\text{optimized}}}

AI 应用场景

深度学习卷积算子优化

阶段：模型部署与推理优化

解决的问题：卷积层如果用朴素循环在 GPU 上跑，大量时间浪费在低效访存上。传统实现要么逐像素遍历，要么频繁访问全局内存，带宽利用率极低。AI 工程师可以用 CUDA kernel 把特征图分块载入共享内存，让线程按 row/col 合并访问，同时避免 warp 发散。

具体案例：某团队优化 Transformer 中的卷积算子，把特征图切分为 16×16 的 tile，每次迭代先在共享内存中缓存邻域数据，再让每个 warp 统一处理同一行数据，最终推理延迟下降 30%。

大规模向量检索与推荐召回

阶段：特征预处理与向量召回

解决的问题：推荐系统每晚要计算千万级向量与用户向量的余弦相似度。传统 CPU 方案耗时严重，而 GPU kernel 需要处理非整数倍的 batch 边界，以及连续内存上的合并访问问题。

具体案例：算法团队用 CUDA 编写向量内积 kernel，启动时按 ceil 计算 block 数，线程内部用 if (tid < n) 做越界保护，并把向量按 cache line 对齐到内存。结果召回耗时从 120ms 降到 15ms，接近一个数量级。

科学计算中的网格模拟

阶段：训练数据生成与数值模拟

解决的问题：流体、电磁、分子动力学等模拟需要对大规模网格做迭代更新。网格边界处的越界风险和多个线程写同一位置的数据竞争，往往让初学者写出崩溃或结果错误的 kernel。

具体案例：某科研团队在流体模拟中用 shared memory 缓存邻域格点，边界处用条件判断跳过无效单元，对跨网格累加结果使用原子操作兜底，最终把模拟速度提升了 20 倍。

Python 演示

python28 行
import numpy as npfrom numba import cuda  @cuda.jitdef vec_add_kernel(a, b, out, n):    tid = cuda.grid(1)    if tid < n:        out[tid] = a[tid] + b[tid]  def main():    n = 1000000    a = np.arange(n, dtype=np.float32)    b = np.full(n, 2.0, dtype=np.float32)    out = np.empty(n, dtype=np.float32)     threads = 256    blocks = (n + threads - 1) // threads    vec_add_kernel[blocks, threads](a, b, out, n)     expected = a + b    assert np.allclose(out, expected)    print("Kernel OK: max diff =", np.abs(out - expected).max())  if __name__ == "__main__":    main()
目的：用 Numba CUDA 实现一个最简单的向量加法 kernel，演示 kernel 声明、执行配置、全局线程号计算和边界保护。
  1. @cuda.jit把 Python 函数编译成 GPU 核函数，相当于给后厨签发菜单。
  2. def vec_add_kernel(a, b, out, n):kernel 声明，参数由主机端传入。
  3. tid = cuda.grid(1)根据 grid/block 计算全局线程号，对应正文的全局索引公式。
  4. if tid < n:边界条件，防止最后一个块越界，就像货不够时提醒帮厨收手。
  5. blocks = (n + threads - 1) // threads向上取整计算线程块数，等价于 gridDim = ceil(n / blockDim)。
  6. vec_add_kernel[blocks, threads](a, b, out, n)主机端启动 kernel，blocks 和 threads 是执行配置。
  7. expected = a + b用 CPU 向量加法验证 GPU 结果正确性。
  8. assert np.allclose(out, expected)校验结果，正确才打印 "Kernel OK"。

常见误区

误区“我启动时的线程数和数据大小是配好的，怎么可能越界？所以 kernel 里不需要 if 判断。”

正解即使 blocks × threads 恰好覆盖数据，在多维索引、共享内存加载、拼接数据中仍然可能越界。更常见的是 blocks = (n + threads - 1) // threads 会多开一些线程，多出来的线程必须用 if (tid < n) 兜底。就像货只剩 5 份却来了 8 名帮厨，多出来的 3 人如果没有提示，就会把隔壁桌的货也拿走。

为何常见：初学者常在小规模且整齐的数据上测试，边界问题被掩盖，换大数据或非整除数据时才暴露。

误区“GPU 核这么多，每个块设置 1024 个线程，并行度不就更强了吗？”

正解blockDim 受硬件上限约束，且块太大可能导致寄存器不足、占用率下降，块内同步等待也更久；块太小又无法隐藏访存延迟。通常从 128、256、512 开始，再结合 occupancy 工具调优。灶台不是越大越好，太大反而会让每个人的操作空间变小、互相等待。

为何常见：把“核多 = 性能高”的直觉直接套到线程块配置上，忽略了资源约束。

误区“共享内存延迟比全局内存低，我在 kernel 一开始就把整个数组复制进 shared memory，速度肯定起飞。”

正解共享内存容量很小，通常只有几十 KB，而且不能跨线程块共享。它只适合“有复用”的数据，比如矩阵分块或卷积邻域。如果没有复用，把数据搬进 shared memory 反而多了一次搬运开销，还要付出同步成本。工作台虽近，但把一箱菜从仓库搬上工作台也要时间，不能让所有人挤着搬一箱菜。

为何常见：把“共享内存快”理解成绝对优势，忽视了载入成本和容量限制。

误区“CPU 上怕分支，GPU 上也一样，所以 kernel 里最好不要写 if。”

正解warp 发散只发生在同一个 warp 内线程走不同分支时。如果条件基于 blockIdx.x，同一个 warp 会整体走同一条路，几乎没有发散；如果条件基于 threadIdx.x，则很容易出现一半 true 一半 false。真正要做的是把条件从 threadIdx.x 上挪开或重排数据，而不是消灭所有 if。

为何常见：把 if 本身当成性能敌人，忽略了 GPU 是以 warp 为调度单位这一核心机制。

面试题

Q1. 在 CUDA 中，哪个限定符声明的函数可由主机端调用并在 GPU 上并行执行？

单选题 · easy

  • A.__host__
  • B.__global__← 正确答案
  • C.__device__
  • D.__shared__

答案：B。__global__ 声明的函数就是 kernel，主机端调用、设备端执行。__device__ 只能在设备端被其他设备函数调用，__host__ 是普通主机函数，__shared__ 是变量限定符，不是函数限定符。

Q2. 数据量 N=1000，blockDim=256，要覆盖所有元素，gridDim 应取多少？

单选题 · easy

  • A.3
  • B.4← 正确答案
  • C.5
  • D.8

答案：B。ceil(1000 / 256) = ceil(3.906) = 4。最后一个线程块只处理 232 个元素，其余 24 个线程需要用边界判断跳过，这正好对应“货不够分时加提示牌”的比喻。

Q3. 线程块内线程编号 threadIdx.x 的起始值是？

单选题 · easy

  • A.0← 正确答案
  • B.1
  • C.32
  • D.256

答案：A。CUDA 的索引体系从 0 开始，块编号 blockIdx.x 和块内线程编号 threadIdx.x 都是从 0 开始的。这与 C 语言数组下标习惯一致，也是写全局线程号公式时必须注意的起点。

Q4. 以下哪种全局内存访问模式最可能实现合并访问？

单选题 · mid

  • A.data[threadIdx.x]← 正确答案
  • B.data[threadIdx.x * 8]
  • C.data[blockIdx.x]
  • D.data[threadIdx.x % 2 ? 1 : 0]

答案：A。合并访问要求同一个 warp 内线程访问连续地址。data[threadIdx.x] 让相邻线程访问相邻地址；threadIdx.x * 8 会造成地址跳跃；data[blockIdx.x] 会让同一块内线程访问同一地址，浪费带宽；取模条件跳转会让一半线程访问同一位置，同样无法合并。

Q5. 一个 warp 通常包含多少个线程？

单选题 · mid

  • A.8 个
  • B.32 个← 正确答案
  • C.64 个
  • D.128 个

答案：B。NVIDIA GPU 的 warp 大小通常固定为 32 个线程。硬件以 warp 为基本调度单位，所以同一个 warp 内线程走不同分支时会发生发散，这也是理解 warp 发散问题的关键前提。

Q6. 下列哪种情况最容易在同一个 warp 内产生分支发散？

单选题 · mid

  • A.if (blockIdx.x % 2 == 0)
  • B.if (threadIdx.x < 16)← 正确答案
  • C.if (blockDim.x > 256)
  • D.if (gridDim.x > 2)

答案：B。发散与否取决于分支条件是否在 warp 内部不一致。threadIdx.x < 16 会让前 16 个线程走 true、后 16 个走 false，必然分裂；而条件基于 blockIdx.x、blockDim.x、gridDim.x 时，同一个 warp 内所有线程取值一样，不会发散。

Q7. 使用共享内存做分块计算时，为什么通常要调用 __syncthreads()？

单选题 · mid

  • A.加快全局内存访问
  • B.保证块内线程读写共享内存完成← 正确答案
  • C.让不同线程块互相通信
  • D.释放共享内存空间

答案：B。共享内存是线程块内部的协作空间，一个线程写入的数据可能被另一个线程读取。__syncthreads() 是块内同步屏障，保证所有线程都到达该点后再继续，防止某个线程提前读取还没写好的数据。它不能跨块通信，也不会释放共享内存。

Q8. 下面哪些是 CUDA kernel 优化的常用手段？

多选题 · hard

  • A.使用全局内存合并访问← 正确答案
  • B.使用共享内存分块复用← 正确答案
  • C.避免 warp 发散← 正确答案
  • D.把所有数据都复制到寄存器

答案：ABC。合并访问、共享内存分块、避免 warp 发散都是 kernel 优化的有效手段。寄存器容量极小，不可能容纳所有数据，强行复制会导致寄存器溢出到局部内存，反而更慢，因此 D 错误。

Q9. 处理矩阵越界问题时，以下哪些做法是正确的？

多选题 · hard

  • A.启动 grid 时按 ceil 向上取整← 正确答案
  • B.在 kernel 内检查 row 和 col← 正确答案
  • C.请求超过物理上限的线程数
  • D.让越界线程直接 return← 正确答案

答案：ABD。按 ceil 配置 grid 能保证线程数覆盖数据，kernel 内检查 row/col 能阻止越界访问，越界线程直接 return 是标准做法。请求超过物理上限的线程数没有意义，也无法真正解决越界，因此 C 错误。

Q10. 在网格跨步循环中，每次迭代的步长应该取多少？

单选题 · hard

  • A.gridDim.x
  • B.blockDim.x
  • C.gridDim.x * blockDim.x← 正确答案
  • D.gridDim.x + blockDim.x

答案：C。网格跨步循环让一个线程处理多个元素，每次跳过的是整个网格的线程总数，也就是 gridDim.x * blockDim.x。这样能减少启动开销并提升数据复用，是处理超大数组时常用技巧。