PTX Parallel Thread Execution ISA 7.4

写在前面的话，
尝试翻译一下PTX参考手册，自己学习一下

Synatax

Source Format

case sensitive

lowercase for keywords

.version 指令 指定PTX语言的版本

.target 指令 指定假定的目标体系结构

Comments

多行注释 /* */

单行注释 //

注释不能出现在字符常量、字符串字面量或其他注释中

注释在PTX中被当作空格对待

Statements

一个PTX语句或者是directive，或者是instruction

以可选的 label 开始，分号结尾

directive关键字以 . 开头

instruction

指令由指令操作码后面跟着一个由0或多个操作数组成的逗号分隔列表组成，并以分号结束。

操作数可以是寄存器变量、常量表达式、地址表达式或标号名称。

指令有一个可选的控制条件执行的保护谓词。

保护谓词紧跟在可选标签之后，在操作码之前，并被写为@p，其中p是谓词寄存器。

保护谓词可以被否定，写成@!p。

identifiers

followsym: [a-zA-Z0-9_$]

identifier: [a-zA-Z]{followsym}* | [_$%]{followsym}+

以字母，后面跟着0或多个字母，数字，下划线，美元符号

以下划线或%开头，后面跟着1或多个字母，数字，下划线，美元符号

Integer Constant

64位 有符号或者无符号 .s64 .u64

hexadecimal literal: 0[xX]{hexdigit}+U?
octal literal: 0{octal digit}+U?
binary literal: 0[bB]{bit}+U?
decimal literal {nonzero-digit}{digit}*U?

字面值是有符号的(.s64)，除非值不能完全用.s64表示或指定了unsigned后缀

Float-Point Constant

浮点数常量用64位双精度值表示，所有浮点常量表达式都使用64位双精度算法计算

唯一的例外是：表示精确单精度浮点值的32位十六进制表示法

浮点数字面量始终以64位双精度浮点数格式表示

PTX包括浮点常量的第二种表示形式，用于使用十六进制常量指定精确的机器表示形式

0[fF]{hexdigit}{8} // 单精度浮点数

0[dD]{height}{16} // 双精度浮点数

例子：

mov.f32 $f3, 0F3f800000;

Constant expression

PTX中的常量表达式不支持整数和浮点之间的强制转换。

运算符优先级和结合性

一元运算符的运算符优先级最高，并且随着图表中的每一行而降低。

同一行上的运算符具有相同的优先级，一元运算符从右向左求值，二元运算符从左向右求值。

整型常量表达式求值

整数常量表达式在编译时根据一组规则进行计算，这些规则确定每个子表达式的类型（signed .s64与unsigned .u64）

求值规则：

1. 除非需要使用unsigned来防止溢出，或者除非字面量使用U后缀，否则字面量是有符号的;
2. 一元加号和减号保留输入操作数的类型;
3. 一元逻辑求反(!)生成值为0或1的有符号结果;
4. 一元逐位补码(~)将源操作数解释为无符号并生成无符号结果;
5. 某些二进制运算符需要对源操作数进行规范化。这种规格化称为通常的算术转换，如果两个操作数中的任何一个都是无符号的，则只需将两个操作数转换为无符号类型;
6. 加法、减法、乘法和除法执行通常的算术转换，并生成与转换后的操作数类型相同的结果。也就是说，如果源操作数中的任何一个是无符号的，则操作数和结果是无符号的，否则是有符号的;
7. 取余符号（%）将操作数解释为无符号;
8. 左移位和右移位将第二个操作数解释为无符号，并生成与第一个操作数类型相同的结果。请注意，右移位的行为由第一个操作数的类型决定：有符号值的右移位是算术运算并保留符号，无符号值的右移位是逻辑运算并在零位移位;
9. AND(&)、OR(|)和XOR(^)执行通常的算术转换，并生成与转换后的操作数类型相同的结果;
10. AND_OP(&&)、OR_OP(||)、Equal(=)和Not _Equal(!=)生成带符号的结果。结果值为0或1;
11. 有序比较（<,<=,>,>=）对源操作数执行通常的算术转换，并生成有符号结果。结果值为0或1;
12. 使用（.s64）和（.u64）强制转换，支持将表达式强制转换为有符号或无符号;
13. 对于条件运算符( ?: )，第一个操作数必须是整数，第二个和第三个操作数要么都是整数，要么都是浮点。通常对第二个和第三个操作数执行算术转换，结果类型与转换后的类型相同。
    

State Spaces, Types, and Variables

GPU上的资源在PTX中通过状态空间和数据类型进行抽象

状态空间

