PTX ISA 7.4 参考手册翻译

阅读量：

文章目录

PTX并行线程执行ISA版本号为7.4
- Syntax
- 源格式
- Notes
- Lines
- Command
identifiers
- 整数常数
- 浮点常数
- 常数表达式
- 整型常数表达式的计算
状态空间及其分类与变量表示
- 状态类别
- 核函数参数设置
- 核参数属性字段指定为：.ptr
  - 类型
  - - 基本类型
- 子字大小的限制使用
Texture Sampling Methods and Surface Type Implementations
统一模式
- 独立模式
- 基于通道的数据类型与排列顺序字段
变量名
- 变量名定义
  - 向量
  - 数组的定义
  - 初始化过程
  - 对齐策略
  - 参数化的变量名称
  - 变量属性描述
  - 变量属性操作指令：指定属性的操作指令：.attribute

指令中的操作数
* * 操作类型的相关信息
* 源端的操作数值
* 目标端的操作数值
* 在内存存储空间中使用数组或向量来存储中间计算结果
* 地址字段作为操作对象
* 全局通用寄存器

复制代码

* 数组做操作数
* 向量做操作数
* 标签和函数名做操作数

  * 类型转换
  * * 标量转换
* 舍入修饰符

  * 操作数成本

* Abstracting the ABI
* * 函数声明和定义
  * * 从PTX 1.x版本以来的变化

  * 可变参数函数
  * Alloca （memory allocated on the stack 栈上分配的内存）

* 内存一致性模型
* * 范围和模型适用性
  * * 系统作用域内原子性的限制

内存操作
* * 重迭 overlap
- 向量数据类型
- Packed Data-types 打包数据类型
- 初始化
状态空间：
- 操作类型：
  - 运算范围：
    - 具有道德强度的操作：
      - 在竞态情况下的表现：
被发现存在数据竞态问题的可能性：
- 在混合规模数据操作中受到限制：
  - 被保证具有严格的互斥性特征：
    - RMW操作的严格互斥性表现：
释放与获取模式 Release and Acquire patterns
- 内存操作序列 Ordering of memory operations
  - 程序流程 program sequence
- 观察流程 Observation sequence
- Fence-SC流程 Fence-SC sequence
- 内存同步 Memory synchronization mechanism
- API同步 API synchronization mechanism

因果次序 Causality Order
- 基本因果次序 Base causality order * 因果关系
  - 一致性顺序 Coherence Order
  - communication order

公理 axioms
- - 一致性 coherence
Fence-SC
原子性 atomicity
- 单个副本的原子性 single-copy atomicity
- 读–修改–写的原子性 RMW atomicity operation's atomicity
  - No Thin Air
  - 每个位置的顺序一致性
  - 因果关系

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

十六进制字面量：采用前缀0x或0X并附加相应数字的表示形式。
八进制字面量：采用前缀0并附加相应数字的表示形式。
二进制字面量：采用前缀0b或0B并附加相应二进制位的表示形式。
十进制字面量：由一个非零首数字开头后跟零或多个后续数字组成的表示形式。

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

Float-Point Constant

每一个浮点常量表达式都采用64位双精度算法计算

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

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

PTX包含用于表示浮点数的第二种格式 ，以便通过十六进制常数明确指定精确的机器数值表达方式

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

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

例子：

复制代码

    mov.f32 $f3, 0F3f800000;
    
    
      
    
    代码解读

Constant expression

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

运算符优先级和结合性

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

同行中的运算符遵循相同的计算顺序规则，在同一行中的一元运算符将按照从右到左的顺序进行计算，在涉及二元运算时则将按照从左到右的方式依次计算

整型常量表达式求值

整数常量表达式在编译阶段 被计算为遵循一组规则来决定每个子表达式的类型：带有符号标志位的有.s64和无符号标志位的有.u64。

求值规则：

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

State Spaces, Types, and Variables

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

状态空间

