PTX ISA 9.4：8. 内存一致性模型

在多线程执行中，每个线程所执行的内存操作的副作用，会以一种部分且互不相同的顺序对其他线程可见。 这意味着任意两个操作可能看起来没有先后顺序，或者对不同线程呈现出不同的顺序。内存一致性模型 所引入的公理，精确规定了不同线程所观察到的各个顺序之间哪些矛盾是不允许出现的。

在任何约束都不存在的情况下，每个读取操作返回由对同一内存位置执行的某次写入操作所提交的值， 包括对该内存位置的初始写入。内存一致性模型实际上限制了这样一个候选写入的集合，读取操作可以从 该集合中返回一个值。

8.1. 模型的范围与适用性#

本模型下所规定的约束适用于任意 PTX ISA 版本号的 PTX 程序，这些程序运行在 sm_70 或更高的架构上。

内存一致性模型不适用于纹理（包括 ld.global.nc）和 表面访问。

8.1.1. 系统作用域下原子性的限制#

当与主机 CPU 通信时，某些具有系统作用域的强操作在有些系统上可能无法以原子方式执行。有关对主机 内存的原子性保证的更多细节，请参阅CUDA 原子性要求。

8.2. 内存操作#

PTX 内存模型中的基本存储单位是字节，由 8 位组成。PTX 程序可用的每个状态空间都是内存中一段 连续字节的序列。PTX 状态空间中的每个字节，相对于所有能够访问同一状态空间的线程而言，都具有一个 唯一的地址。

每条 PTX 内存指令都指定一个地址操作数和一个数据类型。地址操作数包含一个虚拟地址，该虚拟地址在 内存访问期间被转换成物理地址。物理地址与数据类型的大小一起定义了一个物理内存位置，它是以物理地址 为起点、按字节延伸达该数据类型大小的一段字节范围。类似地，地址操作数与数据类型的大小一起定义了该 内存操作所访问的虚拟内存地址范围。

内存一致性模型的规范使用“地址”或“内存地址”这些术语来指代一个虚拟地址，并使用术语“内存位置”来 指代一个物理内存位置。

每条 PTX 内存指令还指定了对应的内存位置中所有字节上要执行的操作——即读取、写入或原子读改写中的一 种。

8.2.1. 重叠#

当某一位置的起始地址处于构成另一位置的那些字节的范围之内时，称这两个内存位置重叠。当两个操作所 访问的虚拟地址范围相交时，称这两个内存操作重叠。当两个内存位置完全相同时，称该重叠为完全重叠；否则 称为部分重叠。

8.2.2. 别名#

当两个不同的虚拟地址映射到同一个内存位置时，称它们是别名。

8.2.3. Multimem 地址#

multimem 地址是一种虚拟地址，它指向跨设备的多个不同内存位置。

只有 multimem.* 操作对 multimem 地址是有效的。也就是说，在任何其他内存操作中访问一个 multimem 地址，其行为是未定义的。

8.2.4. 向量数据类型的内存操作#

内存一致性模型将执行于带有标量数据类型的内存位置上的操作联系起来，标量数据类型的最大大小和 对齐均为 64 位。带有向量数据类型的内存操作被建模为一组等价的、带有标量数据类型的内存操作，这些 操作以未指定的顺序在向量中的各个元素上执行。

8.2.5. 打包数据类型的内存操作#

一个打包数据类型由同一标量数据类型的两个值组成，如打包数据类型所述。这些值在相邻的 内存位置中被访问。针对一个打包数据类型的内存操作被建模为一对针对该标量数据类型的等价内存操作，这些 操作以未指定的顺序在打包数据的每个元素上执行。

8.2.6. 初始化#

内存中的每个字节都由一次假想的写入 W0 初始化，该写入在程序中启动任何线程之前执行。如果该 字节包含在某个程序变量中，且该变量具有一个初始值，那么 W0 为该字节写入相应的初始值；否则假定 W0 已向该字节写入了一个未知但恒定的值。

8.3. 状态空间#

内存一致性模型中定义的关系独立于状态空间。特别地，因果序在跨所有状态空间的所有内存操作之上 闭合。但某一状态空间中一次内存操作的副作用，只能被同样能够访问该状态空间的操作直接观察到。这在 作用域之外进一步限制了一次内存操作的同步效果。例如，PTX 指令 ld.relaxed.shared.sys 的同步效果与 ld.relaxed.shared.cluster 的同步效果完全相同，因为在同一 cluster 之外的任何线程都无法执行一个访问相同内存位置的操作。

