小区要办百家宴了。每户人家的厨房里都热气腾腾：有的炖肉，有的包饺子，有的蒸鱼。最后居委会要统计“今晚一共做了多少道菜？”。你不能让所有家庭同时冲到登记台前自己拿笔写数字——你一笔我一笔，总会有人的数字被盖掉、写丢；可如果让所有人排好队再慢慢数，又太浪费时间。于是聪明的做法出现了：要么在登记台安排一位接待员，大家排成一队逐个报数，每报一个他立即把总数往上加一笔——这叫原子操作；要么给每家发一张卡片，各家把菜数写在自己的卡片上、投进自己的信箱，等全部投完，再由一位总统计员挨个信箱取卡片加总——这叫两阶段规约。Triton 中的跨 Block 规约，处理的正是这件“各家先算、统一汇总”的事，而原子操作，就是那个绝不丢数的接待员。

岗位关联

理解跨 Block 规约与原子操作，是写出正确、高效 GPU 算子的分水岭；掌握这两种“汇总手段”，是 GPU 并行计算工程师从入门迈向实战的必经之路。

目录
  1. 可视化演示
  2. 一、块内归约：每家厨房先清点自家盘数
  3. 二、跨块协调：把数字送上社区公告板
  4. 三、原子操作：居委会的“单人接待柜台”
  5. 四、两阶段规约：卡片与信箱的收集法
  6. 五、方案取舍：排队报数与收卡汇总的权衡
  7. 六、工程实战：一场百家宴的完整流程
  8. AI 应用场景
  9. Python 演示
  10. 常见误区
  11. 面试题

可视化演示

展示GPU跨Block两阶段规约策略、原子操作并发竞争与双缓冲重叠优化，切换模式掌握Triton并行计算核心技术。

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

两阶段跨Block规约策略 阶段一各Block利用共享内存独立完成局部规约，将大量数据缩减为部分结果 阶段二通过原子操作或二次kernel将各Block部分结果汇总为最终规约值 Phase1: block_local_reduce || Phase2: global_atomic_add 两阶段策略突破共享内存限制实现跨Block规约

一、块内归约：每家厨房先清点自家盘数

办百家宴之前，每家都得先把自己厨房里做的菜数清楚。你总不能等总统计员上门时，才发现锅里还炖着一只鸡没算进去。GPU 里的情况很像：每个线程块（block）负责一段数据，在做跨 Block 汇总之前，必须先把自己这块数据压缩成一个数值，这个动作在并行计算里叫归约（Reduce）。

Triton 中，归约通过 tl.reduce(x, axis, combine_fn) 表达，而 tl.sum 是它最常用的“求和特例”。tl.sum(x, axis) 会沿着指定轴把当前 block 内的元素逐个累加，最终返回一个标量或压缩后的张量。注意，这一步的作用范围仅限于当前 block 内部，结果还“出不了家门”。它保证了后续跨 Block 阶段拿到的输入是一份干净、确定的局部和。

以第 \(b\) 个 block 为例，它负责的元素记为 \(x_{b,0}, x_{b,1}, \dots\)，块内归约就是把这些元素全部加起来：

“自家清点”这一步做得越快、越干净，后面跨块汇总就越省心。块内归约的质量，直接决定了整场百家宴统计的上限。

sb=∑i=0nb−1xb,i\displaystyle s_b = \sum_{i=0}^{n_b-1} x_{b,i}
这个公式说的是：一个线程块把自己负责的 \(n_b\) 个元素从头加到尾，得到局部和 \(s_b\)。它是跨 Block 规约的“原材料”。
  1. 块编号 \(b\)表示当前是第几个线程块，每个块独立计算自己那份数据，互不干扰。
    b∈{0,1,…,B−1}\displaystyle b \in \{0, 1, \dots, B-1\}
  2. 元素下标 \(i\)遍历块内元素的序号，\(n_b\) 是本块负责的元素总数。
    i∈{0,1,…,nb−1}\displaystyle i \in \{0, 1, \dots, n_b-1\}
  3. 局部和 \(s_b\)把当前块的所有元素累加成一个数，这就是该块的块内归约输出。
    sb=xb,0+xb,1+⋯+xb,nb−1\displaystyle s_b = x_{b,0} + x_{b,1} + \cdots + x_{b,n_b-1}