State space is a storage region characterized by specific attributes. All variables are contained within a specific state space. The attributes of the state space include its size, addressability, access speed, access permissions, and the degree of sharing among threads.

Kernel Function Parameters

由每个核函数定义所包含的一个可选参数集合构成。这些参数具有寻址能力，在(param)状态空间中的不可变变量。

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

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

核函数的参数可能为正常的值或指针，在constant、global、local或shared state spaces中的对象地址；kernel parameter attribute instructions用于在PTX级别提供相关信息

Kernel Parameter Attribute: .ptr

语法

复制代码

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

用于确定状态空间的同时也可以设定指针类型内核参数所指内存的对齐策略。若未明确设置状态空间则默认该指针指向的是全局静态本地或共享内存中的通用地址若无特别说明，默认该内存块将与4字节边界对齐

类型

基本类型

如果两个基本类型在基本类型和大小上一致，则它们是兼容的。
带有符号整数类型的任何数据字节与其都是兼容的。
带位大小类型的任何基本数据字节与其都是兼容的。

子字大小的限制使用

.u8、.s8以及.b8指令类别主要应用于ld、st以及cvt指令。对于.f16浮点类型而言，在与其对应的.f32及.f64类型的转换过程中以及在半精度浮点操作与纹理提取指令中均被允许使用。类似地，在涉及半精度浮点运算时（如.f16x2），此类操作通常只会在半精度浮点算术运算及纹理提取指令中得到应用。
为了简化操作，在ld/ST/CVT等指令中允许源数据操作数与目标数据操作数的宽度超出当前指定的尺寸范围。例如，在涉及不同数据类型的存储与加载时（如宽度为8位或16位的数据），它们可以直接存放在宽度更大的寄存器（如32位或64位）中。

Texture Sampler and Surface Types

内置的不透明类型

其布局、字段顺序、基地址以及总体大小采用了PTX技术来处理

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

在管理纹理与采样器时

unified manner: texture and sampler information is retrieved via a single .texref handle.

independent mode： texture and sampler information are assigned unique handles individually, enabling them to be defined separately before being integrated at the point of use within 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变量中则不再维护这些信息。此外，在independent模式下新增了一个可选的采样器属性(force_unnormalized_coords)，它允许用户在需要时进行设置。

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

现在，在对变量进行初始化操作时，仅允许对常量以及全局状态空间中的变量进行操作。当这些类型的变量未被指定初始值时，默认会被设置为零值。对于非局部变量而言，在进行声明时禁止对其初始值进行隐式设定。

在初始化过程中产生的这些名称用于指示变量内存地址的位置；其中每个名称都可以作为静态初始化某个指针所使用的基础。

初始化也可以包括var+offset表达式,其中offset是被加至var地址的字节偏移量.**仅在(global)和(const)状态空间中的变量可用于初始化.

通常情况下，默认的结果地址是变量状态空间中的偏移值（类似于使用mov指令获取变量地址的方式）。generic()运算符用于生成一个通用地址作为变量初始化的基础。

自PTXISA7.1版本起,该运算符得以引入,其中掩码(mask)为一个整型常量数值。该运算符仅允许使用整型常量作为掩码,并可表示目标变量地址符号表达式。该运算符的功能是通过掩码值mask确定提取的数量n及其起始位置来从被初始化变量中提取连续的n个比特,并将这些比特直接放置于目标变量的最低有效位位置上。在PTXISA7.1版本中仅允许从变量地址处提取位于字节边界上的单个字节;而到了PTXISA7.3版本则增加了对整数常量表达式作为掩码的支持,以实现更为灵活的操作功能。

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

注释：PTX 3.1 对初始化器中全局变量的默认寻址进行了重新定位，将默认寻址范围从通用地址转换为全局状态空间中的偏移量。传统 PTX 代码被视作在初始化程序使用的每个全局变量上都隐式的设置了 generic() 运算符。PTX 3.1 的代码建议应在初始化程序中明确包含显式的 generic() 运算符，并建议在运行时通过 cvta.global 形成通用地址，或者通过 ld.global 来自非通用地址加载数据。

