NVIDIA GPU PTX - Memory Consistency Model

在multi-threaded执行模型中，每个Thread所执行内存操作的side-effect，会以不完整且非一致的顺序对其他Thread可见。
没有Memory Consistency Model的情况下，读取操作会返回对同一内存位置的写入操作所提交的值的集合中的任意一个。

假设我们有两个Thread和两个初始值为 0 的共享变量 x 和 y。

Thread1	Thread 2
x = 1;	y = 1;
r1 = y;	r2 = x;

没有Memory Consistency Model的情况下，在某一时刻去观测，可能的结果如下(x的值的集合：{0， 1}， y的值的集合：{0， 1}，r1的值的集合：{y}, r2的值的集合：{x})：

• r1 = 0, r2 = 0
• r1 = 1, r2 = 0
• r1 = 0, r2 = 1
• r1 = 1, r2 = 1

Memory Consistency Model约束了读取操作返回值的候选集合。

对上述示例，如果使用Sequential Consistency模型，则r1= 0, r2 = 0的结果不可能出现。

和State Space的关系：memory consistency model的定义独立于state space。但Memory Operation在一个State Space中的Side-Effect只会被能访问该State Space的其他Operation观测到，这在Scope之外进一步限制了同步效果。例如ld.relaxed.shared.sys和ld.relaxed.shared.cluster的同步效果是一样的，因为cluster外的线程不能执行一个访问shared memory的Memory Operation。

Memory operations

一个PTX Memory instruction包含：

• Operation：操作类型，包括：

Operation Type	Instruction/Operation
atomic operation	atom or red instruction.
read operation	ld指令的所有变种和atom指令 (不包含red).
write operation	st指令的所有变种和产生写操作的atomic指令
memory operation	read + write
volatile operation	.volatile修饰的指令
acquire operation	.acquire或.acq_rel修饰的指令
release operation	.release或.acq_rel修饰的指令
mmio operation	.mmio修饰的指令
memory fence operation	membar, fence.sc, fence.acq_rel
proxy fence operation	fence.proxy，membar.proxy
strong operation	memory fence operation, 或者.relaxed, .acquire, .release, .acq_rel, .volatile, .mmio修饰的memory operation
weak operation	.weak修饰的指令
synchronizing operation	barrier instruction, fence operation, release operation，acquire operation.

• 1个Address Operand：包含一个VA（Virtual Address），会在真正访存时转换为PA（Physical Address）。其中multimem Address是一种特殊的VA，指向了多个PA。只有multimem.* operations 可以操作multimem Address。
• Data Type。对Vector Data Types/Packed Data Types，这两种Data Type的Memory operation被建模为一组等价的Scalar类型的Memory Operation，元素间的Memory Order是不确定的。

Operation types

mmio Operation

mmio operation用.mmio修饰符指示，用来进行IO操作。
从Memory Consistency Model的视角，是一种特殊的strong operation，具有额外的属性：

• Write不会被合并且总会被执行。
• Read总会被执行，且不会forward，prefetch，combine，cache hit。

volatile Operation

volatile operation用.volatile 修饰符指示，用来进行IO操作。等价于system scope的relaxed，但有额外的约束：

• 编译器保证volatile instruction数量保持不变。
• 硬件可以合并volatile operations。
  PTX volatile operations主要用于lowering CUDA C++程序中的volatile，相比直接使用strong operations，性能会差一些。

Scope

Strong Operation必须指定一个Scope，Scope是一组Threads，直接与该Operation交互并建立Memory Consistency Model中描述的关系。Scope有四种：

Scope	Description
.cta	和当前Thread在同一个CTA中的Threads
.cluster	和当前Thread在同一个Cluster中的Threads
.gpu	当前Program中和当前Thread在同一个Device中的Threads, 包含其他kernel的grids
.sys	当前Program中所有Device上的所有Threads，和Host Program的所有Threads。

Proxies

proxy（memory proxy）是一个访存方法的abstract label。
两个memory operation使用不同的访存方法，就称作不同的proxy。
在Operation types中定义的Memory operations，属于generic proxy，textures和surfaces属于不同的proxy。
proxy fence用于同步不同的proxy的memory operation。
此外，Virtual Aliases(一个PA对应多个VA)虽然都使用generic proxy，但还是需要proxy fence来保证内存一致性。

示例如下：

// 同一物理内存的两个虚拟地址
__device__ float* virtual_ptr1;
__device__ float* virtual_ptr2; // 指向与virtual_ptr1相同的物理内存

__global__ void virtual_alias_example() {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    
    if (tid < 256) {
        // 通过虚拟地址1写入
        virtual_ptr1[tid] = tid * 2.0f;
        
        // 代理栅栏 - 确保不同虚拟地址间的可见性
        __threadfence_system();
        
        // 通过虚拟地址2读取 - 现在能保证看到更新
        float verified_value = virtual_ptr2[tid];
    }
}