8.4. 操作类型#

为简便起见，本文档的其余部分引用以下操作类型，而不是提及产生这些类型的具体指令。

8.4.1. mmio 操作#

一次mmio 操作是指定带有 .mmio 限定符的内存操作。它通常在一个映射到对等 I/O 设备控制寄存器的内存位置上进行。它也可用于线程之间的通信，但相对于非mmio 操作，其性能较差。

mmio 操作的语义含义无法被精确定义，因为它是由底层 I/O 设备定义的。若从内存一致性模型的角度 对 mmio 操作的语义进行形式化规定，它等价于一次强操作的语义。但如果它在所指定的作用域处满足 CUDA 原子性要求，那么它会遵循若干实现相关的属性：

• 在所指定的作用域内，写入总是被执行，且从不被合并。

• 在所指定的作用域内，读取总是被执行，且不会被转发、预取、合并，也不允许命中任何 缓存。

  • 作为例外，在某些实现中，邻近的位置也可能被一并加载。在这种情况下，所加载的数据量 取决于具体实现，其大小在 32 到 128 字节之间变化。

8.4.2. volatile 操作#

一次volatile 操作是指定带有 .volatile 限定符的内存操作。 volatile 操作的语义等价于一次具有系统作用域的 relaxed 内存操作，但带有以下额外的、实现相关的约束：

• 程序所执行的 volatile 指令（而非操作）的数量被保留。硬件可能会将多个不同 volatile 指令所发出的 volatile 操作合并与归并，也就是说，程序中 volatile 操作 的数量不被保留。

• volatile 指令不会相对于其他 volatile 指令被重新排序，但这些指令所执行的内存 操作可能相对于彼此被重新排序。

8.5. 作用域#

每个强操作都必须指定一个作用域，作用域是可以直接与该操作交互并建立内存一致性模型中所述 任何关系的那些线程的集合。共有四个作用域：

请注意，warp 不是一个作用域；在内存一致性模型中，CTA 是能够构成一个作用域的最小线程集合。

8.6. 代理#

一个内存代理，或一个代理，是应用于一种内存访问方式的抽象标记。当两个内存操作使用不同的 内存访问方式时，称它们属于不同的代理。

需要在不同代理之间同步内存操作时，需要一个代理栅栏。尽管虚拟别名使用的是 generic 内存 访问方式，但由于使用不同的虚拟地址表现得如同使用了不同的代理，因此它们需要一个代理栅栏来建立内存 顺序。

除非另有说明，如操作类型中定义的内存操作使用 generic 内存访问方式，即一个generic proxy。使用不同于 generic 方式的访问方式的操作包括：

• 纹理和表面访问，

• 通过同一代理、使用不同虚拟内存地址对同一位置的访问，

• 由 .async.bulk 操作进行的 async-proxy 和 tensormap-proxy 访问，

• 由 fabric 操作进行的 fabric-proxy 访问，

• 由 ld.proxy::readonly 操作进行的 readonly-proxy 访问。

readonly proxy 可用于加载只读数据。有关细节，请参阅 数据移动与转换指令：ld。

8.6.1. 强代理访问#

通过以下代理进行的强修改，

• generic-proxy，

• async-proxy，

• tensormap-proxy，

• fabric-proxy，

最终可被该列表中通过另一个代理对同一位置所执行的强访问观察到，前提是修改与访问是：

• 系统作用域——即使它们使用了不同的内存地址——或者

• 当所涉及的不同代理属于下列列表时，为 GPU 作用域：

  • async-proxy，

  • generic-proxy。

8.7. 道德强操作#

当两个操作满足以下所有条件时，称它们相对于彼此是道德强的：

1. 这些操作在程序顺序中相关（即，它们都由同一线程执行），或者每个操作都是强操作，并指定一个 包含执行另一操作的那个线程的作用域。

2. 两个操作都通过同一个代理执行。

3. 如果两者都是内存操作，则它们完全重叠。

内存一致性模型中的大多数（但并非全部）公理都依赖于道德强操作之间的关系。

8.7.1. 冲突与数据竞争#