在初始化阶段中使用的设备函数名表示该设备函数在其执行第一条指令时所需的起始地址。这个值通常会被用来构建指向那些需要间接调用的其他功能的指针表。自 PTXISA 3.1版本起可用作构建基址表的基础之一。例如说，在这种模式下核函数可用于构建基址表并结合动态并行性技术时即可从GPU启动内核。

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

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

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

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

对齐

每个可寻址变量在其存储区域中的字节边界都可以在变量声明时进行精确配置。通过紧跟状态空间说明符后使用的可选 .align 字节数目标记符来指定内存对齐策略。这样配置后，该变量将与特定的一个内存边界对齐，并且这个内存边界必须是整数倍的字节间隔。需要注意的是，在这种情况下所选择的位移值的字节数必须是2的幂次方。对于数组类型的数据来说，在编译阶段指定的是整个数组起始位置相对于内存块边界的位移值而非单个元素相对于其所在的存储单元的位置。

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

请留意：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

该属性表明：变量会被分配到同一个统一的虚拟内存环境中，并且系统中的主机以及所有设备都能直接访问这个变量。

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

指令操作数

操作数类型信息

每一个操作数都具有其声明中的已知类型。每个操作数的类型必须与指令模板及指令类型的指定类型的兼容性要求得以满足。 类型间不允许存在自动转换。

尺寸相关的元素不仅能够与其他所有与其尺寸一致的不同类型的元素均能相互配合，在同一尺寸整数类型的各个实例之间也均能相互配合。对于那些虽然尺寸不同于指令类型的其他类别但仍可被隐式转换为指令类型的元素（即操作数），它们会被无差别地转换到指定的操作码空间中。

源操作数

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

PTX 的描述涉及到了一个用于加载存储机器的设计方案，在该方案中定义了操作数的具体存储位置以确保数据完整性与操作一致性

该指令支持多种操作数类型和大小作为输入，并将其转换为其他任意操作数类型和大小作为输出其主要功能就是实现不同操作数之间的相互转换

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

ld指令将数据从可寻址的数据源空间转移至寄存器中；st指令则将数据从寄存器中逆转至目标存储位置；mov指令用于在寄存器之间进行数据复制。

大多数指令都具备一个可选设置用于管理条件判断过程；个别指令则会引入额外的操作数据源用于处理相关功能。这些变量通常被标记为p, q, r, s等。

目的操作数

执行单一目标的 PTX 指令将在指令体说明中指定的目的字段中存储其运算结果。该运算值属于寄存器状态空间中的单一数值或向量数据类型。

使用地址，数组和向量

地址做操作数

每个内存指令都包含一个地址操作数，并且其中可寻址的操作数包括：

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

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

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

地址运算主要采用整数运算与逻辑指令进行处理。具体实例涵盖指针运算与比较操作。所有内存地址及其相关计算均基于字节运算；该系统完全排除C语言风格的指针操作方式。

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

通用地址

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

该状态分为 const、local 和 shared 三类，在统一的通用地址空间中形成特定区域。每个区域则由其起始偏移地址及与其相对应的空间类型所限定的空间大小共同确定。统一虚拟地址通常映射至全局物理内存中；但若该虚拟地址落在特定区域（如常量块、本地区域或共享块）内的相应位置，则不会发生这种映射关系。

对于每个窗口，在其对应的全局地址中进行计算时，会将该全局地址与当前窗口基址进行相减运算以获得底层状态空间中的局部地址

数组做操作数

各类别类型的数组均具备声明的可能性，在类型系统中，数组标识符被视为其状态空间地址常量的映射。在程序设计中，默认情况下所有数组的大小均以固定数值存在。

通过计算字节地址来访问数组元素是可行的方式；另外一种方式是利用方括号表示法来索引数组元素。其中方括号内的内容可能是整型常数、一个寄存器变量或是包含固定偏移量的简单寄存器；这些寄存器中的偏移量通常由另一个寄存器变量加上或减去一个常量表达式来确定。