状态空间是具有特定特征的存储区域。所有变量都驻留在某个状态空间中。状态空间的特征包括其大小、可寻址性、访问速度、访问权限以及线程之间的共享级别。

Kernel Function Parameters

每个核函数定义包含一个可选的参数列表。这些参数是可寻址的，声明在.param状态空间的只读变量

kernel parameter variables 被一个网格中的所有线程块所共享

可以使用ld.param指令来访问这些参数的值

核函数的参数可能是正常的值，也可能是指针（在constant，global，local或shared state spaces内的对象的地址）；指针的情形，编译和运行时系统需要知道哪些参数是指针以及它们指向哪个状态空间。内核参数属性指令（kernel parameter attribute directives）用于在PTX级别提供此信息

Kernel Parameter Attribute: .ptr

语法

.param .type .ptr .space .align N varname
.param .type .ptr        .align N varname  .space = { .const, .global, .local, .shared };  

用于指定状态空间，还可以指定指针类型内核参数指向的内存对齐方式。对齐值N（如果存在）必须是2的幂。如果未指定状态空间，则假定指针是指向常量、全局、本地或共享内存之一的通用地址。如果未指定对齐方式，则假定指向的内存与4字节边界对齐。

类型

基本类型

如果两个基本类型具有相同的基本类型且大小相同，则它们是兼容的。
如果有符号和无符号整数类型的大小相同，则它们是兼容的。位大小类型与具有相同大小的任何基本类型兼容。

子字大小的限制使用

.u8、.s8和.b8指令类型仅限于ld、st和cvt指令。.f16浮点类型仅在与.f32、.f64类型的转换、半精度浮点指令和纹理提取指令中允许。.f16x2浮点类型仅在半精度浮点算术指令和纹理提取指令中允许。
为方便起见，ld、st和cvt指令允许源数据操作数和目标数据操作数的宽度大于指令类型大小，因此可以使用规则宽度寄存器加载、存储和转换窄值。例如，当加载、存储或转换为其他类型和大小时，8位或16位值可以直接保存在32位或64位寄存器中

Texture Sampler and Surface Types

内置的不透明类型

其layout, field ordering, base address, and overall size都被一个PTX程序隐藏

三个内置类型：.texref, .samplerref, 和 .surfref

对于处理纹理和采样器，PTX 有两种操作模式：unified模式 和 independent 模式

unified mode： texture and sampler information is accessed through a single .texref handle

independent mode： texture and sampler information each have their own handle, allowing them to be defined separately and combined at the site of usage in the program.

unified mode

width, height, depth : 纹理或者表面每一维度元素数量上的大小

channel_data_type， channel_data_order: 用对应于源语言API的枚举类型指定纹理或表面的这些属性；比如，OpenCL

normalized_coords: 指明纹理或表面使用标准化的坐标[0.0, 1.0)，如果没有被显示指定，默认由运行时系统基于源语言设置

filter_mode: 如何根据输入纹理坐标计算纹理读取返回的值

addr_mode_{0,1,2}: 定义每一位的寻址模式，确定如何处理超出范围的坐标

independent mode

在independent 模式，采样器属性包含在一个独立的 .samplerref 变量中，.texref变量中的这些域失效。一个额外的采样器属性，force_unnormalized_coords 在independent mode可用

force_unnormalized_coords 域允许采用器覆盖纹理头(texture header)的normalized_coords属性；当为 True 时，纹理头设置被覆盖并使用非规范化坐标；当为 False 时，使用纹理头设置。

这个属性被用在编译OpenCL

Channel Data Type and Channel Order Fields

目前OpenCL是唯一能定义这些域的源语言

变量

PTX中，一个变量的定义包括其 变量类型 和 其 状态空间

变量定义

变量声明命名变量所在的空间、变量的类型和大小、名称、可选的数组大小、可选的初始值设定项和可选的变量固定地址。

谓词变量(predicate variables)只能声明在寄存器状态空间中

向量

PTX支持限制长度的向量类型。

长度为2或4的任意非谓词基本类型的向量可以在类型前添加 .v2, .v4 前缀来定义。

向量必须基于一种基本类型

向量长度不能超过128比特，比如，.v4 .f64 是不允许的

3个元素的向量可以用 .v4 向量来定义，这是一种常用模式在三维网格，纹理，等等

默认情况下，向量变量与其整体大小的倍数（向量长度乘以基类型大小）对齐，以启用需要与访问大小的倍数对齐的地址的向量加载和存储指令。

数组定义

为了声明一个数组，变量名后面是维度声明，类似于 C 中的固定大小数组声明。每个维度的大小是一个常量表达式。