当两次重叠的内存操作中至少有一次是写入时，称它们冲突。

当两次冲突的内存操作在因果序中不相关，并且它们彼此并非道德强时，称它们处于一次数据竞争 之中。

8.7.2. 混合大小数据竞争的限制#

两次完全重叠的操作之间的一次数据竞争称为统一大小数据竞争，而两次部分重叠的操作之间的 一次数据竞争称为混合大小数据竞争。

如果一个 PTX 程序包含一个或多个混合大小数据竞争，则内存一致性模型中的公理不适用。但对于只存在 统一大小数据竞争 的 PTX 程序，这些公理足以描述其行为。

混合大小 RMW 操作的原子性

在任何程序（无论是否包含混合大小数据竞争）中，对于每一对重叠的原子操作 A1 和 A2，只要每个操作 都指定一个包含另一个操作的作用域，以下性质都成立：由 A1 指定的读改写操作要么在 A2 被发起之前完全 执行完毕，要么反之。无论这两个操作 A1 和 A2 是部分重叠还是完全重叠，该性质都成立。

8.8. 释放与获取模式#

某些指令序列会形成参与后文所述内存同步的模式。释放模式使当前线程中先前的操作¹对其他线程的某些操作可见。获取模式使其他线程的某些操作对当前线程中之后的操作可见。

对位置 M 的释放模式由以下之一构成：

1. 对 M 的释放操作

   例如：st.release [M]; 或 atom.release [M]; 或 mbarrier.arrive.release [M];

2. 或者对 M 的释放或获取-释放操作，后接按程序顺序对 M 的强写入

   例如：st.release [M]；st.relaxed [M];

3. 或者释放或获取-释放内存栅栏，后接按程序顺序对 M 的强写入

   例如：fence.release; st.relaxed [M]; 或 fence.release; atom.relaxed [M];

4. 或者释放或获取-释放内存栅栏，后接按程序顺序在 M 上执行强写入的异步操作

   例如：fence.release; cp.async.bulk.global.shared.relaxed.sys.b128 [M];

5. 或者在 M 上执行强写入的释放异步操作

   例如：st.async.release.sys [M];

由释放模式建立的任何内存同步，只影响该模式中首条指令之前、按程序顺序发生的操作。

对位置 M 的获取模式由以下之一构成：

1. 对 M 的获取操作

   例如：ld.acquire [M]; 或 atom.acquire [M]; 或 mbarrier.test_wait.acquire [M];

2. 或者对 M 的强读取，后接按程序顺序对 M 的获取操作

   例如：ld.relaxed [M]; ld.acquire [M];

3. 或者对 M 的强读取，后接按程序顺序的获取内存栅栏

   例如：ld.relaxed [M]; fence.acquire; 或 atom.relaxed [M]; fence.acquire;

4. 或者先观察到在 M 上执行强读取的异步操作完成，随后这一观察之后按程序顺序接一个获取内存栅栏

   例如： cp.async.bulk.mbarrier::complete_tx::bytes.relaxed.sys.b128 [dst], [M], size, [barrier]; // 在 M 上强读取 mbarrier.try_wait.relaxed p, [barrier]; // 观察到在 M 上执行强读取的异步操作完成 @p fence.acquire; // 在观察到完成后按程序顺序的获取栅栏

由获取模式建立的任何内存同步，只影响该模式中末条指令之后、按程序顺序发生的操作。

请注意，尽管原子归约在其读-改-写序列中概念上执行一次强读取，但该强读取并不构成获取模式。

例如：red.add [M], 1; fence.acquire; 不构成获取模式。

¹ 对于释放和获取两种模式，这一效果都会通过因果序的传递性进一步扩展到其它线程中的操作。

8.9. 内存操作的顺序#

每个线程所执行的操作序列被表示为程序顺序，而跨线程的内存同步被表示为因果序。内存操作的副作用对其他内存操作的可见性被表示为通信序。内存一致性模型规定了一侧的通信序与另一侧的因果序和程序顺序之间所不允许的矛盾。

8.9.1. 程序顺序#

程序顺序将某个线程执行的所有操作，与顺序处理器执行相应 PTX 源代码中指令的先后次序关联起来。它是一个传递关系，在该线程所执行的操作上构成全序，但不关联来自不同线程的操作。

8.9.1.1. 异步操作#