向量做操作数

该指令集包含用于处理向量操作的指令类型，并具体涵盖了mov、ld、st以及tex这几种指令类型。当执行函数调用操作时，则可将相应的向量作为输入参数进行传递并使用

从具有后缀.x、.y等标记的向量中获取向量元素，并配合典型的色彩通道字段.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指令来将函数地址加载至寄存器，并在此基础上实现间接调用功能。

自PTXISA 3.1版本起可采用mov指令来获取内核函数地址并将其传递给通过GPU启动内核函数所需的系统调用请求。这一特性属于CUDA动态并行功能体系中的一部分

类型转换

所有算术、逻辑及数据移动指令的操作数都必须满足相同类型的条件，并且只有当这些操作属于指令定义的范围之内时（即更改大小或类型不属于常规操作），才允许改变其大小与/或数据类型。对于不同类型的数据（即不同数值域或变量名），在执行运算之前需进行统一转换以确保运算的有效性与准确性。

标量转换

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

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

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

舍入修饰符

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

操作数成本

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

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

在多线程架构下运行时，在每个处理器/核心中都会发送内存操作指令，并根据需要迅速转而进行其他操作。
程序应尽早地发射load指令，在等待下一个指令需要结果之前尽可能快速完成这一操作。

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

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

Abstracting the ABI

ABI 为机器指令提供了访问数据结构和计算程序的方法，并被描述为一种低级且依赖硬件的操作规范。与此相对地，在源代码中定义这种访问的方式则是 API（应用程序接口），通常被设计为较为高级且独立于硬件水平，并且多用于人机交互以确保操作的一致性和可理解性。其一常用特征即是调用约定，在此框架下明确指定了如何将数据传递给计算程序作为输入以及从其获取结果的方式。

PTX 未透露特定调用约定、堆栈结构以及应用程序二进制接口（ABI）的具体细节内容；相反地，则提供了一个更高层次的抽象机制，并支持多种ABIs的兼容。

函数声明和定义

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

该函数被定义为包含一个可选返回值集合（Function Declaration）：包含函数名、输入参数集合；这些共同指定了该功能的接口或称为原型。

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

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

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

使用 ABI 在时，在 .reg 状态的空间参数尺寸上必须达到至少 32 的要求。源语言中的子字标量对象需提升为 PTX 中的 32-bit 寄存器，并采用以下所述的方式设置 .param 状态空间字节数组。

类似于C语言中的struct或union类型的数据成员，在PTX架构中展平为存储到寄存器或字节数组区域中，并且采用(param段的空间来表示内存区域)

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

对于调用者：

该方法既用于收集将传递给调用函数的参数值，并用于接收调用函数返回的结果值；常见地采用一个.param字节数组来整合并组织这些按值传递的数据字段。

对于被调用者：

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

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

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

对于被调用者：

输入/输出参数仅限于.param或.reg变量。
.param参数必须对齐为1/2/4/8/16字节的倍数。
.reg参数其大小必须至少为32位。
.reg状态空间用于接收/返回基类型标量值和向量值，并支持非 ABI 模式下编译时的子字大小对象；同时该状态空间还提供legacy支持。

采用.reg或.param状态空间不会影响最终参数被传递至物理寄存器或栈的位置 。这些参数在物理寄存器与栈之间的映射关系主要由ABI（ Application Binary Interface）定义以及参数的排列顺序所决定，并受到大小与对齐规则的影响。

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

该版本限定形参仅可使用位.reg的状态空间，并不能接受数组作为参数。类似于在C语言中将结构体对象展开为单个对象后通过多个寄存器来进行传递或返回。该版本能够处理多个返回值的原因也是基于上述特点。

PTX2.0版本开始引入形参、时.reg或.param状态空间，并且(param字段允许数组形式)。对于目标机器sm_20及以上型号，在遵循单返回值规则的前提下(.param字节数组用于返回不适合直接存放在寄存器中的对象类型)，而目标机器sm_1x系列则继续支持多个寄存器作为函数返回结果。

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