若声明时初始化，则数组的第一维可以被省略

初始化

声明变量时可以使用类似C/C++的语法来为变量指定初始值，变量名，等号，初始值

像C语言一样，数组初始化可以不完整，即初始化元素的数量可能比对应数组维度的范围小，剩余数组位置被指定数组类型的默认值初始化

目前，变量初始化只支持常量和全局状态空间。常量和全局状态空间的变量在没有初始化的情况下，默认被初始化为0。外部变量声明不允许使用初始化。

初始化时出现的变量名代表变量的地址，这个可以用来静态初始化一个变量的指针

初始化也可以包含var+offset表达式，offset是被加到var的地址上的字节偏移量。只有.global和.const状态空间中的变量可以被用在初始化中。

默认情况下，结果地址是变量的状态空间的偏移量（就像使用 mov 指令获取变量的地址一样）。generic()运算符，被提供用来创建一个通用地址用作变量初始化。

从PTX ISA 7.1版本开始，mask()运算符被提供，其中mask是一个整型立即数。mask()运算符中允许的表达式只有整型常量表达式和代表变量地址的符号表达式。mask()运算符从初始化中所用的表达式中抽取连续的n个比特，然后把这些比特插入到被初始化变量的最低位。数量n和被抽取比特的开始位置由整型立即数mask指定。PTX ISA 7.1版本只支持从变量的地址中提取从字节边界开始的单个字节。PTX ISA 7.3版本支持整数常量表达式作为 mask() 运算符中的操作数。

mask支持值是 0xFF, 0xFF00, 0XFF0000, 0xFF000000, 0xFF00000000, 0xFF0000000000, 0xFF000000000000, 0xFF00000000000000

注：PTX 3.1 重新定义了初始化器中全局变量的默认寻址，从通用地址到全局状态空间中的偏移量。传统 PTX 代码被视为对初始化程序中使用的每个全局变量都有一个隐式 generic() 运算符。 PTX 3.1 代码应该在初始化程序中包含显式 generic() 运算符，使用 cvta.global 在运行时形成通用地址，或使用 ld.global 从非通用地址加载。

出现在初始化中的设备函数名代表函数中第一条指令的地址，这个可以被用来初始化用于间接调用的函数指针表。PTX ISA 3.1版本开始，核函数名可被用于初始化，例如，初始化核函数指针表，与 CUDA Dynamic Parallelism 一起使用以从 GPU 启动内核。

标签(label)不能被用在初始化中

保存变量或函数地址的变量应该是.u8 .u32 .u64类型

.u8类型只能被用在mask()运算符中

除了 .f16 、 .f16x2 和 .pred 之外的所有类型都允许使用初始化器

对齐

所有可寻址变量的存储字节对齐可以在变量声明时被指定。使用紧跟在状态空间说明符之后的可选 .align 字节计数说明符指定对齐方式。该变量将与一个地址对齐，该地址是字节数的整数倍。对齐值字节计数必须是 2 的幂。对于数组，alignment 指定整个数组的起始地址的地址对齐，而不是单个元素的地址对齐。

标量和数组变量的默认对齐方式是基类型大小的倍数。矢量变量的默认对齐方式是整体矢量大小的倍数。

请注意，PTX的所有访存指令都要求地址与访问大小的倍数对齐。内存指令的访问大小是内存中被访问的总字节数。例如ld.v4.b32的访问大小是16字节，而atom.f16x2的访问大小是4字节。

参数化变量名

由于 PTX 支持虚拟寄存器，因此编译器前端生成大量寄存器名称是很常见的。

PTX支持创建一组有相同前缀字符串跟着整型后缀的变量集合，而不必显示指定每一个变量名。

例如：生命100个.b32类型的变量，变量名分别为 %r0, %r1, %r2, … , %r99。这100个变量定义如下：

.reg .b32 %r<100>;   // declare %r0, %r1, ..., %r99

这种速记语法可以与任何基本类型和任何状态空间一起使用，并且可以以对齐说明符开头。

数组变量不能以这种方式声明，也不允许初始化器

变量属性

.attribute指令用于指定变量的一些特殊属性

关键字 .attribute 后跟括号内的属性说明。多个属性用逗号分隔。

变量属性指令: .attribute

特殊属性：.managed

该属性指明：变量将会被分配在统一虚拟内存环境（unified virtual memory environment）中的一个位置，意味着系统中的主机和其他设备都可以直接访问这个变量。

这个属性只能被用于在.global状态空间的变量上。

指令操作数

操作数类型信息

