假设你是一场 GPU 交响乐团的指挥。上千名乐手（线程）齐刷刷坐在谱架前，面前摊着同一份共享总谱。指挥家轻轻举起指挥棒，发出一个“全体停拍”的手势，整个乐团必须同时安静下来，等最后一位翻完谱子的乐手跟上节拍，才能进入下一小节。CUDA 线程同步，就是这门“让上千个乐手对齐节奏”的艺术——没有它，乐手抢拍、串音、乱作一团；滥用它，音乐会不断停顿、毫无效率。从块内屏障 __syncthreads() 的大挥手势，到 warp 级的眼神示意，再到 Cooperative Groups 的“小编制合奏”，每一种同步原语都对应指挥家手中不同的手势。本文带你站在指挥台上，看懂这些手势背后的原理、代价与陷阱。

岗位关联

对 AI 从业者而言，CUDA 内核开发与 GPU 性能优化是 AI 硬件加速的核心技能。掌握线程同步，能让你的并行程序从“偶对偶错”走向“稳如磐石”，也是高级工程师面试与工程实战中的高频考点。

目录
  1. 可视化演示
  2. 同步概览：没有指挥的乐团
  3. 块内同步：全体停拍
  4. Warp 级同步：给四重奏一个眼神
  5. Cooperative Groups：小编制合奏
  6. 同步开销：停拍的成本
  7. 死锁避免：僵住的指挥台
  8. AI 应用场景
  9. Python 演示
  10. 常见误区
  11. 面试题

可视化演示

交互式展示CUDA从Warp到Grid级的多层同步原语、syncthreads屏障机制与同步开销优化策略。

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

CUDA线程同步层次结构 从线程级到设备级五层同步，每层粒度与开销递增 层次：Thread → Warp → Block → Grid → Device 选择合适的同步粒度是GPU性能优化的关键。

同步概览：没有指挥的乐团

想象一支没有指挥的乐团排练：弦乐手和打击乐手共用一个谱架，两人同时伸手改写同一页乐谱——一个刚写下音符，另一个就把整页涂掉了。几轮下来，谱面只剩最后一个音符，整段旋律支离破碎。CUDA 世界里的“共享谱架”就是共享内存或全局内存，而“同时改写”正是竞态条件（Race Condition）。

CUDA 程序动辄启动成千上万个线程并发执行同一个内核函数，它们经常读写同一份数据。当多个线程同时对一个变量做“读-改-写”操作，后写者会覆盖先写者的结果，造成丢失更新。这时就需要同步——让线程们在关键动作前“对齐节奏”，保证数据一致性。线程的层次结构决定了同步的范围：线程是乐手，线程块是整个弦乐组，warp 是四人重奏小组，网格则是整支乐队。

回到乐团：没有指挥的乐队必然乱套，但有了指挥也不是万事大吉——指挥手势不当，同样会拖垮整场演出。这正是接下来要讲的各类同步手势。

xfinal=xinit+1而非xinit+N(N 个线程无同步并发读改写)x_{\text{final}} = x_{\text{init}} + 1 \quad \text{而非} \quad x_{\text{init}} + N \quad (N\ \text{个线程无同步并发读改写})
这个公式说明：N 个线程同时把共享变量加 1，由于互相覆盖，最终只“活下来”一次增量，远小于应有的 N 次。
  1. 最终值 x_final所有线程竞争结束后，共享变量停留在的最终数值。xfinalx_{\text{final}}
  2. 初始值 x_init竞争开始前，共享变量拥有的起点数值。xinitx_{\text{init}}
  3. 丢失更新N 个线程各自的“读-改-写”互相覆盖，最终只保留一次增量。xfinal=xinit+1x_{\text{final}} = x_{\text{init}} + 1
  4. 理论正确值若同步得当，N 次增量应全部生效，结果应为初始值加 N。xinit+Nx_{\text{init}} + N

块内同步：全体停拍

指挥家举起指挥棒，在空中定住——所有乐手必须立刻收弓停声，等最后一位翻完谱子的乐手跟上，才一起进入下一小节。__syncthreads() 就是 CUDA 里这个“全体停拍”的手势：它让线程块内的所有线程在屏障处汇合，先到的原地等待，直到每一个线程都到达后才统一放行。