在PTX 3.0版本之前允许将reg与local状态空间相关的变量定义为模块作用域内的元素。自PTX 3.0及其后续版本起，在采用ABI进行编译的情况下，在模块作用域内禁止对reg与local变量进行定义，并且仅限于函数作用域内应用它们。当未采用ABI模式进行编译时，则仍可如同以往般对reg与local变量给予支持。对于那些包含来自先前于PTX 3.0之前的代码中模块作用域中的reg或local变量的情况而言，在进行编译时编译器会默认跳过ABI的支持。

可变参数函数

PTX 6.0版本能够允许无尺寸定义的数组参数通过用于处理可变数量的参数的方式被传递给相关函数。

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

PTX支持在运行时动态管理每个线程本地内存堆栈的空间分配。当alloca指令执行完毕并成功分配栈内存后，在返回的指针指引下可以通过ld.local和st.local指令完成对这些栈内存区域的操作。

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

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

内存一致性模型

多线程执行时，在每个线程中内存操作所引发的潜在影响会对其他所有线程产生可见影响，并且这些操作是按不同的顺序部分进行的。这表明任何两个操作都可能无序或以不同的顺序发生。内存一致性模型所建立的原则明确禁止不同线程所观察到的顺序出现矛盾。

范围和模型适用性

该系统中指定的约束仅在支持所有PTXISA版本的前提下，并针对采用sm_70及以上架构的设备适用。对于基于硬件API设计的PTX程序而言。
内存一致性模型无法保证对纹理和表面数据的一致性。

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

在与主机 CPU 的通信过程中，在某些系统中可能存在无法以原子方式完成具有全局范围的 $64$ 位强操作的情况

内存操作

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

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

在PTX的状态空间中每个字节相对于被访问的相同状态空间的所有线程都具有一个唯一的地址

每一个PTX内存指令包含了地址与数据类型两部分的信息。这两个因素共同决定了这个内存的位置。这个范围因此延伸至数据类型所占据的字节数。每个这样的指令还规定了相应的操作内容。

重叠 overlap

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

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

向量数据类型

该模型将通过在内存位置上的操作与标量数据类型建立关联。其最大存储容量及对齐策略定为64位。使用一组等效的标量操作来模拟包含向量数据类型的内存行为，其元素的具体处理顺序未做具体规定。

Packed Data-types 打包数据类型

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

在标量数据类型.f16上建立了一对等价的内存操作模型，并将其应用于打包数据类型的每个元素上，在这种情况下操作执行的顺序是不确定的

初始化

每个内存字节均基于启动程序中所有线程在前进行假设性操作而被W₀初始化。若该字节包含于程序变量之中且具有初始值，则W₀会将其相应地赋初值；否则默认情况下认为W₀已将固定但未知的值赋给该字节。

状态空间

内存一致性模型中基于关系的独立性并不依赖于状态空间的构建。特别地，在所有状态下阻止各种内存操作以保持一致性的机制更加严格

但是在一个状态空间中内存操作所引发的影响仅限于能够访问相同状态空间的操作，并且这些影响可以直接被察觉出来。在作用域之外的情况下这种现象会进一步限制内存操作的同步效果

操作类型

作用域

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

请注意以下要点：首先,warp并非属于作用域,而是与之有所区别；其次,CTA（即线程block）作为符合内存一致性模型所定义的作用域,其本质是一个最小化的线性集合

Morally strong operations

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

操作关联于程序执行顺序的相关（例如, 都由同一线程完成）或者每个操作均具有强原子性, 并指定了该操作域（其中包含执行其他操作的线程）；
两个均为内存操作, 并且完美重叠

内存一致性架构中的大部分核心原则基于具有道德强度的操作之间的影响关系

冲突和数据竞争

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

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

混合大小数据竞争的限制

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

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