指令中的所有操作数都具有其声明中的已知类型。每个操作数的类型必须与指令模板和指令类型决定的类型兼容。类型间没有自动转换。

位大小的类型与其它具有相同大小的每个类型都兼容。具有共同大小的整数类型彼此兼容。具有与指令类型不同但与指令类型兼容的类型的操作数被静默地转换为指令类型。

源操作数

源操作数在指令描述中由名称 a、b 和 c 表示。

PTX 描述了一个加载存储机器，因此 ALU 指令的操作数必须全部位于 .reg 寄存器状态空间中声明的变量中。对于大多数操作，操作数的大小必须一致。

cvt（转换）指令采用多种操作数类型和大小，因为其功能就是把操作数从几乎任意指令类型和大小转换到其他任意指令类型和大小。

ld, st, mov, cvt 指令 把数据从一个位置复制到另一个位置。

ld， st指令从可寻址的状态空间移动数据到寄存器，st指令则相反；mov指令在寄存器间拷贝数据

大多数指令都有一个可选的谓词保护（predicate guard）来控制条件执行，少数指令有额外的谓词源操作数。谓词操作数由名称 p、q、r、s 表示。

目的操作数

产生单个结果的 PTX 指令将结果存储在指令描述中由 d（用于目的地）表示的字段中。结果操作数是寄存器状态空间中的标量或向量变量。

使用地址，数组和向量

地址做操作数

所有内存指令都有一个地址操作数，用于指定要访问的内存位置。可寻址操作数包括：

1. [var] 一个可寻址变量var的名字
2. [reg] 一个包含一个字节地址的整数或者位大小（bit-size）类型的寄存器reg
3. [reg+immOff] 一个包含字节地址的寄存器加上一个整型常量字节偏移量（有符号，32位）的和
4. [var+immOff] 一个可寻址变量的地址加上一个整型常量字节偏移量（有符号，32位）的和
5. [immAddr] 一个立即数绝对字节地址（无符号，32位）

包含地址的寄存器比喻被声明为位大小类型或整型

地址大小或者32位，或者64位。可0扩展或者截断

地址算术是使用整数算术和逻辑指令执行的。示例包括指针算术和指针比较。所有地址和地址计算都是基于字节的；不支持 C 风格的指针算法。

mov指令可以被用来将一个变量的地址移动到指针。

通用地址

如果一个内存指令没有指定命名空间，那么操作数被认为是通用地址。

const, local, shared 状态空间 是 通用地址空间中的窗口。每一个窗口通过窗口基地址和与对应状态空间相等的窗口大小定义。通用地址映射到全局内存，除非它位于常量、本地或共享内存的窗口内。

在每个窗口内，通用地址通过从通用地址中减去窗口基址映射到底层状态空间中的地址

数组做操作数

所有类型的数组都可以被声明，数组标识符被看作数组被声明的状态空间的地址常量。程序中的数组大小是一个常数。

数组元素可以通过显示计算的字节地址访问，也可以通过使用方括号表示法索引到数组。方括号内的表达式或者是一个整型常数，一个寄存器变量，或者一个具有常量偏移表达式的简单寄存器，其中偏移量是一个寄存器变量加或者减去一个常量表达式。

向量做操作数

部分指令支持向量操作数，包括 mov, ld, st, tex。在调用函数时也可以传递向量做参数

可以从带有后缀 .x、.y、.z 和 .w 的向量中提取向量元素，以及典型的颜色字段 .r、.g、.b 和 .a。

花括号括起来 拼接 向量

例：

.reg .v4 .f32 V;
.reg .f32     a, b, c, d;
mov.v4.f32 {a,b,c,d}, V;

向量加载和存储可用于实现宽加载和存储，这可以提高内存性能。加载/存储操作中的寄存器可以是向量，也可以是大括号括起来的类似类型标量列表。

ld.global.v4.f32 {a,b,c,d}, [addr+16];
ld.global.v2.u32 V2, [addr+8];

大括号拼接的向量中的元素，例如 {Ra, Rb, Rc, Rd}，可按如下方式提取：

Ra = V.x = V.r
Rb = V.y = V.g
Rc = V.z = V.b
Rd = V.w = V.a

标签和函数名做操作数

标签和函数名只能分别在 bra/brx.idx 和 call 指令中使用。函数名称可用于 mov 指令以将函数的地址放入寄存器，以用于间接调用。

PTX ISA 3.1版本开始，mov 指令可用于获取内核函数的地址，以传递给从 GPU 启动核函数的系统调用。这个特性是CUDA 动态并行支持的一部分

类型转换