二、跨块协调：把数字送上社区公告板

各家的菜数清点完了，总统计员怎么知道每家数字呢？需要一个“公告板”——大家把数字写到公告板上，统计员才看得到。可比方说，如果所有人都往公告板的同一格子上写字，字迹一定会叠在一起看不清，所以要给每家分配一个独立格子。

Triton 中，block 之间无法直接访问对方的寄存器或 shared memory（块内共享内存），它们唯一能共同读写的地方是全局内存（global memory）——就像小区门口的公告板，谁都能看、谁都能写。因此，跨 Block 规约的第一步，就是让每个 block 把块内归约结果写到全局内存中属于自己的槽位上。槽位索引通常直接用当前 block 的编号 pid（即 program id）计算。

写成代码就是：先声明一个全局缓冲区 buf，形状为 (B,)，然后让第 pid 个 block 执行 buf[pid] = local_sum。因为每个 block 只写自己编号对应的位置，天然不存在写冲突，完全不需要加锁或原子操作保护。

公告板上的格子彼此独立，这正是两阶段规约能够既并行又安全的前提。接下来，就轮到统计员登场了。

buf[b]=sb,b=0,1,…,B−1\displaystyle \mathrm{buf}[b] = s_b, \quad b = 0, 1, \dots, B-1
这个式子表示：每个线程块把自己算好的局部和 \(s_b\) 放进公告板上的第 \(b\) 个格子，各写各的，谁也不会踩到别人。
  1. 缓冲区 \(\mathrm{buf}\)指向全局内存的数组，长度等于 block 总数，相当于公告板上的一排格子。
    buf∈RB\displaystyle \mathrm{buf} \in \mathbb{R}^{B}
  2. 索引 \(b\)用 block id 定位，保证第 \(b\) 个块只写第 \(b\) 个格子，写入位置唯一。
    b=program_id\displaystyle b = \text{program\_id}
  3. 写入时机该赋值语句必须出现在块内归约完成之后，否则写进去的就不是最终局部和。
    buf[b]=sb\displaystyle \mathrm{buf}[b] = s_b

三、原子操作：居委会的“单人接待柜台”

如果不想用公告板和信箱，而是让各家直接到居委会的“总账本”上报数，账本只有一册，接待员必须让所有人排队，一次只处理一个人的“读数字—加数字—写回数字”。这个不可分割的完整动作，就叫原子操作。

Triton 中，最常用的原子操作是 tl.atomic_add(pointer, value)。它把“读出旧值、加上 value、写回新值”三步绑定成一个不可分割的原子事务，硬件保证不会有两个线程交错执行。函数返回的是写入前的旧值；如果业务上不需要旧值，可以直接忽略返回值。它的语义可以写成：

接待员虽然一次只能服务一个人，但绝对不会漏记任何一笔。这就是为什么多个 block 并发写同一个全局计数器时，必须使用原子操作而不是普通加法。

old=A[p];A[p]=old+v\displaystyle \text{old} = A[p]; \quad A[p] = \text{old} + v
这段伪代码的意思是：先取出全局数组 \(A\) 中位置 \(p\) 的当前值，把它加上 \(v\) 再写回去；关键点是“先读、再加、再写”整个过程不会被其他线程打乱。
  1. 读取旧值把指针位置当前存储的值读出来，作为本次原子事务的输入。
    old=A[p]\displaystyle \text{old} = A[p]
  2. 计算新值在旧值基础上加上 \(v\)，得到要写回的新值。
    A[p]new=old+v\displaystyle A[p]_{\text{new}} = \text{old} + v
  3. 不可分割写回硬件保证这三步紧密打包，其他线程不可能看到中间状态，也不会在写回前再次修改同一地址。
    A[p]←A[p]new\displaystyle A[p] \leftarrow A[p]_{\text{new}}