内存一致性模型的公理无法应用于PTX程序中包含多个混合大小数据竞争的情形；但这些公理能够完整描述仅限于一致大小数据竞争的PTX程序的行为。

混合大小RMW操作的原子性

无论是否存在混合大小数据竞争的情况，在所有涉及重叠且相互作用（overlap and mutually exclusive）的关系中的程序中，在这种情况下对于所有重叠且相互作用（overlap and mutually exclusive）的关系中的每一对原子操作（atom operation）A1和A2来说，在这种情况下每个操作都明确包含了另一个的操作域范围：即在启动另一方的操作（operation）之前完成自己的全部行为（behavior）。特别地，在这种情况下对于每一对这样的原子操作来说，在其启动另一方的操作之前完成自己的全部行为（behavior）。

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

这个是干嘛的？？

Release and Acquire patterns

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

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

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

对于释放-获取模式的影响范围基于因果顺序的传递特性得以延伸至其他线程的操作流程中。

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

M上的释放操作

复制代码

    st.release [M];

    或者
    atom.acq_rel [M];
    
    
         
         
         
    代码解读

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

复制代码

    st.release [M];

    st.relaxed [M];
    
    
         
         
    代码解读

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

复制代码

    fence;

    st.relaxed [M];
    
    
         
         
    代码解读

任何基于释放锁定机制实现的内存同步仅作用于在该模式执行的第一条指令之前按照程序执行次序发生的操作

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

M上的acquire操作

复制代码

    ld.acquire [M];

    或者
    atom.acq_rel [M];
    
    
         
         
         
    代码解读

在M位置上执行的强读随后进行的是一个获取（acquire）操作，在程序流程中依次安排。

复制代码

    ld.relaxed [M];

    ld.acquire [M];
    
    
         
         
    代码解读

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

复制代码

    ld.relaxed [M];

    fence;
    
    
         
         
    代码解读

Each memory synchronization established when using the acquire mode only affects operations that occur after the last instruction of this mode in the program sequence.

Ordering of memory operations 内存操作的顺序

每个线程执行的操作序列被抽象为程序顺序（program order），而跨线程的内存同步行为被视为因果关系（causality order）。内存操作对其他内存操作产生的副作用的影响范围被定义为空间依赖关系（communication order）。该模型旨在构建一个协调系统，在该系统中允许存在通信依赖关系的同时明确区分并禁止混合使用因果关系与程序流程之间的矛盾。

The memory consistency model specifies contradictions that are not permitted between communication order for one, and causality order and program order for the other.

这句话啥意思？？？

程序顺序 program order

程序顺序与其在相应PTX源中执行指令的顺序之间建立了紧密的关联关系。这种关联表现为一种传递性质，在每个线程的操作上构成全序集合，并且这些操作仅限于同一线程内部。

观察顺序 Observation order

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

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

R and W are morally powerful, and R reads from the value written by W.
Alternatively, for some atomic operations Z, in observation order, W occurs before Z, which in turn occurs before R.

Fence-SC 顺序

Fence-SC序列是由运行时因素决定的一个无环偏序关系，并与每一对具有道德强度高的fence-sc操作相关联

内存同步

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

当X在其Fence-SC序列中领先于Y时，则执行一个fence-sc操作的同时也执行另一个fence-sc的操作。
在同一个屏障下进行的bar-sync、bar-red或bar$arrive的动作与其在该屏障上同时进行的bar-sync或bar-red的动作相互协调。
当一个write动作在其观察序列中领先于read动作时，并且该write动作的第一步和read动作的最后一部均为moralistically强的动作时，则释放模式为该write作用而acquire模式为该read作用。

API同步

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

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

因果顺序 Causality Order

通过同步机制描述内存操作如何使得在多线程环境中能够被观察到或识别。基于公理的因果顺序对一组写入操作施加限制，并允许从这些区域中进行读取。

在因果关系顺序中所涉及的关系主要由其基本构成，在基于预先设定的传递机制下形成

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

Base causality order 基本因果关系顺序

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