所有算术、逻辑和数据移动指令的所有操作数必须具有相同的类型和大小，但更改大小和/或类型是指令定义的一部分的操作除外。不同大小或类型的操作数必须在操作之前进行转换。

标量转换

cvt指令使用的给定不同类型的操作数的精度和格式

超出浮点数范围的浮点转换用最大浮点值表示（f32 和 f64 为 IEEE 754 Inf，f16 为 ~131,000）

扩展类型（零扩展或者符号扩展）基于目标格式

舍入修饰符

共四种整型舍入修饰符，四种浮点修饰符

操作数成本

来自不同状态空间的操作数影响操作的速度。寄存器最快，全局内存最慢。

许多内存延迟可以通过多种方式隐藏。

1. 多线程执行，硬件可以发出内存操作然后切换到其他执行
2. 尽可能早地发出load指令，因为执行不会被阻塞直到所需结果被接下来的指令使用时

存储操作中的寄存器可用得更快

使用不同类型内存的预估时间开销：

Abstracting the ABI

ABI 定义了如何在机器代码中访问数据结构或计算例程，这是一种低级、依赖于硬件的格式。相比之下，API 在源代码中定义了这种访问，这是一种相对高级、独立于硬件、通常是人类可读的格式。 ABI 的一个常见方面是调用约定，它确定如何将数据作为计算例程的输入提供，或作为计算例程的输出读取。

PTX 没有公开特定调用约定、堆栈布局和应用程序二进制接口 (ABI) 的详细信息，而是提供了一个稍微更高级别的抽象并支持多个 ABI 实现。

函数声明和定义

PTX中用.func directive来定义和声明函数

函数声明指定一个可选的返回参数列表，函数名，可选的输入参数列表；这些一起指明了函数的接口或称原型。

函数定义指定了函数的接口和函数体。

函数被调用前必须提前进行声明或者调用

标量和向量基本类型的输入和返回参数可以简单地表示为寄存器变量。在调用时，参数可以是寄存器变量或常量，返回值可以直接放入寄存器变量中。调用中的参数和返回变量的类型和大小必须与被调用者的相应形式参数匹配

使用 ABI 时，.reg 状态空间参数的大小必须至少为 32 位。源语言中的子字标量对象应提升为 PTX 中的 32 位寄存器，使用接下来描述的 .param 状态空间字节数组。

类似C语言中的结构体或者联合体(union)的对象在PTX中被展开（flatten）放入到寄存器或者字节数组中，并且使用.param空间的内存表示。

.param 状态空间在设备函数中的使用：

对于调用者：

被用于设置将要被传递给被调用函数的值，或者从被调用函数接收返回值。通常，一个.param字节数组被用来将按值传递的结构体的字段收集在一起

对于被调用者：

接收参数值，或者 将返回值传回调用者

参数传递中的限制：
对于调用者：

1. 参数必须是.param变量或者.reg变量
2. 对于形参是.param字节数组的情形，参数必须也是.param字节数组带有匹配的类型，大小和对齐。.param参数必须被声明在调用者的局部作用域(local scope)内
3. 对于形参是.param的基本类型标量或者向量时，对应的参数应为.param或.reg状态空间同时具有匹配的类型和大小，或者是可以用形式参数的类型表示的常量
4. 对于形参是.reg状态空间时，对应的参数应为.param或.reg状态空间同时具有匹配的类型和大小，或者是可以用形式参数的类型表示的常量
5. 对于形参是.reg状态空间时，寄存器的大小最少是32位
6. 所有用于向函数调用传递参数的 st.param 指令必须紧跟在相应的调用指令之前，用于收集返回值的 ld.param 指令必须紧跟在调用指令之后，没有任何控制流程改变。用于参数传递的 st.param 和 ld.param 指令不能被断言（cannot be predicated）。这将启用编译器优化并确保 .param 变量不会在调用者的框架中消耗超出 ABI 所需空间的额外空间。 .param 变量只是允许在调用点在可能位于多个位置的数据之间进行映射（例如，调用者操作的结构位于寄存器和内存中）到可以作为参数或返回值传递的内容给被调用者。

对于被调用者：

1. 输入和返回参数只能是.param或.reg变量
2. .param的参数必须被对齐到1，2，4，8或者16字节的倍数
3. .reg参数的大小必须至少32位
4. .reg 状态空间可用于接收和返回基类型标量和向量值，包括在非 ABI 模式下编译时的子字大小对象。 .reg 状态空间提供legacy支持

参数传递时选择.reg或.param状态空间对最终参数传递到物理寄存器还是栈没有影响。参数映射到物理寄存器和栈的位置依赖于ABI定义和参数的顺序，大小和对齐方式。