除了“等人”，它还有第二重身份——内存栅栏。就像指挥棒落下的瞬间，所有乐手必须确认自己已经翻到正确的页码：__syncthreads() 之后的线程，一定能看到屏障之前块内其他线程写入共享内存的数据。这保证了“先写后读”的顺序，防止线程读到半新不旧的脏数据。

“全体停拍”虽然可靠，却有个苛刻前提：块内每个线程都必须走到同一个屏障。如果某个乐手提前离场，这个停拍手势就永远等不到他——这正是下一节和死锁章节要反复强调的陷阱。

trelease=max⁡i ∈ Block tiarrive,Block={0,1,…,blockDim.x−1}t_{\text{release}} = \max_{i\ \in\ \text{Block}}\ t_i^{\text{arrive}}, \qquad \text{Block} = \{0, 1, \ldots, \text{blockDim.x} - 1\}
屏障的放行时刻由最慢的线程决定——先到的线程必须等待最慢的那位，就像乐手们等最后一位翻谱的同事。
  1. 放行时刻 t_release屏障统一放行所有线程的那一刻。treleaset_{\text{release}}
  2. 取最大值 max所有到达时刻中的最大值，即最慢线程到达的时刻，体现“木桶效应”。max⁡i∈Block\max_{i \in \text{Block}}
  3. 到达时刻 t_i^arrive第 i 个线程抵达屏障的时刻，每个线程到达时间各不相同。tiarrivet_i^{\text{arrive}}
  4. 线程块 Block块内线程编号从 0 到 blockDim.x-1，同步只在本块内生效，跨块互不影响。Block={0,1,…,blockDim.x−1}\text{Block} = \{0, 1, \ldots, \text{blockDim.x} - 1\}

Warp 级同步：给四重奏一个眼神

不是每次排练都需要全场停拍。当只有某个弦乐四重奏节奏不稳时，指挥只需递过去一个眼神，四个人对齐即可，其余乐手继续演奏。Warp 级同步就是这样的“眼神”——它的同步范围只有 32 个线程。

Warp 是 GPU 硬件调度的天然单位：32 个线程组成一个 warp，在单一指令流下同时执行（SIMT 模型）。在无分支分歧时，warp 内的线程本来就走同一条指令路径，天然“隐式同步”。但内存访问顺序、显式的数据交换仍需保证可见性，此时可用 __syncwarp(mask) 做轻量级屏障：它只同步 warp 内指定掩码的线程，开销远小于 __syncthreads()。

一个眼神解决一个小范围的对齐问题，比每次全体停拍高效得多。当程序需要更灵活的同步粒度时，就该请出“小编制合奏”的方案了。

laneID=threadIdx.x mod 32,mask=∑l=031bl⋅2l\text{laneID} = \text{threadIdx.x} \bmod 32, \qquad \text{mask} = \sum_{l=0}^{31} b_l \cdot 2^{l}
laneID 是线程在 warp 内的座位号；mask 则像一份“谁参与同步”的点名册，每一位对应一个座位。
  1. laneID 通道号线程在 warp 内的编号，范围是 0 到 31。laneID\text{laneID}
  2. 对 32 取模因为一个 warp 恰好 32 个线程，用线程编号对 32 取余即可得到座位号。 mod 32\bmod 32
  3. 同步掩码 mask一个 32 位的整数，第 l 位为 1 表示编号 l 的通道参与本次同步。mask\text{mask}
  4. 位加权求和把每个参与位 b_l 乘上 2 的 l 次方再累加，正好把 32 个开关打包成一个整数。∑l=031bl⋅2l\sum_{l=0}^{31} b_l \cdot 2^{l}

Cooperative Groups：小编制合奏

大型演出中，指挥常把乐团分成若干个小编制组合分别排练：弦乐四重奏、铜管五重奏各配一位小指挥，互不干扰；需要全场齐奏时再统一合练。Cooperative Groups（协同组）就是 CUDA 提供的这套“小编制”编程模型。

通过 cg::this_thread_block() 可以拿到当前线程块这个“大组”；cg::tiled_partition<32>() 则把块内线程切成若干个固定大小的子组（tile），子组内部可以调用 sync() 做细粒度同步，不会阻塞子组外的线程。更进一步，cg::grid_group 配合 cooperative launch 还能实现网格级同步——相当于让整支乐队在乐章末尾整齐收拍，而不必回到 CPU 端重新发起内核。