四、两阶段规约：卡片与信箱的收集法

居委会还可以换个思路：给每家发一张卡片，各家把自家菜数写在卡片上、投进自己的信箱；等所有人都投完，再由总统计员取回所有卡片，逐张加出总数。这样既不用排队，也不会发生写冲突。这正是两阶段规约的朴素版本。

在 Triton 中实现时，第一阶段每个 block 把局部和写入全局缓冲区中自己的槽位，因为各写各的，所以不需要原子保护；第二阶段由一个“收尾 block”（或单独启动的第二个 kernel）读回所有槽位，用 tl.sum 一次性汇总。这里只出现一个全局同步点，而不是让每个 block 都去抢同一个全局计数器。

信箱模式的好处是“各行其道、最后统一收口”。每家投信时互不打照面，总统计员只需跑一趟信箱，就能拿到所有数据。

S=∑b=0B−1buf[b]\displaystyle S = \sum_{b=0}^{B-1} \mathrm{buf}[b]
该公式描述了两阶段规约的第二阶段：收尾块把公告板上所有格子里的局部和再相加，最终得到全局总和 \(S\)。
  1. 槽位值 \(\mathrm{buf}[b]\)第 \(b\) 个 block 写入的局部和，来自块内归约。
    buf[b]=sb\displaystyle \mathrm{buf}[b] = s_b
  2. 求和范围\(b\) 从 0 取到 \(B-1\)，覆盖全部 block，一个都不能漏。
    b∈{0,1,…,B−1}\displaystyle b \in \{0, 1, \dots, B-1\}
  3. 全局结果 \(S\)所有局部和的累加，就是跨 Block 规约的最终答案。
    S=buf[0]+buf[1]+⋯+buf[B−1]\displaystyle S = \mathrm{buf}[0] + \mathrm{buf}[1] + \cdots + \mathrm{buf}[B-1]

五、方案取舍：排队报数与收卡汇总的权衡

原子操作像“排队报数”，两阶段规约像“发卡收卡”，各有适合的场景。十家人报数，排队也就几秒钟；可要是一千家人同时涌到柜台，接待员就要忙疯了。但如果发卡收卡，虽然多了一道准备卡片的工序，高峰时反而更从容。

从性能角度粗略分析：原子加法每次都需要对一个全局地址做“读改写”，当 block 数量 \(B\) 很大、并发度很高时，硬件必须串行化对同一个地址的访问，产生大量排队等待，这就是锁竞争带来的性能损失。两阶段规约的写入阶段完全并行，只有最终收尾块需要读一遍所有槽位，整体吞吐更高。但代价是额外占用一块全局显存，并且要多一次收尾调度。

选择哪种方案，本质上就是在“省事的简单”和“高峰期的流畅”之间做权衡。工程上，没有绝对最优，只有场景适配。

Tatomic≈B⋅(trw+tc),Ttwo-phase≈B⋅tw+B⋅tr\displaystyle T_{\text{atomic}} \approx B \cdot (t_{rw} + t_c), \qquad T_{\text{two-phase}} \approx B \cdot t_w + B \cdot t_r
这两项估算对比了两种方案的总耗时：原子方案中每个块都经历一次读改写并承受竞争惩罚 \(t_c\)；两阶段方案中每个块只写一次、收尾块只读一遍，读写时间就是全部开销。
  1. 原子总耗时 \(T_{\text{atomic}}\)\(B\) 个块各自执行一次全局原子读改写，每次还要额外付出排队竞争的代价。
    Tatomic=B(trw+tc)\displaystyle T_{\text{atomic}} = B(t_{rw} + t_c)
  2. 两阶段总耗时 \(T_{\text{two-phase}}\)\(B\) 个块并行写入槽位，总共耗时 \(B t_w\)；收尾块逐个读取 \(B\) 个槽位求和，耗时 \(B t_r\)。
    Ttwo-phase=B tw+B tr\displaystyle T_{\text{two-phase}} = B\,t_w + B\,t_r
  3. 决策点当 \(t_c\) 远大于 \(t_r\)（即竞争很严重）时，两阶段明显更优；当 block 数量很少、\(t_c\) 微不足道时，原子加法又简单又够快。
    若 tc≫tr, 选两阶段；否则可用原子\displaystyle \text{若 } t_c \gg t_r,\ \text{选两阶段；否则可用原子}