某些 PTX 指令（cp.async、cp.async.bulk、cp.reduce.async.bulk、wgmma.mma_async 的所有变体）所执行的操作相对于执行该指令的线程是异步的。这些异步操作排在同一线程中先前指令之后（wgmma.mma_async 的情况除外），但它们不属于该线程程序顺序的一部分。取而代之，它们提供指令说明中所记载的较弱排序保证。

例如，作为 cp.async 的一部分所执行的加载与存储彼此有序，但与同一线程发起的任何其他 cp.async 指令的加载与存储并无次序，也与该线程随后发出的任何其他指令（cp.async.commit_group 或 cp.async.mbarrier.arrive 除外）并无次序。由 cp.async.mbarrier.arrive 指令执行的异步 mbarrier arrive-on 操作，与同一线程发起的所有先前 cp.async 操作所执行的内存操作有序，但与线程所发出的任何其他指令的内存操作并无次序。作为 cp.async.bulk 和 cp.reduce.async.bulk 指令所有变体的一部分的隐式 mbarrier complete-tx 操作，仅与同一异步指令所执行的内存操作有序，特别地，它不会传递性地建立相对于发出线程先前指令的次序。

8.9.2. 观察序#

观察序通过一串可选的原子读-改-写操作，将写入 W 与读取 R 关联起来。

若满足以下条件，写入 W 在观察序中先于读取 R：

1. R 与 W 是道德强的，且 R 读取到 W 所写的值，或者

2. 对于某个原子操作 Z，W 在观察序中先于 Z，且 Z 先于 R。

8.9.3. fence.sc 序#

Fence-SC 序是一个在运行时确定的无环偏序，它关联每一对道德强 fence.sc 操作。

8.9.4. 内存同步#

由不同线程执行的同步操作在运行时按此处所述彼此同步。这种同步的效果是建立跨线程的因果序。

1. 若 X 在Fence-SC 序中先于 Y，则fence.sc 操作 X 与fence.sc 操作 Y 同步。

2. bar{.cta}.sync 或 bar{.cta}.red 或 bar{.cta}.arrive 操作同步于在同一栅栏上执行的 bar{.cta}.sync 或 bar{.cta}.red 操作。

3. barrier.cluster.arrive 操作与 barrier.cluster.wait 操作同步。

4. 若 X 中的写入操作在观察序中先于 Y 中的读取操作，且 X 中的首个操作与 Y 中的末个操作是道德强的，则释放模式 X 与获取模式 Y 同步。

API 同步

某些 CUDA API 也能建立同步关系。

1. 入队到某个 CUDA stream 中的任务的完成，与同一 stream 中后继任务（如果存在）的启动同步。

2. 就上述而言，在某个 stream 中记录或等待一个 CUDA event，或者由于 cudaStreamLegacy 而导致插入一个跨 stream 栅栏，都会把任务入队到相关的 stream 中，即使没有直接的副作用。event 记录任务与匹配的event 等待任务同步，栅栏到达任务与匹配的栅栏等待任务同步。

3. CUDA kernel 的启动与 kernel 中所有线程的启动同步。kernel 中所有线程的结束与 kernel 的结束同步。

4. CUDA graph 的启动与图中所有源节点的启动同步。CUDA graph 中所有汇节点的完成与图的完成同步。图节点的完成与所有具有直接依赖关系的节点的启动同步。

5. 为将任务入队而发起的 CUDA API 调用的启动，与该任务的启动同步。

6. 某个 stream 中最后入队任务（如果存在）的完成，与从 cudaStreamSynchronize 返回同步。最近入队的匹配event 记录任务（如果存在）的完成，与从 cudaEventSynchronize 返回同步。同步某个 CUDA 设备或上下文，其行为如同同步该上下文中的所有 stream，包括那些已被销毁的 stream。

7. 从查询某个 CUDA 句柄（例如 stream 或 event）的 API 返回 cudaSuccess，其行为与从匹配的同步 API 返回相同。

除了建立同步关系之外，上述 CUDA API 同步机制还参与代理保持的基础因果序；但tensormap 代理除外——它在 CUDA kernel 启动时并未从通用代理获取，因此需要时必须使用 fence.proxy.tensormap::generic.acquire 显式获取。

8.9.5. 因果序#