从PTX 1.x版本以来的变化

PTX1.x版本限制形参只能位.reg状态空间，不支持数组做参数。类似C语言中结构体的对象被展开，用多个寄存器来传递或返回。PTX1.x支持多个返回值，也是基于上述原因。

PTX2.0开始，形参或者时.reg或者是.param状态空间，同时.param支持数组。对于目标机器sm_20及以上，PTX限制函数具有单个返回值，.param字节数组被用于返回不适合寄存器的对象。目标机器sm_1x继续支持多个返回寄存器。

PTX 仅为目标 sm_20 或更高版本实现基于堆栈的 ABI。

PTX3.0之前的版本允许.reg和.local状态空间的变量被定义在模块作用域。PTX3.0及以后的版本在使用ABI编译时不允许在模块作用域中定义.reg和.local变量，限制只在函数作用域中使用它们。不使用ABI编译时.reg和.local变量还像以前被支持。当编译遗留PTX代码（3.0之前）包含模块作用域的.reg或.local变量时，编译器静默取消使用ABI。

可变参数函数

PTX 6.0 版支持将未定义大小的数组参数传递给可用于实现可变参数函数的函数。

Alloca （memory allocated on the stack 栈上分配的内存）

PTX提供alloca指令来在运行时在每个线程的本地内存堆栈上分配存储空间。被分配的栈内存可通过alloca返回的指针使用ld.local和st.local指令访问

stacksave 允许读取局部变量中堆栈指针的值

stackrestore 可以用保存的值恢复堆栈指针

内存一致性模型

多线程执行时，每个线程的内存操作的副作用对其他线程可见，而且是以部分和不相同的顺序。这意味着对于不同的线程，任何两个操作都可能无序或以不同的顺序发生。内存一致性模型引入的公理明确规定了不同线程观察到的顺序之间禁止出现哪些矛盾。

范围和模型适用性

此模型下指定的约束适用于具有任何PTX ISA 版本号、在 sm_70 或更高版本架构上运行的 PTX 程序。
内存一致性模型不适用于纹理和表面访问。

系统作用域内原子性的限制

与主机 CPU 通信时，在某些系统上可能无法原子地执行具有系统范围的 64 位强操作。

内存操作

PTX内存模型的基本存储单元是字节，8比特。

PTX 程序可用的每个状态空间都是内存中的一个连续字节序列。

PTX状态空间中每一个字节相对于访问相同状态空间的所有线程都有一个独一无二的地址。

每个PTX内存指令指定一个内存地址和一个数据类型。两者一起定义了一个内存位置，字节范围从地址扩展到数据类型的字节大小。每个内存指令也指定操作

重叠 overlap

两个内存操作重叠：相应的内存位置有重叠

两个内存操作完全重叠：两个内存位置相同

向量数据类型

内存一致性模型将在内存位置上执行的操作与标量数据类型相关联，标量数据类型的最大大小和对齐为 64 位。具有向量数据类型的内存操作被建模为一组具有标量数据类型的等效内存操作，以未指定的顺序对向量中的元素执行。

Packed Data-types 打包数据类型

.f16x2 包含 两个向量内存位置的.f16值

在打包数据类型.f16x2上的内存操作被建模成一对等价的在标量数据类型.f16上的内存操作，以未指定的顺序在打包数据的每个元素上执行

初始化

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

状态空间

内存一致性模型中定义的关系独立于状态空间。特别是，因果关系顺序关闭所有状态空间中的所有内存操作。

但是在一个状态空间中内存操作的副作用只能通过也可以访问相同状态空间的操作直接观察到。除了作用域之外，这进一步限制了内存操作的同步效果。

操作类型

作用域

每个强操作都必须指定一个作用域，这是一组可以直接与该操作交互并建立内存一致性模型中描述的任何关系的线程。

请注意，warp 不是作用域； CTA （线程block）是符合内存一致性模型作用域的最小线程集合。

Morally strong operations

同时满足下面两个条件的操作成为morally strong relative：

1. 操作按程序顺序相关（比如，都由同一个线程执行）或者每个操作都是strong，并指定了一个作用域，其中包括执行另一个操作的线程；
2. 两个都是内存操作，而且完全重叠

内存一致性模型的大多数公理依赖于morally strong operations之间的联系

冲突和数据竞争

两个重叠的内存操作是冲突的（conflict）：至少一个是写操作

两个冲突内存操作是数据竞争的：如果它们在因果关系上没有关联，并且不是morally strong