X与Y同步运行，
对某些特定的Z操作而言，
1. 在程序顺序上,X领先于Z,而基于基本因果关系,Z又领先于Y,
2. 基于基本因果关系,X领先于Z,与此同时,在程序顺序上,Z又领先于Y,
3. 基于基本因果关系,X领先于Z,并且基于同样的逻辑,Z又遵循程序顺序领先于Y

因果关系

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

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

在基于基本因果关系的排序中,X位于Y之前

一致性顺序 Coherence Order

确定由运行时决定的与重叠写入操作相关联的实体传输路径，并定义为一致传输顺序

当两个相互作用的写操作在moral上具有强关联或存在因果关系时，则它们具有具有一致性关系。

如果两个共享资源处于资源争用中，则它们在一致性顺序上就无关了；这从而影响了该一致性顺序的一些特性。

无法直接评估一致性顺序, 因为它完全由写操作构成. 可以通过限制读取操作仅能访问其候选写入集的方式来间接推断出这一特性.

communication order

通信过程由运行时确定且无传递性，在此过程中执行的操作会与那些在内存中进行但可能同时进行的操作相关联。

A 在通信顺序上早于一个被W写的重叠读R；当且仅当R返回值是由W所写的任意字节。
A 在一致性方式下早于另一个被W'写的动作；前提是W在一致性方式下早于W'。
A 在读取R时，在通信顺序上早于一个被W写的动作；条件是对于所有同时访问过R和W的字节，在这种情况下R返回值是由被一致性的保证所确定的一个特定前驱动作所决定的。

在通信顺序下记录内存操作间的可达性，在此情况下，则称若一个内存操作X在通信顺序下先于另一个内存操作Y执行，则称X对Y具有可达性。

公理Axioms

一致性 coherence

当执行者W依照因果顺序先于另一执行者W'时，则该执行者必须遵循一致顺序

Fence-SC

按照逻辑关系,Fence-SC序列必须与因果序列保持一致.考虑任意两个具有道德强度的fence-sc操作项 F1 和 F2,若操作项 F1依照因果先后关系位于操作项 F2之前,则该系统必须确保操作项 F1 的 Fence-SC序列同样位于 F2之前.

原子性

单拷贝（Single-Copy）原子性

冲突的操作如果是strongly进行，并采用单拷贝原子性执行。当读操作R和写操作W均为strongly进行时，在它们都被访问过的字节集合上，在同一个执行单位内通信不能同时发生：

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

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

When atomic operations A and write operation W are concurrent and strongly consistent, for the data bytes accessed by both A and W, the next two communications cannot be performed in the same transactional scope within a single invocation of a transactional point (TP).

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

No Thin Air

值可能不会凭空产生：执行无法非推测性地生成值，并因此导致基于指令依赖链及线程间通信的自足性假设失效。这也符合程序员对预期与实际行为的一致认知，在形式分析过程中则需对此进行明确陈述

Litmus Test: Load Buffering

每个位置的顺序一致性

在任何一组配对具有道德强度的重叠内存操作中, 通信流程必须与程序流程保持同步. 具体而言, 重叠操作以及其与通信流程相关的道德强度联系中的一系列程序流程必须避免形成循环. 此外, 这种机制确保了配对具有道德强度的所有进程块均保持严格的时序一致性.

Litmus Test: CoRR

因果关系

通信顺序的关联不能违背因果顺序。该读操作可读的候选写操作集合被这一限制所限定。

当读操作R在时序上早于共享的写操作W时，则该读操作无法基于一致性模型基于该写操作获取数据。
当 write 操作 W 在时序上早于 overlapping 的 read 操作，则对于任意被这两个 operation 访问的数据块：这些数据块无法基于一致性模型基于早于该 write 的另一个潜在修改获取其值。

Litmus Test: Message Passing

Litmus Test: Store Buffering

全部评论 (0)

还没有任何评论哟~

PTX ISA 7.4 参考手册翻译

文章目录 PTXParallelThreadExecutionISA7.4 Synatax SourceFormat Comments Statements instruction identifie...