六、工程实战：一场百家宴的完整流程

开宴前一天，总统计员必须设计好流程：是让大家排队报数，还是发卡收卡？真实 Triton 内核也是如此。一个完整的跨 Block 归约任务，通常可以分为三步走：先写好内核，确定 grid 大小与 block 数；再在每个 block 内做块内归约，得到局部和；最后通过原子操作或两阶段方式汇总所有局部和。

以计算一个超大向量所有元素之和为例：第一个内核启动 \(B\) 个 block，每个 block 用 tl.sum 算出局部和并写入 buf[pid]；如果走两阶段路线，再启动一个单 block 的小内核读回全部槽位求和；如果走原子路线，则直接让每个 block 执行 tl.atomic_add(global_ptr, local_sum)。这套思路同样适用于 softmax 的分母归约——先求全局最大值保证数值稳定，再求指数和。

有了“块内清点 + 跨块汇总”这套模板，任何需要全局统计的任务——求和、求最大值、softmax 分母——都能像办百家宴一样，先各自精算，再有序汇总。完整的 tl.atomic_add 参数说明可以参考 [Triton Language 官方文档](1)。

pi=e xi−m∑j=1ne xj−m,m=max⁡1≤j≤nxj\displaystyle p_i = \frac{e^{\,x_i - m}}{\displaystyle \sum_{j=1}^{n} e^{\,x_j - m}}, \qquad m = \max_{1 \le j \le n} x_j
这就是带数值稳定的 softmax 公式。先算所有元素的最大值 \(m\)，把每个元素减去 \(m\) 再取指数，最后把所有指数加起来作为分母，既能防止指数爆炸，又能兼容跨 Block 场景。
  1. 最大值 \(m\)先在整个数据范围（可能是跨 Block 的）内求出最大值，用于把指数函数的输入拉到安全区间。
    m=max⁡jxj\displaystyle m = \max_j x_j
  2. 分子 \(e^{x_i - m}\)原始得分减去最大值后取指数，结果被限制在 \((0, 1]\) 内，数值稳定。
    exi−m∈(0,1]\displaystyle e^{x_i - m} \in (0, 1]
  3. 分母（全局归约和）所有分子的累加和，需要跨 Block 归约得到，这正是前五节方法的用武之地。
    ∑j=1nexj−m\displaystyle \sum_{j=1}^{n} e^{x_j - m}

AI 应用场景

大模型 Softmax 跨块归约

阶段：模型推理 / 算子优化

解决的问题：在长序列注意力（如 FlashAttention）中，每行注意力分数的最大值和分母和可能跨越多个线程块。传统做法是一个 block 算完一整行，当序列长度超过单块容量时就会失败。借助跨 Block 规约，先分块求局部最大与局部和，再用原子或两阶段方式合并，从而支持任意长的上下文。

具体案例：某大语言模型在推理 32K 长文本时，注意力得分行被切分到 8 个 block 上计算。每个 block 先算局部最大值，通过原子最大值 tl.atomic_max 拿到全局最大值，再用两阶段归约汇总指数和，最终得到正确的 attention 输出。

分布式训练梯度累加

阶段：模型训练 / 梯度同步

解决的问题：反向传播时，梯度张量可能按数据分片分散在多个 block 中，需要累加成一个完整的梯度。若多个 block 同时写同一个梯度地址，普通加法会丢失更新。原子加法让多 block 安全地把梯度累加到参数梯度缓冲中，且实现极简。