混合大小数据竞争的限制

完全重叠操作间的数据竞争被称为 统一大小的数据竞争

部分重叠操作间的数据竞争被称为 混合大小的数据竞争

内存一致性模型的公理不适用于PTX程序中包含一个或更多混合大小数据竞争的情形。但这些公理足够描述只含一致大小的数据竞争的PTX程序的行为。

混合大小RMW操作的原子性

在任何有或没有混合大小数据竞争的程序中，以下属性适用于每对重叠的原子操作 A1 和 A2，这样每个操作都指定了一个包含另一个的作用域：A1 指定的读-修改-写操作是在启动 A2 之前完全执行，反之亦然。

无论两个操作 A1 和 A2 是部分重叠还是完全重叠，此属性都成立

这个是干嘛的？？

Release and Acquire patterns

一些指令序列会产生参与内存同步的模式

release模式使得现在线程的之前操作对其他线程的一些操作可见

acquire模式使得其他线程的一些操作对现在线程的后续操作可见

对于释放和获取模式，这种影响通过因果顺序的传递性质进一步扩展到其他线程中的操作

位置 M 上的释放release模式由以下之一组成：

1. M上的释放操作

   st.release [M];
   或者
   atom.acq_rel [M];
   

2. 一个M上的释放操作跟着一个在M位置上的strong写，按照程序顺序

   st.release [M];
   st.relaxed [M];
   

3. 一个fence跟着一个在M上的strong写，按照程序顺序

   fence;
   st.relaxed [M];
   

任何由release模式建立的内存同步仅影响发生在这个模式的第一条指令之前按程序顺序发生的操作

位置 M 上的acquire模式由以下之一组成：

1. M上的acquire操作

   ld.acquire [M];
   或者
   atom.acq_rel [M];
   

2. 一个在M位置上的strong读跟着一个在M位置上的acquire操作，按照程序顺序

   ld.relaxed [M];
   ld.acquire [M];
   

3. 一个M位置上的strong读跟着一个fence，按照程序顺序

   ld.relaxed [M];
   fence;
   

任何由acquire模式建立的内存同步仅影响发生在这个模式的最后一条指令之后按程序顺序发生的操作

Ordering of memory operations 内存操作的顺序

每个线程执行的操作序列被捕获为程序顺序（program order），而跨线程的内存同步被捕获为因果顺序（causality order）。内存操作对其他内存操作的副作用的可见性被捕获为通信顺序（communication order）。内存一致性模型定义了一方面通信顺序，另一方面因果顺序和程序顺序之间不允许的矛盾

The memory consistency model defines contradictions that are disallowed between
communication order on the one hand, and causality order and program order on the other.

这句话啥意思？？？

程序顺序 program order

程序顺序将线程执行的所有操作与顺序处理器在相应 PTX 源中执行指令的顺序相关联。它是一种传递关系，它在线程执行的操作上形成全序，但不关联来自不同线程的操作。

观察顺序 Observation order

观察顺序通过一个可选的原子 读-修改-写 操作序列把写W关联到R。

一个写W以观察顺序先于读R，若：

1. R和W是morally strong 并且 R读由W写的值
2. 或者，对一些原子操作Z，在观察顺序上，W先于Z并且Z先于R

Fence-SC 顺序

Fence-SC顺序是一个运行时决定的无环偏序，关联每一对morally strong的 fence.sc 操作

内存同步

同步的效果是建立线程间的因果关系（causality order）

1. 如果X以Fence-SC顺序先于Y，则一个fence.sc 操作X 与一个fence.sc 操作Y 同步
2. bar.sync 或 bar.red 或 bar.arrive 操作与在同一屏障上执行的 bar.sync 或 bar.red 操作同步
3. 如果一个写操作X以观察顺序先于一个读操作Y，并且X的第一个操作和Y的最后一个操作是morally strong的，那么释放模式X于acquire模式Y同步

API同步

同步关系也可以由特定的CUDA API建立

1. CUDA 流中排队的任务的完成与同一流中后续任务的开始同步（如果有）
2. 出于上述目的，记录或者等待流中的CUDA事件，或者由于cudaStreamLegacy插入的跨流的屏障，即使没有直接副作用，也会将任务排入相关流。事件记录任务与匹配的事件等待任务同步，屏障到达任务与匹配的屏障等待任务同步
3. CUDA 内核的启动与内核中所有线程的启动同步。内核中所有线程的结束与内核的结束同步
4. CUDA图的启动与图中所有节点的开始同步。CUDA图中所有汇节点的完成与图的完成同步。图节点的完成与其直接依赖的所有节点的开始同步
5. 调用CUDA API来入队任务与开始任务同步
6. 流中最后一个排队任务完成，如果有，与从cudaStreamSynchronize返回同步。最近排队的匹配事件记录任务的完成（如果有）与从 cudaEventSynchronize 返回的同步。同步 CUDA 设备或上下文的行为就像同步上下文中的所有流，包括已销毁的流
7. 从 API 返回 cudaSuccess 以查询 CUDA 句柄（例如流或事件）的行为与从匹配的同步 API 返回的行为相同