小编制的妙处在于“各扫门前雪”：子组内同步不等整块线程，减少了无谓等待。但灵活是有代价的——同步指令本身并不免费，频繁的停拍会让音乐变得拖沓。

sub_id=⌊thread_ranktile_size⌋,tile_size∈{1,2,4,8,16,32}\text{sub\_id} = \left\lfloor \frac{\text{thread\_rank}}{\text{tile\_size}} \right\rfloor, \qquad \text{tile\_size} \in \{1, 2, 4, 8, 16, 32\}
这个公式决定每个线程属于哪个小编制：把线程的线性编号按小组大小向下取整分组。
  1. 子组编号 sub_id线程所属 tile 子组的编号，子组内可独立调用 sync()。sub_id\text{sub\_id}
  2. 线程线性编号 thread_rank线程在块内的连续编号，相当于乐手的工号。thread_rank\text{thread\_rank}
  3. 向下取整除法把连续的 tile_size 个工号分到同一个小组，向下取整保证不留缝隙。⌊⋅⌋\left\lfloor \cdot \right\rfloor
  4. 合法小组大小 tile_size只能是 2 的幂且不超过 32，对应硬件支持的同步宽度。{1,2,4,8,16,32}\{1, 2, 4, 8, 16, 32\}

同步开销：停拍的成本

指挥每喊一次“停”，乐手们要收弓、静默、重新对准节拍；若每两个音符就停一次，音乐会支离破碎，甚至比不排练还慢。同步在 CUDA 里同样不是免费的午餐，它的开销来自三个部分：等待最慢线程造成的空闲、屏障指令本身的执行成本、以及维护内存一致性的栅栏开销。

性能分析中，Amdahl 定律揭示了同步的“天花板”：程序中可并行部分占比越高、同步串行部分占比越低，加速比越理想。每一次 __syncthreads() 都是一段隐形的串行代码，把整组线程捆在一起，再快的线程也不得不向最慢者看齐。因此，工程实践中讲究“粗粒度同步”：在同步点之间尽量多做事，减少同步次数；能用 __syncwarp 或 tile 子组同步就不用块级屏障。

好的指挥懂得少喊停、喊得准。但比“开销”更可怕的是“僵局”——当指挥与乐手互相等待，音乐会永远卡在同一小节。

S=1(1−p)+pNS = \frac{1}{(1 - p) + \frac{p}{N}}
当串行占比 p 固定时，即使线程数 N 无限增加，加速比 S 也逼近 1/p——同步开销就是那个"p"。
  1. 加速比 S并行化后相比串行执行的性能提升倍数。SS
  2. 串行占比 p程序中被同步或串行化的时间比例，p 越大天花板越低。pp
  3. 线程数 N参与并行的线程数量，N 越大并行能力越强。NN
  4. 极限加速比当 N 趋于无穷时，S 收敛于 1/p，串行部分钉死了上限。lim⁡N→∞S=1(1−p)\lim_{N \to \infty} S = \frac{1}{(1-p)}

死锁避免：僵住的指挥台

想象这样的对峙：指挥让弦乐组停下，等待打击乐手翻谱；打击乐组的负责人却命令组员停下，等待弦乐组先动。两边都在等对方先妥协，音乐会永远卡在这一小节。CUDA 的死锁（Deadlock）与此如出一辙——线程们互相等待永远无法满足的条件。

最常见的死锁成因是“分支发散”：把 __syncthreads() 放进 if/else 分支，或让不同线程走不同循环次数路径。例如块内 8 个线程进入 if 分支执行屏障，其余 24 个线程跳过分支直接继续——屏障永远等不到那 24 位“离席”的乐手，于是整个块卡死。同步的本质是“全员到齐”，而全员到齐要求所有线程必须走相同的同步路径、以相同的次数到达每个屏障。

避免死锁的铁律很简单：让每个乐手都知道下一步该在哪一拍汇合——把同步放在所有线程都可到达的统一路径上，避免在 return、break、continue 之后放置依赖它的同步，并用 compute-sanitizer 的 --tool=synccheck 工具自动检测非法的同步布局。