因果序刻画内存操作如何通过同步操作跨线程变得可见。公理“因果性”利用这一顺序来约束读取操作可以从中读取值的写入操作集合。

因果序中的关系主要由基础因果序¹ 中的关系构成，它是一种在运行时确定的传递序。

基础因果序

若满足以下条件，操作 X 在基础因果序中先于操作 Y：

1. X 在程序顺序中先于 Y，或者

2. X 与 Y 同步，或者

3. 对于某个操作 Z，

   1. X 在程序顺序中先于 Z，且 Z 在基础因果序中先于 Y，或者

   2. X 在基础因果序中先于 Z，且 Z 在程序顺序中先于 Y，或者

   3. X 在基础因果序中先于 Z，且 Z 在基础因果序中先于 Y。

代理保持的基础因果序

若内存操作 X 在基础因果序中先于 Y，且满足以下条件，则 X 在代理保持的基础因果序中先于内存操作 Y：

1. X 与 Y 使用通用代理针对同一地址执行，或者

2. X 与 Y 使用同一代理、由同一线程块针对同一地址执行，或者

3. X 与 Y 是别名，且从 X 到 Y 的基础因果路径上存在一个别名代理栅栏。

因果序

因果序按如下方式将基础因果序与若干非传递关系结合起来：

若满足以下条件，操作 X 在因果序中先于操作 Y：

1. X 在代理保持的基础因果序中先于 Y，或者

2. 对于某个操作 Z，X 在观察序中先于 Z，且 Z 在代理保持的基础因果序中先于 Y。

¹ 基础因果序的传递性解释了同步操作的“累积性”。

8.9.6. 一致性序#

存在一个在运行时确定的、关联重叠写入操作的传递偏序，称为一致性序¹。若两次重叠写入操作是道德强的，或者它们在因果序中相关，则它们在一致性序中相关。若两次重叠写入处于数据竞争之中，则它们在一致性序中不相关，这就产生了一致性序的偏序性质。

¹ 一致性序无法被直接观察，因为它完全由写入操作构成。它可以被间接观察——通过它约束读取操作可从中读取的候选写入集合的作用。

8.9.7. 通信序#

通信序是在运行时确定的非传递序，它将写入操作与其他重叠的内存操作关联起来。

1. 若 R 返回由 W 所写的任一字节的值，则写入 W 在通信序中先于重叠的读取 R。

2. 若 W 在一致性序中先于 W’，则写入 W 在通信序中先于写入 W’。

3. 若对于 R 与 W 都访问的任一字节，R 返回由在一致性序中先于 W 的某个写入 W’ 所写的值，则读取 R 在通信序中先于重叠的写入 W。

通信序刻画内存操作的可见性——当内存操作 X1 在通信序中先于内存操作 X2 时，就称 X1 对 X2 可见。

8.10. 公理#

8.10.1. 一致性#

如果写入 W 在因果序中先于与之重叠的写入 W’，那么 W 必须在一致性序中先于 W’。

8.10.2. Fence-SC#

Fence-SC 序不得与因果序相矛盾。对于一对道德强的 fence.sc 操作 F1 和 F2，如果 F1 在因果序中先于 F2，那么 F1 必须在 Fence-SC 序中先于 F2。

8.10.3. 原子性#

单副本原子性

相互冲突的道德强操作以单副本原子性执行。当一次读取 R 与一次写入 W 都是道德强 时，对于 R 和 W 共同访问的字节集合，下列两种通信不可能在同一次执行中同时存在：

1. R 从 W 读取任一字节。

2. R 从在一致性序中先于 W 的任一写入 W’ 读取任一字节。

读改写（RMW）操作的原子性

当一个原子操作 A 与一次写入 W 重叠并且二者都是道德强时，对于 A 和 W 共同访问 的字节集合，下列两种通信不可能在同一次执行中同时存在：

1. A 从在一致性序中先于 W 的写入 W’ 读取任一字节。

2. A 在一致性序中后于 W。

石蕊测试 1

.global .u32 x = 0;

A1: atom.sys.inc.u32 %r0, [x];

A2: atom.sys.inc.u32 %r0, [x];

FINAL STATE: x == 2

当这些操作是道德强时，原子性是得到保证的。

石蕊测试 2

.global .u32 x = 0;

A1: atom.cta.inc.u32 %r0, [x];

