PTX ISA 9.4：6. 指令操作数

6.1. 操作数类型信息#

指令中的所有操作数都从其声明中已知类型。每个操作数类型必须与由指令模板和指令类型所确定的类型相兼容。类型之间不存在自动转换。

位宽类型与所有具有相同位数的类型相兼容。具有相同位宽的整数类型彼此兼容。类型与指令类型不同但与之兼容的操作数会被静默地转换为指令类型。

6.2. 源操作数#

源操作数在指令说明中用名称 a、b 和 c 表示。PTX 描述的是一种装载-存储（load-store）机器，因此 ALU 指令的操作数必须全部位于 .reg 寄存器状态空间中声明的变量里。对于大多数运算，操作数的大小必须保持一致。

cvt（转换）指令接受多种操作数类型和大小，因为它的任务是将几乎任何数据类型转换为任何其他数据类型（以及大小）。

ld、st、mov 和 cvt 指令将数据从一个位置复制到另一个位置。 ld 和 st 指令在可寻址状态空间与寄存器之间搬移数据。mov 指令在寄存器之间复制数据。

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

6.3. 目标操作数#

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

6.4. 使用地址、数组与向量#

将标量变量用作操作数十分直接。有趣的功能从地址、数组和向量开始。

6.4.1. 作为操作数的地址#

所有访存指令都接受一个指明所访问内存位置的地址操作数。该可寻址操作数是以下之一：

[var]

可寻址变量 var 的名称。

[reg]

包含字节地址的整数类型或位宽类型寄存器 reg。

[reg+immOff]

包含字节地址的寄存器 reg 加上一个常量整数字节偏移（有符号，32 位）之和。

[var+immOff]

包含字节地址的可寻址变量 var 的地址加上一个常量整数字节偏移（有符号，32 位） 之和。

[immAddr]

一个立即绝对字节地址（无符号，32 位）。

var[immOff]

如 作为操作数的数组 所述的一个数组元素。

包含地址的寄存器可以声明为位宽类型或整数类型。

访存指令的访问大小（access size）是在内存中访问的字节总数。例如，ld.v4.b32 的访问大小为 16 字节，而 atom.f16x2 的访问大小为 4 字节。

地址必须按访问大小的倍数自然对齐。如果地址未正确对齐，则产生的行为是未定义的。例如，除其他情况外，访问可能会通过静默地屏蔽掉低位地址比特来实现正确的舍入而继续执行，或者该指令可能发生错误（fault）。

地址大小可以是 32 位或 64 位。不支持 128 位地址。地址会按需零扩展到指定宽度，如果寄存器宽度超过目标架构的状态空间地址宽度，则会被截断。

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

mov 指令可用于将变量的地址搬移到指针中。该地址是变量所处声明状态空间中的一个偏移。 访存指令所指定的状态空间必须与地址操作数的状态空间相匹配，否则其行为是未定义的。装载和存储操作在寄存器与可寻址状态空间中的位置之间搬移数据。 其语法与许多汇编语言中使用的类似：标量变量只需直接命名，而地址则通过将地址表达式用方括号括起来进行解引用。地址表达式包括变量名、 地址寄存器、地址寄存器加字节偏移，以及在编译时求值为常量地址的立即地址表达式。

以下是几个示例：

.shared .u16 x;
.reg    .u16 r0;
.global .v4 .f32 V;
.reg    .v4 .f32 W;
.const  .s32 tbl[256];
.reg    .b32 p;
.reg    .s32 q;
ld.shared.u16   r0,[x];
ld.global.v4.f32 W, [V];
ld.const.s32    q, [tbl+12];
mov.u32         p, tbl;

6.4.1.1. 通用寻址#

如果某条访存指令未指定状态空间，则该操作使用通用寻址（generic addressing）来执行。状态空间 .const、Kernel 函数参数 （.param）、.local 和 .shared 被建模为 通用地址空间内的窗口。每个窗口由一个窗口基址和一个窗口大小定义，其窗口大小等于相应状态空间的大小。通用地址映射到 global 内存，除非它落在 const、local 或 shared 内存的窗口之内。 Kernel 函数参数（.param）窗口包含 在 .global 窗口之内。在每个窗口内，通用地址通过将窗口基址从通用地址中减去来映射到底层状态空间中的一个地址。

6.4.2. 作为操作数的数组#

可以声明所有类型的数组，并且标识符在数组声明所在的空间中成为一个地址常量。数组的大小在程序中是一个常量。

数组元素既可以使用显式计算出的字节地址来访问，也可以使用方括号记法通过索引数组来访问。方括号内的表达式要么是常量整数、寄存器变量，要么是一个简单的 带常量偏移的寄存器 表达式，其中偏移是一个常量表达式，被加到或减去一个寄存器变量上。如果需要更复杂的索引，则必须在使用之前将其写成一个地址计算。示例如下：

ld.global.u32  s, a[0];
ld.global.u32  s, a[N-1];
mov.u32        s, a[1];  // move address of a[1] into s

6.4.3. 作为操作数的向量#

向量操作数可以被指定为指令的源操作数和目标操作数。然而，当它被指定为目标操作数时，向量表达式中的所有元素必须唯一，否则行为未定义。 向量也可以作为参数传递给被调用的函数。

可以使用后缀 .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

6.4.4. 作为操作数的标号与函数名#

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

自 PTX ISA 3.1 版本起，mov 指令可用于取内核函数（kernel functions）的地址，以传递给一个从 GPU 发起内核启动的系统调用。此特性是对 CUDA 动态并行（CUDA Dynamic Parallelism） 支持的一部分。详情请参阅 CUDA Dynamic Parallelism Programming Guide。

6.5. 类型转换#

所有算术、逻辑和数据搬移指令的全部操作数必须具有相同的类型和大小，除非那些改变大小和/或类型本身就是指令定义一部分的运算。大小或类型不同的操作数必须在运算之前进行转换。

6.5.1. 标量转换#

表 15 和 表 16 展示了 在给定不同类型的操作数时 cvt 指令所使用的精度和格式。例如，如果一条 cvt.s32.u16 指令获得一个 u16 源操作数以及作为目标操作数的 s32， 那么 u16 会被零扩展为 s32。

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

注释

sext = 符号扩展；zext = 零扩展；chop = 仅保留能够容纳的低位；

s2f = 有符号数转浮点数；f2s = 浮点数转有符号数；u2f = 无符号数转浮点数；

f2u = 浮点数转无符号数；f2f = 浮点数转浮点数。

¹ 如果目标寄存器比目标格式更宽，则在截断之后，结果会被扩展到 目标寄存器的宽度。扩展的类型（符号扩展或零扩展）基于 目标格式。例如，针对 32 位寄存器的 cvt.s16.u32 会先截断为 16 位，然后 符号扩展为 32 位。

6.5.2. 舍入修饰符#

转换指令可以指定舍入修饰符。在 PTX 中，有四种整数舍入 修饰符和六种浮点舍入 修饰符。表 17 和 表 18 对舍入修饰符进行了汇总。

6.6. 操作数开销#

来自不同状态空间的操作数会影响操作的速度。寄存器最快，而 全局内存最慢。访问内存的大部分延迟可以通过多种方式隐藏。第一种 方式是拥有多个执行线程，这样硬件可以发出一个内存操作，然后 切换到其他执行。隐藏延迟的另一种方式是尽早发出加载指令， 因为在期望的结果被后续（时间上靠后）的指令使用之前，执行不会被阻塞。存储操作 中的寄存器可用得快得多。表 19 给出了 使用不同类型内存的开销估计。