release  ⟺  ⋀i=1kri=true,∃ i: unreachable(ri)⇒deadlock\text{release} \iff \bigwedge_{i=1}^{k} r_i = \text{true}, \qquad \exists\ i:\ \text{unreachable}(r_i) \Rightarrow \text{deadlock}
屏障放行等价于块内 k 个线程的“到达标志”全部为真；只要有一个标志永远无法置真，死锁就必然发生。
  1. 逻辑与 ∧放行条件要求所有到达标志同时为真，缺一不可。⋀i=1kri\bigwedge_{i=1}^{k} r_i
  2. 到达标志 r_i线程 i 是否到达屏障的布尔量，到达则为 true。ri=truer_i = \text{true}
  3. 存在不可达 ∃只要存在某个线程因分支发散或提前返回而永远到不了屏障，条件恒为假。∃ i: unreachable(ri)\exists\ i:\ \text{unreachable}(r_i)
  4. 死锁结论条件永远为假时，所有等待中的线程互相僵持，程序挂起。deadlock\text{deadlock}

AI 应用场景

深度学习梯度归约内核

阶段：模型训练与推理加速

解决的问题：训练大模型时，每个线程块并行算出部分梯度，若直接写回全局内存，不同块会互相覆盖导致梯度丢失。传统做法是多次启动内核做规约，通信开销巨大；AI 解决方案是在块内用 __syncthreads() 先同步局部归约结果，再统一写回，减少内核启动次数。

具体案例：在 Megatron 风格的大规模语言模型训练中，梯度 All-Reduce 内核在每个线程块内先做树形归约，屏障同步后再把块级部分和写进全局缓冲区，训练吞吐量提升 15% 以上。

大规模科学仿真的网格级同步

阶段：HPC 并行计算核心循环

解决的问题：分子动力学、流体仿真每推进一个时间步，都要求全场数据交换完毕才能进入下一步。传统方案每步启动一个新内核，启动开销巨大；用 Cooperative Groups 的 grid_group 在网格内直接同步，实现“一步一同步”的流水线。

具体案例：某 GPU 加速的 N-body 模拟程序，用 cooperative launch 让 512 个线程块在每步结束处网格同步，省去了每步的内核启动与设备同步，整体仿真速度提升约 30%。

FlashAttention 式算子中的 tile 同步

阶段：Transformer 注意力内核实现

解决的问题：注意力计算要把 QK 矩阵分块（tile）后流水处理，块间数据依赖需要轻量级同步。若一律用 __syncthreads() 会拖慢 warp 间流水；用 tiled_partition 把线程块分成多个 warp 级子组，各自同步、并行推进。

具体案例：实现 FlashAttention-2 风格的前向内核时，每个 tile 内部用 __syncwarp 保证共享内存读写顺序，配合 tile 间的粗粒度屏障，使注意力算子在 A100 上达到理论峰值带宽的 85%。

Python 演示

python82 行
import sysimport threading sys.setswitchinterval(1e-5)  # 调小线程切换间隔, 放大竞态  class SimpleBarrier:    """模拟 CUDA __syncthreads() 的可复用屏障"""     def __init__(self, num_threads):        self.num_threads = num_threads        self._count = 0        self._generation = 0        self._cond = threading.Condition()     def sync(self):        with self._cond:            gen = self._generation            self._count += 1            if self._count == self.num_threads:                self._count = 0                self._generation += 1                self._cond.notify_all()            else:                while gen == self._generation:                    self._cond.wait()  N = 8ROUNDS = 500counter = 0results = [0] * N  def race_worker():    """无同步: 读-改-写 竞态导致丢失更新"""    global counter    for _ in range(ROUNDS):        tmp = counter        for _ in range(200):   # 模拟计算, 放大切换概率            pass        counter = tmp + 1  def barrier_worker(bar, tid):    """有同步: 写私有结果后越过屏障, 保证对全员可见"""    global results    s = 0    for _ in range(ROUNDS):        s += 1    results[tid] = s    bar.sync()                 # 屏障之后, 所有私有结果都已落定  def run_race():    global counter    counter = 0    threads = [threading.Thread(target=race_worker) for _ in range(N)]    for t in threads:        t.start()    for t in threads:        t.join()    return counter  def run_barrier():    bar = SimpleBarrier(N)    threads = [threading.Thread(target=barrier_worker, args=(bar, i))               for i in range(N)]    for t in threads:        t.start()    for t in threads:        t.join()    return sum(results)  if __name__ == "__main__":    expected = N * ROUNDS    r1 = run_race()    r2 = run_barrier()    print("无同步(竞态):", r1, "期望:", expected, "丢失:", expected - r1)    print("屏障同步:", r2, "期望:", expected, "正确:", r2 == expected)
目的：用 8 个 Python 线程模拟 8 个 CUDA 线程，对比“无同步读改写”与“屏障同步后归约”的结果差异，直观展示竞态条件与 __syncthreads() 的修复作用。
  1. self._count += 1每来一个线程就登记一次，记录已到达屏障的线程数。
  2. self._generation += 1最后一个线程到达时把代数加一，通知所有等待者进入下一轮。
  3. while gen == self._generation: self._cond.wait()不是最后一个到达就挂起等待，直到代数变化被唤醒。
  4. tmp = counter无同步线程先读出旧值，这是竞态的起点。
  5. counter = tmp + 1写回时可能覆盖别人刚写的新值，造成丢失更新。
  6. results[tid] = s各自把私有结果写进共享列表，写完之后才能越过屏障。
  7. bar.sync()屏障放行后，任何线程都能读到完整的 results，体现内存可见性保证。