具体案例：在数据并行训练中，一个 batch 被拆到 4 个 GPU 上计算，每个 GPU 内的 kernel 又用多个 block 并行算出局部梯度，最后通过 tl.atomic_add 将所有局部梯度累加到全局梯度桶中，确保梯度值不丢失。

推荐系统特征共现计数

阶段：特征工程 / 数据预处理

解决的问题：CTR（点击率）预估模型中，需要统计用户行为序列里特征对的共现次数。不同 block 处理不同日志片段时，可能命中同一个特征 ID，需要安全地对该 ID 的计数加一。原子加法是解决高频并发计数最直接的手段。

具体案例：某推荐平台统计“用户点击物品 A 后 1 小时内又点击物品 B”的共现矩阵。每条日志由一个 block 处理，命中特征对后执行 tl.atomic_add(matrix_ptr + pair_id, 1)，最终矩阵计数完整无丢失。

Python 演示

python59 行
"""Triton 跨 Block 归约的两种方式：Python 线程模拟教学版"""import threadingimport time  class GlobalCounter:    """用锁模拟硬件原子性：原子操作 = 不可分割的读-改-写"""     def __init__(self):        self.value = 0        self.lock = threading.Lock()     def atomic_add(self, v):        with self.lock:          # 锁保证同一时刻只有一个人操作账本            old = self.value            self.value = old + v        return old  def atomic_reduce(block_sums):    """原子归约：所有 block 并发往同一个全局计数器累加"""    counter = GlobalCounter()    threads = []    for s in block_sums:        t = threading.Thread(target=counter.atomic_add, args=(s,))        threads.append(t)        t.start()                 # 模拟多个 block 同时启动    for t in threads:        t.join()    return counter.value  def two_phase_reduce(block_sums):    """两阶段归约：各写各的槽位，最后统一汇总"""    buffer = [0] * len(block_sums)    for i, s in enumerate(block_sums):        buffer[i] = s             # 写独立槽位，天然无冲突    return sum(buffer)            # 收尾块读回全部槽位求和  def main():    block_sums = [3, 5, 7, 9, 11, 13]    expected = sum(block_sums)     t0 = time.perf_counter()    r1 = atomic_reduce(block_sums)    t1 = time.perf_counter()     t2 = time.perf_counter()    r2 = two_phase_reduce(block_sums)    t3 = time.perf_counter()     print(f"原子归约    结果={r1} 耗时={t1 - t0:.5f}s")    print(f"两阶段归约  结果={r2} 耗时={t3 - t2:.5f}s")    print(f"期望值      {expected}")  if __name__ == "__main__":    main()
目的：用 Python 线程并发的视角，生动模拟 Triton 中“原子加法”与“两阶段规约”两种跨 Block 汇总方式的执行过程与正确性。
  1. with self.lock:锁模拟硬件原子事务，保证同一时刻只有一个线程对全局计数做读改写。
  2. counter.atomic_add(s)每个线程就像每个 block 把自己的局部和累加到全局账本上。
  3. buffer[i] = s两阶段中每个块只写自己的槽位，天然无冲突，不需要任何锁。
  4. return sum(buffer)模拟“收尾块”读回所有槽位完成最终归约。
  5. threading.Thread(target=..., args=(s,))一个线程模拟一个线程块，同时启动模拟 Block 并发执行。

常见误区

误区我以为在 Triton 内核里直接写 `buf[0] += x` 就行，多个 block 同时加到同一个变量上也不会出问题，反正最后结果差不多。

正解普通加法是“读、加、写”三步分开执行的。多个 block 并发操作同一地址时，两个线程可能同时读到旧值、各自加完再分别写回，后写的覆盖先写的，导致更新丢失。必须用 tl.atomic_add 把三步绑成原子事务，或改用两阶段独立槽位写入。

为何常见：日常顺序编程里 += 是天然安全的，很容易把这种直觉带进并发场景，忽略了硬件层面读改写的非原子性。