A2: atom.gpu.inc.u32 %r0, [x];

FINAL STATE: x == 1 OR x == 2

如果这些操作不是道德强的，则原子性得不到保证。

8.10.4. 无薄空气#

值不得“从薄空气中出现”：一次执行不能以这样的方式推测性地产生一个值，使得该推测通过一串指令 依赖和线程间通信而自我满足。这与程序员的直觉和硬件的实际情况都相符，但在执行形式化分析时有必要 明确陈述这一点。

石蕊测试：带真实依赖的加载缓冲

.global .u32 x = 0;
.global .u32 y = 0;

A1: ld.global.u32 %r0, [x];
B1: st.global.u32 [y], %r0;

A2: ld.global.u32 %r1, [y];
B2: st.global.u32 [x], %r1;

FINAL STATE: x == 0 AND y == 0

这个被称为“LB+deps”（带依赖的加载缓冲，Load Buffering with dependencies）的石蕊测试检查这类可能从薄空 气中产生的被禁止的值。两个线程 T1 和 T2 各从一个变量读取，并把观察到的结果复制到第二个变量，其中第一个 和第二个变量在两个线程之间互换。若每个变量的初始值为零，则最终结果也应为零。如果 A1 从 B2 读取且 A2 从 B1 读取，那么经由本例中内存操作传递的值形成一个环：A1->B1->A2->B2->A1。只有值 x == 0 和 y == 0 被 允许来满足这个环。如果本例中任何一个内存操作推测性地把一个不同的值与相应的内存位置关联起来，那么这样的 推测就会变成自我实现的，从而被禁止。

石蕊测试：不带依赖的加载缓冲

.global .u32 x = 0;
.global .u32 y = 0;

A1: ld.global.u32 %r0, [x];
B1: st.global.u32 [y], 1;

A2: ld.global.u32 %r1, [y];
B2: st.global.u32 [x], 1;

FINAL STATE: x == 1 AND y == 1

这个石蕊测试与上面那个的区别在于，它无条件地向 x 和 y 存储 1。 在这个石蕊测试中，最终状态 x == 1 且 y == 1 是允许的。这次执行并不与 前一个石蕊测试所演示的要求相矛盾。这里不存在自我实现的环——该石蕊测试 将始终且无条件地向 x 和 y 存储 1，因此这里的环没有被自我实现，推测是有效的。

这里依赖的缺失是显而易见的，但实现可以执行任意一串推理来判断 某次存储不依赖于先前的载入，从而打破那些否则会被无薄空气公理 显然禁止的自我实现的环。

这种形式的加载缓冲在 PTX 内存一致性模型中是被有意允许的。

8.10.5. 每个位置的顺序一致性#

在任意一组两两道德强且相互重叠的内存操作中，通信序不得与程序顺序相矛盾， 也就是说，重叠操作之间的程序顺序与通信序中道德强关系的拼接不得形成一个环。这确保 了每个由重叠、两两道德强操作构成的程序切片都严格顺序一致。

石蕊测试：CoRR

.global .u32 x = 0;

W1: st.global.relaxed.sys.u32 [x], 1;

R1: ld.global.relaxed.sys.u32 %r0, [x];
R2: ld.global.relaxed.sys.u32 %r1, [x];

IF %r0 == 1 THEN %r1 == 1

石蕊测试“CoRR”（一致性读-读，Coherent Read-Read）演示了这一保证的一个推论。线程 T1 在位置 x 上执行一次写入 W1，而线程 T2 在同一位置 x 上执行两次（或无限次）读取 R1 和 R2。除模拟 初始值的那一次写入外，x 上不执行其他写入。操作 W1、R1 和 R2 两两道德强。如果 R1 从 W1 读取， 那么随后的读取 R2 也必须观察到相同的值。如果 R2 反而观察到 x 的初始值，那么在通信序中就会 形成一串道德强关系 R2->W1->R1，它与线程 T2 中的程序顺序 R1->R2 相矛盾。因此，在这样的 一次执行中，R2 不能读取 x 的初始值。

8.10.6. 因果性#

通信序中的关系不得与因果序相矛盾。这约束了读操作可能从中读取的候选写操作集合：

1. 如果一次读取 R 在因果序中先于一次与之重叠的写入 W，那么 R 不能从 W 读取。