常见误区

误区我以为 `__syncthreads()` 就是让快的线程等一下慢的线程，等齐了继续跑，跟内存没什么关系。

正解它同时是内存栅栏。就像指挥举棒的瞬间，乐手不仅要停下，还必须确认谱子已经翻到正确页码——屏障之前的共享内存写入，在屏障之后对所有块内线程可见。若只“等人”不“刷内存”，线程可能读到屏障前写入的半旧数据。

为何常见：很多教材把同步画成“排队等待图”，只强调汇合语义，忽略了内存可见性这一半职责。

误区我觉得把 `__syncthreads()` 放在 if 的条件分支里没问题，反正编译器会帮我处理不同分支的线程。

正解这是死锁的头号来源。屏障要求块内所有线程都到达同一个屏障；若 8 个线程进 if、24 个线程走 else，走进 if 的线程在屏障前干等，永远等不到那 24 位，整个线程块挂死。同步必须放在所有线程都能执行的统一路径上。

为何常见：写串行代码时“条件等待”很自然，但 CUDA 屏障是“全员到齐”语义，容易用单线程思维误判。

误区既然 warp 里 32 个线程天然同时执行同一条指令，那 warp 内部肯定不存在同步问题，不用管。

正解无分歧时指令流确实同步，但内存访问顺序和显式数据交换仍需保证。比如一个线程写共享内存、另一个线程立刻读，可能需要 __syncwarp() 做轻量级栅栏；分支分歧时 warp 会串行执行各路径，隐式同步也不成立。

为何常见：把“指令级同步”误当成“内存级同步”，忽略了 SIMT 模型下分支与访存顺序的复杂性。

误区为了保险，我在内核里到处加 `__syncthreads()`，觉得多同步几次总不会出错，性能差点无所谓。

正解每次同步都要付出三重代价：等待最慢线程的空闲、屏障指令开销、内存栅栏的一致性维护成本。过度同步会让并行程序退化成串行，甚至比不同步还慢。正确做法是减少同步次数、用 warp/tile 级细粒度同步替代块级屏障。

为何常见：初学阶段“调通”优先于“调快”，容易把同步当成万能安全网，忽视了 Amdahl 定律对串行部分的惩罚。

面试题

Q1. __syncthreads() 的主要作用是什么？

单选题 · easy

  • A.让线程在屏障处汇合并保证内存可见← 正确答案
  • B.只同步线程块内最快的若干线程
  • C.让全部线程的共享内存自动清零
  • D.让不同线程块之间的线程通信

答案：A。__syncthreads() 是块内屏障，要求块内所有线程到齐才放行，同时充当内存栅栏，保证屏障前的写入对屏障后的读取可见。它只同步本块线程，不能跨块通信，也不会清零内存。

Q2. 当一个线程块的所有线程都调用 __syncthreads() 时会发生什么？

单选题 · easy

  • A.块内全部线程到达后一起放行← 正确答案
  • B.只要一个线程到达就继续执行
  • C.线程块会自动分裂成多个小块
  • D.所有线程立即退出当前内核

答案：A。屏障的语义是“全员到齐才放行”，先到的线程原地等待最慢的线程。这就像指挥举棒后，乐手们等最后一位翻完谱的同事。它不会分裂线程块，更不会让线程退出。