误区我在 kernel 里用了 `tl.sum`，感觉已经把数据归约完了，为什么还需要再讨论跨 Block？是不是多余一步？

正解tl.sum 只归约当前 block 内的数据，作用域出不了家门外。跨 Block 归约解决的是“多个 block 的局部结果如何汇成一个全局结果”，它依赖全局内存、原子操作或两阶段缓冲区，跟块内 reduce 是两个层面的事。

为何常见：初学并行编程时往往把“归约”理解成一个统一的语义，忽略了 Triton 中块内与块间共享的存储介质完全不同。

误区我听说原子操作有锁竞争，所以任何场景都该用两阶段规约，原子加法能不用就不用。

正解两阶段需要额外的全局缓冲区，还要等待所有 block 写完后再启动收尾逻辑，多了一次 kernel 调度或全局内存读回。在 block 数量少、并发冲突不严重时，原子加法实现更简单、总耗时也更短。正确做法是按 block 规模与写冲突频率做基准测试。

为何常见：技术文章常强调原子操作的争用缺点，读者容易把“有瓶颈”理解成“一定慢”，忽视了问题规模的影响。

误区`tl.atomic_add(ptr, 1)` 返回旧值，我每次写完都直接丢掉，觉得返回值是多余的设计。

正解返回的旧值在很多场景是核心能力：实现自增唯一 ID（拿旧值当 ID）、实现互斥锁（拿旧值判断谁拿到锁）、实现直方图回填（用旧值定位写入位置）都依赖它。忽略返回值只是因为我们暂时不需要，不代表返回值没用。

为何常见：多数入门例子里只关心最终累加结果，很少展示需要旧值的并发算法，导致学习者低估返回值的作用。

误区我以为所有 block 启动以后会自动步调一致，一个 block 写的数据，另一个 block 立刻就能看到，像线程束一样天然同步。

正解Triton 中不同 block 的执行顺序和进度完全独立，块之间没有内置栅栏同步。两阶段规约必须等所有 block 写完槽位，才能安全地启动收尾 block，通常通过 kernel 边界或单独的收尾 kernel 保证可见性。

为何常见：GPU 的 SIMT 模型让初学者误以为所有并行单元是整齐划一的，实际上 block 之间是自由调度的。

面试题

Q1. 在 Triton 中，对当前线程块内的数据求和，应该使用哪个 API？

单选题 · easy

  • A.tl.sum 是块内求和入口← 正确答案
  • B.tl.add 仅做逐元素加法
  • C.tl.reduce 配合求和函数
  • D.tl.atomic_add 只能跨块用

答案：A。tl.sum 是块内归约求和的直接入口；tl.add 只做逐元素相加，不是归约；tl.reduce 是通用归约，但更常见的是直接用 tl.sum；tl.atomic_add 用于跨块原子累加，并非块内求和工具。

Q2. 跨 Block 之间不能直接通过什么进行通信？

单选题 · easy

  • A.全局内存中的数组
  • B.共享内存 shared memory← 正确答案
  • C.一个独立的全局计数器
  • D.全局缓冲区槽位数组

答案：B。shared memory（共享内存）的作用域限于单个 block 内部，其他 block 根本看不到，因此不能用来跨块通信。全局内存、全局计数器、全局缓冲区都是所有 block 可见的公共区域，可以承担跨块通信的职责。

Q3. tl.atomic_add(ptr, v) 的返回值表示什么？

单选题 · mid

  • A.写入前该地址的旧值← 正确答案
  • B.写入后该地址的新值
  • C.本次累加操作的数量
  • D.原子操作的全局序号

答案：A。原子加法返回的是写操作发生之前的旧值，这个旧值可用于生成唯一 ID 或实现锁。新值可以通过旧值加 v 推导出来；返回值并不表示操作数量或序号，因此 B、C、D 均为错误选项。

Q4. 两阶段规约的第一阶段应完成什么操作？