因果顺序 Causality Order

因果关系顺序捕获内存操作如何通过同步操作在线程间变得可见。公理用因果顺序来约束一组写操作，读操作可以从中读取一个值。

因果关系顺序中的关系主要由基本因果关系顺序中的关系组成，即在运行时确定的传递顺序。

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

Base causality order 基本因果关系顺序

操作X以基本因果关系顺序先于操作Y，若

1. X与Y同步，或者
2. 对一些Z操作
   1. X以程序顺序先于Z，且Z以基本因果关系先于Y，或
   2. X以基本因果关系先于Z，且 Z以程序顺序先于Y，或
   3. X以基本因果关系先于Z，且 Z以基本因果顺序先于Y

因果关系

因果关系结合基本因果关系和一些如下的非传递关系

X操作以因果关系 先于操作Y，若

1. X以基本因果顺序先于Y，或
2. 对一些操作Z,X以观察顺序先于Z，且
   1. Z以基本因果顺序先于Z，或
   2. Z以程序顺序先于Y，且Z与Y重叠

一致性顺序 Coherence Order

存在在运行时确定的与重叠写入操作相关的部分传递顺序，称为一致性顺序

两个重叠的写操作如果它们是morally strong或者因果顺序相关的，那么就是一致性顺序相关的。

如果两个重叠写入处于数据竞争中，则它们在一致性顺序上是不相关的，这导致了一致性顺序的部分性质。

无法直接观察一致性顺序，因为它完全由写操作组成。可以通过使用它来约束读取操作可以从中读取的候选写入集来间接观察到这一点。

communication order

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

1. A 写 W以通信顺序先于一个重叠读R，如果R返回的值由任意字节是由W写的
2. A写W以通信顺序先于一个写W’，如果W以一致性顺序先于W’
3. A读R以通信顺序先于一个重叠写W，如果对于R和W都访问的任意字节，R返回的值由以一致性顺序先于W的W‘写

通信顺序捕获内存操作间的可见性，当一个内存操作X以通信顺序先于内存操作X2，那么说X1是对X2可见的

公理Axioms

一致性 coherence

如果写操作W以因果顺序先于重叠的写操作W’，那么W必须以一致性顺序先于W’

Fence-SC

Fence-SC 顺序不能与因果顺序矛盾。对于一对morally strong的fence.sc操作F1和F2，如果F1 以因果顺序先于F2，那么F1必须以Fence-SC顺序先于F2

原子性

单拷贝（Single-Copy）原子性

冲突的morally strong操作以以单拷贝原子性执行。当读操作R和写操作W是morally strong的，那么对R和W都访问的字节集合，接下来的通信不能同时在一次执行中出现：

1. R从W读任意字节
2. R从以一致性顺序先于W的W’读任何字节

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

当原子操作A和写操作W重叠且morally strong，那么对于A和W都访问的字节集合，接下来的两个通信不能在一次执行中同时存在：

1. A从以一致性顺序先于W的W’读任意字节
2. A以一致性顺序排在W后

No Thin Air

值可能不会“凭空出现”：执行不能推测性地产生值，使得推测通过指令依赖链和线程间通信变得自我满足。这符合程序员的直觉和硬件现实，但在执行形式分析时有必要明确说明。

Litmus Test: Load Buffering

每个位置的顺序一致性

在任意一组成对morally strong的重叠内存操作中，通信顺序不能与程序顺序冲突，例如 重叠操作和通信顺序morally strong的联系间的一系列程序顺序不能成环。
这个保证了成对morally strong的重叠内存操作的每个程序片是严格顺序一致的

Litmus Test: CoRR

因果关系

通信顺序的关系不能顶撞因果顺序。这个限制了读操作可以读的候选写操作集合：

1. 若读操作R以因果顺序先于一个重叠的写操作W，那么R不能从W读；
2. 若写操作W以因果顺序先于一个重叠的读操作，那么对于R和W都访问的任何字节，R不能从任何以一致性顺序先于W的写操作W’处读

Litmus Test: Message Passing

Litmus Test: Store Buffering