Q3. 在条件分支中只有部分线程调用 __syncthreads() 会导致什么结果？

单选题 · mid

  • A.死锁或出现未定义的行为← 正确答案
  • B.获得明显的性能收益
  • C.编译器自动修复该问题
  • D.只同步到达的线程即可

答案：A。屏障要求块内所有线程都到达。若部分线程跳过屏障，屏障永远等不到它们，线程块卡死，属于未定义行为。编译器无法自动修复，性能更无从谈起。这也是死锁章节强调的头号陷阱。

Q4. 一个 warp 通常由多少个线程组成？

单选题 · easy

  • A.32 个线程← 正确答案
  • B.64 个线程
  • C.128 个线程
  • D.取决于线程块大小

答案：A。GPU 以 32 个线程为一个 warp 作为硬件调度单位，它们在同一指令流下执行。线程块大小是程序员指定的逻辑概念，与 warp 的固定 32 线程硬件尺寸无关。

Q5. tiled_partition<32>() 会把一个线程块划分成怎样的结构？

单选题 · mid

  • A.每 32 个线程分成一个子组← 正确答案
  • B.每 1 个线程构成一个子组
  • C.每 64 个线程分成一个子组
  • D.整个线程块变成一个子组

答案：A。tiled_partition<T>() 的模板参数 T 是子组大小，T=32 表示每 32 个线程一组，子组内部可独立调用 sync() 同步。T 只能取 1、2、4、8、16、32 这些 2 的幂，且不超过 warp 大小。

Q6. 关于 __syncwarp() 的说法正确的有哪些？

多选题 · mid

  • A.只在单个 warp 内部同步← 正确答案
  • B.开销比 __syncthreads 更小← 正确答案
  • C.支持跨线程块同步数据
  • D.通过掩码选择参与同步的通道← 正确答案

答案：ABD。__syncwarp() 是 warp 级轻量屏障，只同步 32 个线程组成的 warp，开销小于块级屏障；它用 32 位掩码指定哪些通道参与同步。它不能跨线程块，块间同步需要 Cooperative Groups 的网格级机制。

Q7. 某线程块内 8 个线程进入 if 分支调用 __syncthreads()，其余 24 个线程走 else 分支直接继续，会发生什么？

单选题 · hard

  • A.未定义行为或死锁卡死← 正确答案
  • B.自动降级为顺序执行
  • C.只同步到达的线程即可
  • D.编译器警告后继续执行

答案：A。屏障要求块内所有线程到达同一位置。24 个线程绕过屏障后可能继续执行、甚至退出内核，if 分支里的 8 个线程永远等不齐，导致死锁或未定义行为。解决方法是把同步移到 if/else 合并后的统一路径上。

Q8. 同步开销的主要来源包括哪些？

多选题 · hard

  • A.等待最慢线程产生的停顿← 正确答案
  • B.屏障指令本身的执行开销← 正确答案
  • C.维护内存一致性的额外成本← 正确答案
  • D.显存容量不足导致的等待

答案：ABC。同步开销有三块：最慢线程拖慢全体（木桶效应）、屏障指令本身的流水线代价、内存栅栏保证一致性的开销。显存容量属于存储瓶颈，与同步指令的串行化开销无关。

Q9. 下列哪项是避免同步死锁的最佳实践？

单选题 · easy

  • A.保证全部线程走同一同步路径← 正确答案
  • B.尽量密集地调用同步原语
  • C.只让主线程执行同步等待
  • D.用原子操作替代所有同步

答案：A。死锁的根源是部分线程到不了屏障，因此铁律是让所有线程以相同路径、相同次数到达每个屏障。密集同步会放大开销，只同步主线程违背全员语义，原子操作也无法替代需要协作的屏障。

Q10. 下列关于 CUDA 线程同步的说法正确的有哪些？

多选题 · extreme

  • A.块内线程可用 __syncthreads 同步← 正确答案
  • B.网格级同步需 Cooperative Groups← 正确答案
  • C.无分歧时 warp 线程天然同步← 正确答案
  • D.同步可以消除一切数据竞争

答案：ABC。__syncthreads() 是块内标准同步原语；网格级同步需要 cooperative launch 配合 grid_group；同一 warp 无分歧时指令天然同步。但同步只保证“到齐+可见”，不保证读改写原子性——要消除特定数据竞争还需原子操作或更细粒度设计。