【参考手册】IMX6ULL参考手册第28章GPIO翻译

第28章GPIO 1\.概述Overview GPIO通用输入/输出外设提供专用通用引脚，可配置为输入或输出。当配置为输出时，可以写入内部寄存器以控制输出引脚上驱动的状态。当配置为输入时，可以通过读...

CUBRID 8.4.3用户参考手册（翻译中）

http://www.cubrid.org/manual/843/en/ 注释解释器是一种SQLstyle标准化查询语言风格的方法；SQLstyle注释以双破折号（）开始，双破折号后的部分被视为注释...

GNU ARM 汇编简明参考手册-个人翻译

GNUARMAssemblerQuickReference GNUARM汇编简明参考手册 AsummaryofusefulcommandsandexpressionsfortheARMarchitec...

WaveForms™ SDK 参考手册(翻译笔记与总结)

本 WaveForms SDK 手册描述了。

ADRV9008、ADRV9009硬件参考手册-UG-1295的翻译过程

最近缺钱缺的凶，接了一个私活，ADRV9008、ADRV9009硬件参考手册UG1295的翻译，兄dei，真的不是我不翻译pcie6.0，真的是别人给的太多了，那就先翻译这个，pcie往后稍一稍。

[译] Dojo 工具包参考手册

说明：此文翻译并非完全按照原文进行逐字逐句翻译。期间会删除一些笔者认为无须介绍的部分，或加上一些需要说明的部分。原文介绍可以参考： <http://dojotoolkit.org/reference...

privoxy3.0.17手册翻译

随便翻译，纯粹为了日后自己看来方便…… \分割线 Privoxy是一个非缓存的代理软件，有着强大的过滤功能，可以保护个人隐私、修改页面数据、http头、控制访问、过滤广告以及其他的网络垃圾。Priv...

芝加哥风格手册：参考文献格式（中文翻译版）

基于芝加哥风格手册，第17版，2017，部分翻译；完整版原文可在GALILEO上找到（推荐查阅完整版），可见： Chapter14Contentschicagomanualofstyle.org注释和...

linux手册翻译 gdb

在linux系统中,通过mangdb可以查看gdb使用手册,手册是英文的,闲来无事,翻译一下,就当是学习了. 名称 gdbGNU调试器使用简介 gdb[OPTIONS][progprogprocID...

是否确定退出登录?

PTX ISA 7.4 参考手册翻译

文章目录

PTX Parallel Thread Execution ISA 7.4

Synatax

Source Format

Comments

Statements

instruction

identifiers

Integer Constant

Float-Point Constant

Constant expression

整型常量表达式求值

State Spaces, Types, and Variables

状态空间

Kernel Function Parameters

Kernel Parameter Attribute: .ptr

类型

基本类型

子字大小的限制使用

Texture Sampler and Surface Types

unified mode

independent mode

Channel Data Type and Channel Order Fields

变量

变量定义

向量

数组定义

初始化

对齐

参数化变量名

变量属性

变量属性指令: .attribute

指令操作数

操作数类型信息

源操作数

目的操作数

使用地址，数组和向量

地址做操作数

通用地址

数组做操作数

向量做操作数

标签和函数名做操作数

类型转换

标量转换

舍入修饰符

操作数成本

Abstracting the ABI

函数声明和定义

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

可变参数函数

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

内存一致性模型

范围和模型适用性

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

内存操作

重叠 overlap

向量数据类型

Packed Data-types 打包数据类型

初始化

状态空间

操作类型

作用域

Morally strong operations

冲突和数据竞争

混合大小数据竞争的限制

混合大小RMW操作的原子性

Release and Acquire patterns

Ordering of memory operations 内存操作的顺序

程序顺序 program order

观察顺序 Observation order

Fence-SC 顺序

内存同步

API同步

因果顺序 Causality Order

Base causality order 基本因果关系顺序

因果关系

一致性顺序 Coherence Order

communication order