2. 如果一次写入 W 在因果序中先于一次与之重叠的读取 R，那么对于 R 和 W 共同访问的 任一字节，R 不能从在一致性序中先于 W 的任一写入 W’ 读取。

石蕊测试：消息传递

.global .u32 data = 0;
.global .u32 flag = 0;

W1: st.global.u32 [data], 1;
F1: fence.sys;
W2: st.global.relaxed.sys.u32 [flag], 1;

R1: ld.global.relaxed.sys.u32 %r0, [flag];
F2: fence.sys;
R2: ld.global.u32 %r1, [data];

IF %r0 == 1 THEN %r1 == 1

这个被称为“MP”（消息传递，Message Passing）的石蕊测试体现了典型同步算法的精髓。绝大多数有用的 程序都可以归结为该模式的有顺序的应用。

线程 T1 先向一个数据变量写入，再向一个标志变量写入，而第二个线程 T2 先从标志变量读取，再从数据 变量读取。对标志的操作是道德强的，每个线程中的内存操作由一道栅栏分隔，并且这些 栅栏是道德强的。

如果 R1 观察到 W2，那么释放模式“F1; W2”与获取模式“R1; F2”同步。这建立了因果序 W1 -> F1 -> W2 -> R1 -> F2 -> R2。于是因果性公理保证 R2 不能从在一致性序中先于 W1 的 任何写入读取。在本例中不存在其他写入的情况下，R2 必须从 W1 读取。

石蕊测试：CoWR

// These addresses are aliases
.global .u32 data_alias_1;
.global .u32 data_alias_2;

W1: st.global.u32 [data_alias_1], 1;
F1: fence.proxy.alias;
R1: ld.global.u32 %r1, [data_alias_2];

%r1 == 1

虚拟别名要求在同步路径上有一条 alias 代理栅栏。

石蕊测试：存储缓冲

这个被称为“SB”（存储缓冲，Store Buffering）的石蕊测试演示了由 fence.sc 所强制的顺序一致性。线程 T1 向第一个变量写入，然后读取第二个变量的值，而第二个线程 T2 向第二个变量 写入，然后读取第一个变量的值。每个线程中的内存操作由 fence.sc 指令 分隔，并且这些栅栏是道德强的。

.global .u32 x = 0;
.global .u32 y = 0;

W1: st.global.u32 [x], 1;
F1: fence.sc.sys;
R1: ld.global.u32 %r0, [y];

W2: st.global.u32 [y], 1;
F2: fence.sc.sys;
R2: ld.global.u32 %r1, [x];

%r0 == 1 OR %r1 == 1

在任何一次执行中，要么 F1 在 Fence-SC 序中先于 F2，要么相反。如果 F1 在 Fence-SC 序中 先于 F2，那么 F1 与 F2 同步。这建立了 W1 -> F1 -> F2 -> R2 的因果序。因果性公理确保 R2 不能从在一致性序中先于 W1 的任何写入读取。在该变量不存在任何其他写入的情况下，R2 必须从 W1 读取。类似地，在 F2 在 Fence-SC 序中先于 F1 的情况下，R1 必须从 W2 读取。如果本例中每一处 fence.sc 都被替换为一条 fence.acq_rel 指令，那么这一结果得 不到保证。可能存在这样一种执行，其中每个线程的写入都未被另一线程观察到，也就是说，存在一种执行， 其中 R1 和 R2 分别返回变量 y 和 x 的初始值“0”。

8.11. 特殊情况#

8.11.1. 归约不构成获取模式#

像 red 这样的原子归约操作不会与获取栅栏构成获取模式。

石蕊测试：使用归约指令的消息传递

.global .u32 x = 0;
.global .u32 flag = 0;

W1: st.u32 [x], 42;
W2: st.release.gpu.u32 [flag], 1;

RMW1: red.sys.global.add.u32 [flag], 1;
F2: fence.acquire.gpu;
R2: ld.weak.u32 %r1, [x];

%r1 == 0 AND flag == 2

这个被称为“MP”（消息传递）的石蕊测试演示了 归约被排除在获取模式之外的推论。 可以观察到 R2 从 x 读取值 0 且 flag 的最终值为 2 的这一结果。 该结果是可能的，因为 T1 中的释放模式并不与 T2 中的任何获取模式同步。 改用 atom 指令代替 red 会禁止这一结果。