单选题 · easy

  • A.每个块写独立槽位← 正确答案
  • B.一个块读取所有数值
  • C.直接原子加到全局和
  • D.对数据做块内算术

答案：A。两阶段规约的第一阶段是“各写各的槽位”，用全局缓冲区为每个 block 分配独立位置，消除写冲突。第二阶段的“读取所有数值”是收尾块做的事；“直接原子加到全局和”是原子方案而非两阶段方案；块内算术描述太宽泛，不够准确。

Q5. 为什么全局计数器使用普通加法在并发下会出错？

单选题 · mid

  • A.读改写三个步骤不原子← 正确答案
  • B.全局内存的容量太小
  • C.计数器不是张量类型
  • D.编译器优化乱了顺序

答案：A。普通加法在硬件层面是“读、加、写”三步，多个线程交错执行时会出现丢失更新：两个线程读到同一个旧值，各自加一写回，最终只增加一次。全局内存容量、张量类型、编译器排序都不是并发丢失更新的根本原因。

Q6. 两阶段归约相比原子操作的优势包括哪些？

多选题 · hard

  • A.写入阶段无锁竞争← 正确答案
  • B.额外需要全局缓冲
  • C.收尾阶段单块汇总← 正确答案
  • D.完全避免全局同步

答案：AC。A 正确，各 block 写独立槽位，不需要争抢同一地址的锁；C 正确，收尾由一个块统一读回并汇总，结构清晰；B 是两阶段的代价而非优势；D 错误，两阶段仍有同步点（收尾块需要等待全部写完成），只是冲突更少。

Q7. 以下哪个场景更适合用原子加法做跨块规约？

单选题 · mid

  • A.block 数量少且逻辑简单← 正确答案
  • B.block 数量多且并发很高
  • C.需要额外缓冲区避免写冲突
  • D.必须保证写入顺序完全严格

答案：A。原子加法实现简单，适合 block 数量少、冲突不严重的场景。block 数量多且并发高时，全局原子热点会导致明显竞争；需要额外缓冲区避免冲突恰恰是两阶段方案的动机；原子操作本身不保证写入顺序，因此 D 也不适用。

Q8. 跨 Block 归约最常见的两种实现方式是什么？

单选题 · mid

  • A.原子加法和两阶段规约← 正确答案
  • B.共享内存锁和信号量
  • C.块内 reduce 和广播
  • D.异步拷贝和流水线

答案：A。跨 Block 归约的两条主流路线就是原子加法（全局计数器直接累加）和两阶段规约（槽位写入 + 收尾汇总）。共享内存锁无法跨块生效，块内 reduce 与广播解决的是块内问题，异步拷贝属于访存优化手段，都不是跨块汇总的标准方法。

Q9. Softmax 跨 Block 归约中先求全局最大值 \(m\) 的主要目的是什么？

单选题 · hard

  • A.防止指数运算数值溢出← 正确答案
  • B.让每个 block 负载均衡
  • C.减少全局内存访问次数
  • D.加速原子操作执行速度

答案：A。\(e^{x}\) 在 \(x\) 较大时快速膨胀甚至溢出。先求出全局最大值 \(m\)，让每个元素都减去 \(m\) 再取指数，指数结果被控制在 \((0,1]\) 区间，从根本上保证数值稳定。它并不负责负载均衡、访存优化或原子加速，因此 B、C、D 均不正确。

Q10. 关于 Triton 中跨 Block 全局内存可见性，正确的说法有哪些？

多选题 · extreme

  • A.块之间没有自动同步← 正确答案
  • B.需要依赖 kernel 边界保证可见← 正确答案
  • C.一个块写后其他块立即可见
  • D.独立槽位可避免写冲突← 正确答案

答案：ABD。A 正确，block 之间独立调度，无内置栅栏；B 正确，两阶段规约常常借助 kernel 边界让收尾块看到全部写入；C 错误，非原子写在不同 block 间并不保证即时可见；D 正确，每个槽位只被一个块写入，从根本上消除写冲突，这也是两阶段规约的基石。