PTX-ISA
官方文档 Parallel Thread Execution ISA 的简单中译
其中PTX版本为8.0
文档正在不断补充，有以下代办事项
Chapter 1. Introduction
1.1.. Scalable Data-Parallel Computing using GPUS
PTX 定义了一套抽象设备层面的 ISA 用于通用的并行编程指令。让开发人员可以忽略掉具体的目标设备指令集差异，进行通用的开发。
1.2. Goals of PTX
- 提供了一套跨越多 GPU 架构的稳定 ISA 。
- 可以提供近似 native 的性能。
- 为 C/C++ 和其他编译器提供与目标设备架构无关的 ISA。
- 为应用和中间件开发人员提供了易用的 ISA。
- 为优化代码的生成器和转换器提供了通用 ISA。
- 简化库、性能内核和体系结构测试的码量。
- 提供了可扩展的编程模型，涵盖多种架构的GPU。
1.3. PTX ISA Version 8.0
8.0版本有如下新特性：
- 添加对目标
sm_90a的支持。 - 添加对异步
warpgroup级矩阵乘法和累加的支持操作wgmma的支持。 - 通过对大数据进行批量操作来扩展异步复制操作。
- 引入压缩整数类型
.u16x2和.s16x2。 - 扩展整数算术指令
add，允许压缩整数类型.u16x2和.s16x2。 - 扩展整数算术指令
min和max，允许压缩整数类型.u16x2和.s16x2. - 添加特殊寄存器
%current_graph_exec。 - 添加对
elect.sync指令的支持。 - 添加对函数和变量的
.unified属性的支持。 - 添加对
setmaxnreg指令的支持。 - 在
barrier.cluster指令中添加对.sem限定符的支持。 - 扩展
fence指令以允许使用op_restrict进行特定于操作码的同步。 - 添加
mbarrier.arrive、mbarrier.arrive_drop、mbarrier.test_wait和mbarrier.try_wait操作。 - 添加对
mbarrier对象的事务计数操作的支持，指定为.expect_tx和.complete_tx限定符。
Chapter 2. Programming Model
2.1. A Highly Multithreaded Coprocessor
GPU 是可以并行执行大量线程的设备，作为主 CPU 或主机的协处理器运行。
应用程序的一部分被执行多次，但独立在不同的数据上——可以隔离成一个内核函数，在 GPU 上执行尽可能多的不同线程，此类函数被编译为 PTX 指令集。
2.2. Thread Hierarchy
执行 GPU 内核函数的线程被划分为 线程网格(Grid)，而 Grid 又下划为 Cluster 与 CTA 。
2.2.1. Cooperative Thread Arrays
在PTX的概念中，CTA是一组可以相互通信的线程所组成的线程块，对应 CUDA 中的 Thread Block。
在CTA中同样有warp的概念，warp是CTA的最小执行线程集合。
2.2.2. Cluster of Coopperative Thread Arrarys
Cluster 是由多个 CTA 组成，Cluster 大小是可选的，默认是 1x1x1 的大小。
有特定的符号可以查询 CTA 的 id 等信息，存放在特殊寄存器中。
ps：目前只在sm_90或以上的硬件架构中才支持这一概念。
2.2.3. Grid of Cluster
Grid是最高的线程等级，包含了多个Cluster。
存在特定的符号可以查询 Cluster 的 id 等，存放在特殊寄存器中。
2.3. Memory Hierarchy
- global memory：可读可写，线程共享；
- constant memory：只读，cached，线程共享；
- texture：只读，cached；
- surface：可读可写，cached；
- shared memory：CTA中线程共享；
- local memory：线程独占；
Chapter 3. PTX Machine Model
3.1. A Set of SIMT Multiprocessors
GPU 硬件模型：
3.2. Independent Thread Scheduling
在Volta架构之前，一个 warp 内的32个线程因为共用一个程序计数器，通过 active mask 区别 active thread。
从Volta架构开始，支持了 warp 内的线程独立调度，每个线程都有自己独立的程序计数器。当出现一个 warp 内的线程分化的时候，允许不同的线程做不同的事情，不再阻塞。
开发者在编写Volta及以上架构的PTX代码时，需要特别留意由于独立线程调度操作引起的 向下兼容性问题。
3.3. On-chip Shared Memory
根据 Figure4 中的信息显示，每个 Multiprocessor 可以利用的片上内存主要分为以下四种：
- 每个 processor 都有一组32-bit的本地寄存器；
- 每个 processor 共享的
shared memory，其拥有并行数据缓存； - 每个 processor 可通过共享的只读 cache ，加速读取设备的指定常量存储区域
constant memory，内存有限； - 每个 processor 可通过共享的只读 cache ，加速读取设备指定的存储区域
texutre，支持多种寻址模式和数据滤波器；
需要注意的是，local memory和global memory没有专用 cache 加速。
Chapter 4. Syntax
PTX 源程序模块带有汇编语法风格的指令操作符和操作数。通过 ptxas 后端编译优化器对 PTX 源模块进行优化、编译并生成对应的二进制对象文件。
4.1. Source Format
- 源模块是以 ASCII 文本形式，以
\n进行换行。 - 所有空格将被忽略，除非在语言中被用于分格标记。
- 接受 C 风格的预处理标记，通过
#标记，如：#include, #define, #if, #ifdef, #else, #endif, #line, #file - PTX 区分大小写，关键字用小写。
- 每个 PTX 模块必须从指定 PTX 语言版本的
.version指令开始，接着是一个.target指令指定假设的目标体系结构。
4.2. Comments
PTX 的注释服从 C\C++ 风格，使用 /* 注释内容 */ 或//均可。
4.3. Statements
PTX 语句既包含预处理 (directive) 也包含指令 (instruction)，以可选的指令标记开头并以分号结尾。
.reg .b32 r1, r2;
.global .f32 array[N];
start: mov.b32 r1, %tid.x;
shl.b32 r1, r1, 2; // shift thread id by 2 bits
ld.global.b32 r2, array[r1]; // thread[tid] gets array[tid]
add.f32 r2, r2, 0.5; // add 1/2
4.3.1. Directive Statements
PTX 中支持的编译器指示如下：
4.3.2. Instruction Statements
- 指令由一个指令操作码和由逗号分隔的零个或多个操作数组成，并以分号结束。操作数可以是寄存器变量，常量表达式、地址表达式或指令标签名称。
- 指令有一个可选的判断条件作控制流的跳转。判断条件在可选的指令标记后面，在操作码前面，并被写成
@p，其中p是一个条件寄存器。判断条件可以取非，写成@!p。 - 指令标记之后的字段，首先是目标操作数，后续是源操作数。
指令关键字：
4.4. Identifiers
- 用户定义的标识符，服从 C++ 的规则，字母或者下划线开头，或者以
$开头。 - PTX 没有指定标识符的最大长度，并表示所有实现至少支持 1024 个字符。
- PTX支持以
%为前缀的变量，用于避免命名冲突。 - PTX以
%为前缀预定义了一个常量和一小部分特殊寄存器，如下表所示：
4.5. Constants
PTX支 持整型和浮点常量和常量表达式。这些常数可用于数据初始化和作为指令的操作数。对于整型、浮点和位大小类型检查规则是相同的。
对于判断类型的数据和指令，允许使用整型常量，即0为False和!0为True。
4.5.1. Integer Constants
整型常量的大小为64位，有符号或无符号，即每个整数常量的类型为.s64或.u64。
在指令或数据初始化中使用时，每个整整型常量会根据使用时的数据或指令类型转换为适当的大小。
整型常量可以写作十六进制、十进制、八进制、二进制，写法同C语言一直，最后加U表示unsigned:
十六进制： 0[xX]{hexdigit}+(U)
十进制： {nonzero-digit}{digit}+(U)
八进制： 0{octal digit}+(U)
二进制： 0[bB]{bit}+(U)
4.5.2. Floating-Point Constants
浮点常量表示为 64 位双精度值，所有浮点常量表达式都使用 64 位双精度算术求值。
需要注意的是如果用十六进制表示，是表示32位单精度浮点。并且可能不会被用在常量表达式中。
浮点数的值：
- 第一种表示，可以用一个可选的小数点和带符号的指数进行表达(1.34e-2)。但和 C\C++ 不同的是，在 PTX 里面不能通过后缀来区分浮点数的的类型，比如：1.0f。
- 第二种表示，可以使用十六进制进行表示，如下：
0[fF]{hexdigit}{8} // single-precision floating point
0[dD]{hexdigit}{16} // double-precision floating point
mov.f32 $f3, 0F3f800000; // 1.0, 表示：$f3 = 1.0;
4.5.3. Predicate Constants
整型常量也可以作为判断数据，0表示False，!0表示True。
4.5.4. Constant Expressions
在 PTX 中，常量表达式是使用 C 中的操作符形成的，并使用与C中类似的规则求值，但通过限制类型和大小、删除大多数强制转换和定义完整语义来简化，以消除 C 中表达式求值依赖于实现的情况。(减去编译器推导数据类型等的负担)
常量表达不支持从整型到浮点数的类型转换。
常量表达式中的优先级顺序从上到下如小表所示，第一行执行优先级最高，同一行的优先级相同，对于多个一元操作求值的话是从右向左的顺序，而二元操作是从左向右：
4.5.5. Integer Constant Expression Evaluation
整型常量表达式，在编译时有一套规则进行推导。这些规则基于 C 中的规则，但它们已被简化为只适用于 64 位整数，并且在所有情况下都完全定义了行为（即不会有二义性的表达式）
- 默认整型常数是
signed除非需要转换为unsigned防止溢出，或者手动添加U后缀如：
42, 0x1234, 0123 are signed.
0xfabc123400000000, 42U, 0x1234U are unsigned
- 一元加减符保留输入操作数的类型，如：
+123, -1, -(-42) are signed.
-1U, -0xfabc123400000000 are unsigned.
- 一元操作中的取非
!操作会产生带符号的0或1。 - 位操作中的取反操作
~默认将源操作数是为unsigned，结果也为unsigned。 - 一些二元操作需要规范化源操作数，如果其中有一个是
unsigned，那么需要将两个源操作数都转换为unsigned进行计算，称为常用算数转换。 - 加减乘除执行计算之后，结果与源操作数的数据类型保持一致，即，有一个为
unsigned则结果也为unsigned，反之则为signed。 - 取余
%的操作会将操作数解释为unsigned，与C不同，C允许负除数。但属于实现定义行为 - 移位操作的第二个源操作数解释为
unsigned，结果数据类型与第一个源操作数一致。如果是signed右移则为算术右移，unsigned为逻辑右移。 - 位与
&，位或|，位异或^操作也服从常用数据转换规则。 - 与
&&，或||，等于==，不等!=操作产生signed结果，值为0或1。 - 大小比较运算符(
<、>、<=、>=)对于源操作数符服从常用转换规则，产生signed结果，值为0或1。 - 可使用
(.s64)或(.u64)将表达式转换为signed或unsigned。 - 对于三元判断符
?:，第一个源操作数必须是整型，但第二个和第三个可以是整型或者浮点型，其结果类型与选择的操作数类型一致。
4.5.6. Summary of Constant Expression Evaluation Rules
下表总结了常量表达式的推导规则：
Chapter 5. State Spaces, Types, and Variables
虽然特殊的资源在不同架构的GPU上可能是不同的，但资源种类是通用的，这些资源通过状态空间和数据类型在 PTX 中被抽象出来。
5.1. State Spaces
状态空间是具有特定特征的存储区域。所有变量都驻留在某个状态空间中。状态空间的特征包括其大小、可寻址性、访问速度、访问权限和线程之间的共享级别。
不同的状态空间如下：
不同状态空间的性质如下：
5.1.1. Register State Space
.reg寄存器读写速度很快，但是数量有限制，并且不同架构的寄存器数量不一样。当寄存器使用超标时，会溢出到内存中，影响读写速度。
寄存器可以是有类型的，也可以是无类型的，但是寄存器大小是被严格限制的，除了 1-bit 的判断符(bool) 寄存器以外，还有宽度为 8-bit\16-bit\32-bit\64-bit 的标量寄存器，以及 16-bit\32-bit\64-bit\128-bit 的矢量寄存器。
8-bit 寄存器最常见用途是和ld、st和cvt指令一起使用，或作为向量组的元素。
寄存器与其他状态空间的区别在于，它们不是完全可寻址的，也就是说，不可能引用寄存器的地址。(可以理解为仅在作用域内有效，即寄存器是栈上存储)
寄存器对于多字的读写可能会需要做边界对齐。
5.1.2. Special Register State Space
.sreg特殊寄存器是预定义的平台特殊寄存器，所有的特殊寄存器都是预定义的，如grid、cluster等相关参数。
5.1.3. Constant State Space
.const常量状态空间是由 host 端初始化的只读内存，通常使用ld.const进行访问，目前常量内存的限制为 64KB。
另外还有一个 640KB 的常量内存，被划分为 10 个 64KB 的区域，驱动程序可以在这些区域上进行初始化数据分配，并通过指针的形式作为 kernel 参数传入。由于这 10 个常量内存区域并不连续，所以驱动程序在分配的时候应该保证每一块常量内存不得超过64KB，不得越界。
静态大小的常量变量有一个可选的变量初始化器。默认情况下，没有显式初始化式的常数变量被初始化为零。驱动程序分配的常量缓冲区由 host 初始化，并将指向这块常量内存的指针作为 kernel 参数传入。
5.1.4. Global State Space
.global全局状态空间是能够被kernel中所有线程都访问到的内存空间，使用ld.global、st.global和atom.global指令访问全局内存。
没有显示初始化的全局变量默认初始化为0。
5.1.5. Local State Space
.local本地状态空间是每个线程私有的内存空间。通常是带缓存的标准内存。其有大小限制，因为必须按每一个线程进行分配。
使用ld.local、st.local进行本地变量的访问。
在编译的ABI的时候，必须将.local声明在函数作用域内，并且内存申请在栈上。
在不支持堆栈的实现中，所有本地内存变量都存储在固定地址中，不支持递归函数调用，并且.local变量可能在模块(module)作用域声明。
在PTX 3.0及以下，module-scope .local将默认被禁用。
5.1.6. Parameter State Space
.param参数状态空间主要用于以下情况：
- 作为从host传入kernel的输入参数；
- 在kernel执行过程中，为调用的device函数声明正式的输入和返回参数；
- 通常可用于声明局部作用域的字节矩阵，主要通过值传递大型的结构体。
kernel 函数参数与 device 函数参数是不同的，一个是内存的访问与共享权限不同 (read-only 对比 read-write，per-kernel 对比 per-thread)。
5.1.6.1. Kernel Function Parameters
每个内核函数定义都包含一个可选的参数列表。这些参数是在.param状态空间中声明的可寻址只读变量。通过使用ld.param指令访问内核参数值。内核参数变量被grid内的所有线程共享。
内核参数的地址可以使用mov指令移动到寄存器中。结果地址在.param状态空间中，可以使用ld.param指令访问。
.entry foo ( .param .b32 N, .param .align 8 .b8 buffer[64] )
{
.reg .u32 %n;
.reg .f64 %d;
ld.param.u32 %n, [N];
ld.param.f64 %d, [buffer];
...
.entry bar ( .param .b32 len )
{
.reg .u32 %ptr, %n;
mov.u32 %ptr, len; // 寄存器%ptr指向len变量的地址
ld.param.u32 %n, [%ptr]; //寄存器%n读取%ptr指针指向的值
5.1.6.2. Kernel Function Parameter Attributes
kernel 函数参数可以用可选的.ptr属性声明，可用来指示参数是指向内存的指针，也可表明指针所指向内存的状态空间和对齐方式。
5.1.6.3. Kernel Parameter Attribute: .ptr
.ptr语法：
.param .type .ptr .space .align N varname
.param .type .ptr .align N varname
.space = { .const, .global, .local, .shared };
其中.space和.align是可选的属性，.space缺失则默认是.const, .global, .local, .shared中的一种（基本属于未定义，所以一般还是不建议省略），.align缺失则默认按照 4 byte 对齐。
.entry foo ( .param .u32 param1,
.param .u32 .ptr.global.align 16 param2,
.param .u32 .ptr.const.align 8 param3,
.param .u32 .ptr.align 16 param4 // generic address
// pointer
) { .. }
5.1.6.4. Device Function Parameters
从 PTX2.0 开始扩展了 device 参数空间的使用，最常见的用法是不按照寄存器大小传值，如传入 8 bytes 大小的结构体参数。
// pass object of type struct { double d; int y; };
.func foo ( .reg .b32 N, .param .align 8 .b8 buffer[12] )
{
.reg .f64 %d;
.reg .s32 %y;
ld.param.f64 %d, [buffer];
ld.param.s32 %y, [buffer+8];
...
}
// code snippet from the caller
// struct { double d; int y; } mystruct; is flattened, passed to foo
...
.reg .f64 dbl;
.reg .s32 x;
.param .align 8 .b8 mystruct;
...
st.param.f64 [mystruct+0], dbl;
st.param.s32 [mystruct+8], x;
call fooo, (4, mystruct);
...
函数的输入参数可以使用ld.param进行读，返回值可以使用st.param进行写。
但是写input参数和读返回值都是不合法的。
除了按值传递结构外，当形式形参的地址在被调用的函数中被取时，还需要.param空间标注。
在 PTX 中，函数输入参数的地址可以使用mov指令移动到寄存器中。注意，如果需要，参数将被复制到堆栈中，因此地址将位于.local状态空间中，并通过ld.local和st.local指令进行访问。
不能使用mov来获取局部作用域的.param空间变量的地址。从PTX ISA 6.0版本开始，可以使用mov指令获取设备函数返回参数的地址。
// pass array of up to eight floating-point values in buffer
.func foo ( .param .b32 N, .param .b32 buffer[32] )
{
.reg .u32 %n, %r;
.reg .f32 %f;
.reg .pred %p;
ld.param.u32 %n, [N];
mov.u32 %r, buffer; // forces buffer to .local state space
Loop:
setp.eq.u32 %p, %n, 0;
@%p: bra Done;
ld.local.f32 %f, [%r];
...
add.u32 %r, %r, 4;
sub.u32 %n, %n, 1;
bra Loop;
Done:
...
}
5.1.7. Shared State Space
.shared共享内存属于执行运算的CTA并且可以被同属一个cluster中的所有CTA的线程读写。
附加的子限定符::cta或::cluster可以在使用.shared的指令中指定状态空间，指示该地址是否属于正在执行的 CTA 或 cluster 中的任何 CTA 的共享内存。(即cluster共享，还是CTA内部共享)
.shared::cta的地址窗口也属于.shared::cluster的地址窗口。如果.shared状态空间中没有指定子限定符，则默认为::cta。例如，ld.shared等价于ld.shared::cta`。
在.shared状态空间中声明的变量引用当前CTA中的内存地址。指令mapa给出了cluster中另一个CTA中对应变量的.shared::cluster地址。
共享内存通常有一些优化来支持共享。一个例子是广播，所有线程从同一个地址读取。另一种是从顺序线程的顺序访问。
5.2. Types
5.2.1. Fundamental Types
在 PTX 中，基本类型反映了目标架构支持的原生数据类型。基本类型同时指定类型和大小。
寄存器变量总是一种基本类型，指令对这些类型进行操作。
基本类型如下：
大多数指令都有一个或多个类型说明符，用于完全指定指令的行为。操作数类型和大小将根据指令类型进行检查，以确保兼容性。
位大小相同的任何基本类型之间都是兼容的。
原则上，所有基本类型(除开 predicate 类型)可以只用位大小但标明具体类型进行声明。
5.2.2. Restricted Use of Sub-Word Sizes
.u8、.s8和.b8被限制在ld、st和cvt指令中使用。
.fp16只能被用在与fp32和fp64的相互转化中，以及半精度浮点指令和纹理获取指令中。
.fp16x2只能被用在半精度浮点指令和纹理获取中。
ld、st和cvt指令允许源操作数和目标数据操作数比指令类型的大小更宽。因此可以使用规则宽度的寄存器加载、存储和转换窄的值。例如，在加载、存储或转换为其他类型和大小时，8 位或 16 位的值可能直接保存在 32 位或 64 位寄存器中。
5.2.3. Alternate Floating-Point Data Formats
PTX 中支持的基本浮点类型具有隐式的位表示，表示用于存储指数和尾数的位数。
- bf16
- e4m3
- e5m2
- ft32
替代数据格式不能用作基本类型。它们被某些指令支持为源格式或目标格式。
5.3. Texture Sampler and Surface Types
PTX 中有一些内建的不透明类型来定义texture、sampler、surface descriptor变量。
这些类型的命名字段类似于结构体，但所有的信息如：布局、字段顺序、基址和总体大小都隐藏在PTX程序中，因此称为不透明。
这些不透明类型的使用有如下限制：
- 变量定义在全局 (module) 作用域和内核参数列表中；
- module-scope 变量的静态初始化使用逗号隔开静态赋值表达式；
- texture\sampler\surface 的引用通过 texture\surface 的 load\save 指令完成
tex,suld,sust,sured。 - 通过查询指令检索指定成员的值；
- 创建指向不透明变量的指针可以使用
mov指令，如：mov.u64 reg, opaque_var。产生的指针可以从内存中读写，也可以通过参数传递给函数，还可以被 texture\surface 的读写查询指令所引用。 - 不透明变量不能出现在初始化中，如：初始化一个指针指向不透明变量。
从 PTX ISA 3.1 版本开始支持使用指向不透明变量的指针间接访问 texture\surface，需要目标架构sm_20及以上。
上述的三种内置的不透明类型是.texref、.samplerref和.surfref。
5.3.1. Texture and Surface Properties
上表中的width、height和depth表示 texture\surface 在每个维度的元素个数(更准确的说可以理解为像素pixel)。
其中每一个像素的属性可以由channel_data_type和channel_order来表示。
OpenCL中的定义是被 PTX 支持的，所以可以参考OpenCL的定义如下：
5.3.2. Sampler Properties
关于 sampler 的属性，有以下意义意义：
normalized_coords：表示坐标是否归一化为[0.0, 1.0f]。如果没有被显式设置，则会在runtime阶段根据源码进行设置(也就是说还有可能默认被设为开启？？没试过，通常是默认非归一化的)filter_mode：表示如何基于坐标值计算texture读取的像素。addr_mode_{0,1,2}： 定义了每个维度的寻址模式，该模式决定了每个维度如何处理out-of-range的坐标。force_unnormalized_coords: 在Independant Mode下独有的属性，字面意思很好理解，会去将texture中的normalized_coords强行改写为unnormalized_coords当其被设置为True时，如果是False，那么就默认使用texture中的设置。
我们在声明这些不透明类型的时候，如果位于 module-scope 中，则需要使用.global状态空间；而如果位于 kernel 参数列表中，则需要使用.param状态空间。
.global .texref my_texture_name;
.global .samplerref my_sampler_name;
.global .surfref my_surface_name;
在.global状态空间中，可以使用静态列表进行初始化，如：
.global .texref tex1;
.global .samplerref tsamp1 = { addr_mode_0 = clamp_to_border,
filter_mode = nearest
};
5.3.3. Channel Data Type and Channel Order Fields
见5.3.1 Table11 Table12
5.4. Variables
在 PTX 中，除了基本的数据类型，还支持简单的聚合数据类型，如 vector 和 array 。
5.4.1. Variable Declarations
所有的存储数据都是通过变量声明来定义的。
标量声明包含，变量所在状态空间，类型和大小，变量命。以及可选的数组大小，可选的初始化方式，可选的变量固定地址。
.global .u32 loc;
.reg .s32 i;
.const .f32 bias[] = {-1.0, 1.0};
.global .u8 bg[4] = {0, 0, 0, 0};
.reg .v4 .f32 accel;
.reg .pred p, q, r;
5.4.2. Vectors
任何长度为 2 或 4 的 non-predicat 基础类型 vector 可以通过.v2或.v4的前缀进行声明。
vector 必须是基础类型，可以被声明为寄存器，长度不能超过 128bit，只包含 3 个元素的矢量也会被创建为.v4矢量，剩余一个元素是padding位。
.global .v4 .f32 V; // a length-4 vector of floats
.shared .v2 .u16 uv; // a length-2 vector of unsigned ints
.global .v4 .b8 v; // a length-4 vector of bytes
默认情况下，矢量的大小是内存对齐的(与长度和类型大小有关)，所以在我们进行矢量读写的时候，应该保证访问的内存大小是对齐到矢量的整数倍的。
5.4.3. Array Declarations
数组的声明和 C 一样，无论是一维还是多维，并且声明的大小必须是常量表达式。
.local .u16 kernel[19][19];
.shared .u8 mailbox[128];
当数组的声明伴随着初始化表达式时，数组的第一维尺寸是可以被省略的，一维的尺寸是由舒适化表达式中的元素决定的。
.global .u32 index[] = { 0, 1, 2, 3, 4, 5, 6, 7 };
.global .s32 offset[][2] = { {-1, 0}, {0, -1}, {1, 0}, {0, 1} };
5.4.4. Initializers
变量的初始化方法与 C/C++ 是类似的。并且时支持不完整的初始化，会默认进行补 0 操作。
.const .f32 vals[8] = { 0.33, 0.25, 0.125 };
.global .s32 x[3][2] = { {1,2}, {3} };
// is equivalent to
.const .f32 vals[8] = { 0.33, 0.25, 0.125, 0.0, 0.0, 0.0, 0.0, 0.0 };
.global .s32 x[3][2] = { {1,2}, {3,0}, {0,0} };
当前，只支持const和global内存空间的变量初始化操作，如果上述两种空间中的变量没有显式的初始化，则默认初始化位0。不允许初始化外部变量。
在初始化式中出现的变量名表示变量的地址，这可用于静态初始化指向变量的指针。
初始化式支持var+offset的表达式，offset表示基于var地址的byte偏移。
PTX 提供了一个操作符generic()，用于获取变量的地址。
从 PTX ISA 7.1 版本开始，提供了一个mask()操作符，其中mask是一个整型的立即数。
mask()操作符中唯一允许的表达式是整数常量表达式和符号表达式，用于表示变量地址。可以理解为通过通过为&操作和移位操作提取出某个byte的数据，并且作为初始化数据。
支持的数有：0xFF, 0xFF00, 0xFF0000, 0xFF000000, 0xFF00000000,0xFF0000000000, 0xFF000000000000, 0xFF00000000000000
.const .u32 foo = 42;
.global .u32 bar[] = { 2, 3, 5 };
.global .u32 p1 = foo; // offset of foo in .const space
.global .u32 p2 = generic(foo); // generic address of foo
// array of generic-address pointers to elements of bar
.global .u32 parr[] = { generic(bar), generic(bar)+4,
generic(bar)+8 };
// examples using mask() operator are pruned for brevity
.global .u8 addr[] = {0xff(foo), 0xff00(foo), 0xff0000(foo), ...};
.global .u8 addr2[] = {0xff(foo+4), 0xff00(foo+4), 0xff0000(foo+4),...}
.global .u8 addr3[] = {0xff(generic(foo)), 0xff00(generic(foo)),...}
.global .u8 addr4[] = {0xff(generic(foo)+4), 0xff00(generic(foo)+4),...}
// mask() operator with integer const expression
.global .u8 addr5[] = { 0xFF(1000 + 546), 0xFF00(131187), ...};
初始化式不支持.fp16、.fp16x32和.pred，其余类型都支持。
.global .s32 n = 10;
.global .f32 blur_kernel[][3]
= {{.05,.1,.05},{.1,.4,.1},{.05,.1,.05}};
.global .u32 foo[] = { 2, 3, 5, 7, 9, 11 };
.global .u64 ptr = generic(foo); // generic address of foo[0]
.global .u64 ptr = generic(foo)+8; // generic address of foo[2]
5.4.5 Alignment
所有可寻址变量的内存字节对齐数，可以在变量生命的时候被定义，使用可选的.align关键字。
关于对齐，与 C/C++ 中的类似
// allocate array at 4-byte aligned address. Elements are bytes.
.const .align 4 .b8 bar[8] = {0,0,0,0,2,0,0,0};
所有访问内存的 PTX 指令都要求地址与访问大小的倍数对齐。内存指令的访问大小是在内存中访问的总字节数。如：ld.v4.b32的访问大小是 16bytes，而atom.fp16x2的访问大小是 4bytes。
5.4.6. Parameterized Variable Names
由于 PTX 支持虚拟寄存器，编译器前端生成大量寄存器名是很常见的。寄存器支持像数组一样的批量声明。
.reg .b32 %r<100>; // declare %r0, %r1, ..., %r99
5.4.7. Variable Attributes
变量可以用可选的.attribute指令来声明，该指令允许指定变量的特殊属性。关键字.attribute后面是在括号内的属性说明。多个属性用逗号分隔。
5.4.8. Variable and Function Attribute Driective: .attribute
.managed: 该属性指定变量将分配到统一虚拟内存中，在该内存中，系统中的host和device可以直接引用该变量。只能用于.global状态空间。
.unified：该属性指定该函数在 host 上和系统中的其他 devics 上具有相同的内存地址。只能用于.global状态空间。
.global .attribute(.managed) .s32 g;
.global .attribute(.managed) .u64 x;
.global .attribute(.unified(19,95)) .f32 f;
.func .attribute(.unified(0xAB, 0xCD)) bar() { ... }
5.5. Tensors
5.5.1. Tensor Dimension, size and format
Tensor 是内存中的多维矩阵结构，有以下属性：
- 维度
- 每个维度的大小
- 单个元素的类型
- Tensor 的维度跨度
PTX 支持对 Tensor 数据进行操作的指令：
- 在全局内存与共享内存之间复制数据
- 减少 Tensor 的数据
PTX Tensor 指令将 Tensor 存为全局内存中一个多维结构和共享内存中的线性数据。
5.5.2 .Tensor Access Modes
- Tensor 的尺寸: 1D, 2D, 3D, 4D, 5D
- Tensor 的位大小：.b32, .b64
- Tensor 的整数类型：.u8, .u16, .u32, .s32, .u64, .s64
- Tensor 的浮点数类型： .f16, .bf16, .tf32, .f32, .f64 (四舍五入到最近的偶数).
- Tensor 的访问模式：平铺模式、Im2col 模式