PTX ISA 9.4：4. 语法

PTX 程序是文本源模块（文件）的集合。PTX 源模块具有类汇编语言的语法，带有指令操作码和操作数。伪操作用于指定符号与寻址管理。ptxas 优化后端编译器对 PTX 源模块进行优化与汇编，以生成相应的二进制目标文件。

4.1. 源文件格式#

源模块为 ASCII 文本。各行以换行符（\n）分隔。

所有空白字符都是等价的；除用于分隔语言中的记号外，空白字符都会被忽略。

可以使用 C 预处理器 cpp 来处理 PTX 源模块。以 # 开头的行是预处理器指示符。以下是常见的预处理器指示符：

#include, #define, #if, #ifdef, #else, #endif, #line, #file

Harbison 与 Steele 所著的 C: A Reference Manual 对 C 预处理器作了很好的介绍。

PTX 区分大小写，关键字使用小写。

每个 PTX 模块都必须以一个 .version 指示符开头，用以指定 PTX 语言版本，其后是一个 .target 指示符，用以指定所假设的目标架构。有关这些指示符的更多信息，请参阅 PTX 模块指示符。

4.2. 注释#

PTX 中的注释遵循 C/C++ 语法，使用不可嵌套的 /* 和 */ 来表示可以跨越多行的注释，并使用 // 开始一段延伸到下一个换行符的注释，该换行符用于结束当前行。注释不能出现在字符常量、字符串字面量或其他注释内部。

PTX 中的注释按空白字符处理。

4.3. 语句#

PTX 语句要么是一个指示符，要么是一条指令。语句以一个可选标号开头，并以分号结束。

示例

        .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. 指示符语句#

指示符关键字以一个点号开头，因此不会与用户定义的标识符发生冲突。PTX 中的指示符列于 表 1 和 并在 状态空间、类型与变量 和 指示符 中作了说明。

4.3.2. 指令语句#

指令由一个指令操作码以及其后一个以逗号分隔的、包含零个或多个操作数的列表构成，并以分号结束。操作数可以是寄存器变量、常量表达式、地址表达式或标号名。指令带有可选的守卫谓词，用于控制条件执行。守卫谓词位于可选标号之后、操作码之前，写作 @p，其中 p 是一个谓词寄存器。守卫谓词可以取反，写作 @!p。

目标操作数在最前，其后是源操作数。

指令关键字列于 表 2。所有指令关键字在 PTX 中都是保留记号。

4.4. 标识符#

用户定义的标识符遵循扩展的 C++ 规则：它们要么以一个字母开头，后接零个或多个字母、数字、下划线或美元字符；要么以下划线、美元符号或百分号字符开头，后接一个或多个字母、数字、下划线或美元字符：

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

PTX 没有规定标识符的最大长度，并建议所有实现都支持至少 1024 个字符的最小长度。

许多高级语言（如 C 和 C++）对标识符名遵循类似的规则，只不过不允许使用百分号。PTX 允许将百分号作为标识符的第一个字符。可以使用百分号来避免命名冲突，例如用户定义的变量名与编译器生成的名称之间的冲突。

PTX 预定义了一个常量以及少数以百分号开头的特殊寄存器，列于 表 3。

4.5. 常量#

PTX 支持整数常量、浮点常量以及常量表达式。这些常量可用于数据初始化，也可用作指令的 操作数。整数、浮点和位宽类型的类型检查规则保持不变。对于谓词类型的数据和指令，允许使用 整数常量，并按 C 的方式解释，即零值为 False，非零值为 True。

4.5.1. 整数常量#

整数常量的位宽为 64 位，且为有符号或无符号，即每个整数常量的类型为 .s64 或 .u64。整数常量的有符号/无符号属性对于正确求值包含 除法与有序比较等操作的常量表达式是必需的，因为这些操作的行为取决于操作数类型。当用于 指令或数据初始化时，每个整数常量会根据其使用处的数据或指令类型被转换为合适的大小。

整数字面量可以用十进制、十六进制、八进制或二进制记法书写。语法遵循 C 的语法。 整数字面量之后可以紧跟字母 U，表示该字面量为无符号。

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）；在这两种情况下字面量为无符号（.u64）。

预定义整数常量 WARP_SZ 指定目标平台每个 warp 的线程数；迄今为止， 所有目标架构的 WARP_SZ 值均为 32。

4.5.2. 浮点常量#

浮点常量以 64 位双精度值表示，并且所有浮点常量表达式都使用 64 位双精度算术求值。 唯一的例外是用于表示精确单精度浮点值的 32 位十六进制记法；此类值保留其精确的 32 位 单精度值，并且不得用于常量表达式。每个 64 位浮点常量都会根据其使用处的数据或指令类型 被转换为合适的浮点大小。

浮点字面量书写时可以带有可选的小数点和可选的带符号指数。与 C 和 C++ 不同，这里没有 用于指定位宽的后缀字母；字面量始终以 64 位双精度格式表示。

PTX 包含浮点常量的第二种表示形式，使用十六进制常量来指定精确的机器表示。若要指定 IEEE 754 双精度浮点值，常量以 0d 或 0D 开头，后跟 16 个十六进制 数字。若要指定 IEEE 754 单精度浮点值，常量以 0f 或 0F 开头，后跟 8 个 十六进制数字。

0[fF]{hexdigit}{8}      // single-precision floating point
0[dD]{hexdigit}{16}     // double-precision floating point

示例

mov.f32  $f3, 0F3f800000;       //  1.0

4.5.3. 谓词常量#

在 PTX 中，整数常量可以用作谓词。对于谓词类型的数据初始化器和指令操作数，整数常量 按 C 的方式解释，即零值为 False，非零值为 True。

4.5.4. 常量表达式#

在 PTX 中，常量表达式像 C 一样使用运算符构成，并采用与 C 类似的规则求值，但通过 限制类型与大小、去除大多数类型转换以及定义完整语义而得到简化，从而消除了 C 中表达式 求值依赖实现的情况。

常量表达式由常量字面量、一元正号与负号、基本算术运算符（加法、减法、乘法、除法）、 比较运算符、条件三元运算符（ ?: ）以及括号构成。整数常量表达式还允许一元 逻辑非（!）、按位取反（~）、求余（%）、移位运算符 （<< 和 >>）、位类型运算符（&、 | 和 ^）以及逻辑运算符（&&、 ||）。

PTX 中的常量表达式不支持整数与浮点之间的类型转换。

常量表达式使用与 C 相同的运算符优先级。 表 4 给出了运算符优先级与结合性。一元运算符的优先级最高，并且优先级随图表中每一行的下降而 降低。同一行上的运算符具有相同的优先级，一元运算符从右向左求值，二元运算符从左向右求值。

4.5.5. 整数常量表达式求值#

整数常量表达式在编译时根据一组规则求值，这些规则确定每个子表达式的类型（有符号 .s64 还是无符号 .u64）。这些规则基于 C 中的规则，但已被简化为仅适用于 64 位整数，并且在所有情况下行为都被完整定义（特别是对于求余和移位运算符）。

• 字面量为有符号，除非需要无符号才能防止溢出，或者字面量使用了 U 后缀。例如：

  • 42、0x1234、0123 为有符号。

  • 0xfabc123400000000、42U、0x1234U 为无符号。

• 一元正号和负号保留输入操作数的类型。例如：

  • +123、-1、-(-42) 为有符号。

  • -1U、-0xfabc123400000000 为无符号。

• 一元逻辑非（!）产生值为 0 或 1 的有符号结果。

• 一元按位取反（~）将源操作数按无符号解释，并产生无符号 结果。

• 某些二元运算符要求对源操作数进行规范化。这种规范化称为 通常算术转换，即当任一操作数为无符号时，将两个操作数都转换为无符号类型。

• 加法、减法、乘法和除法执行通常算术转换，并产生与转换后操作数类型 相同的结果。也就是说，若任一源操作数为无符号，则操作数与结果均为无符号，否则为有符号。

• 求余（%）将操作数按无符号解释。注意这与 C 不同，C 允许 负的除数，但将其行为定义为实现相关。

• 左移和右移将第二个操作数按无符号解释，并产生与第一个操作数类型相同的结果。 注意右移的行为由第一个操作数的类型决定：对有符号值右移是算术移位并保留符号，对无符号值 右移是逻辑移位并移入零位。

• 与（&）、或（|）和异或（^）执行通常算术转换，并产生与转换后操作数类型相同的结果。

• AND_OP（&&）、OR_OP（||）、Equal（==）和 Not_Equal（!=）产生有符号结果。结果值为 0 或 1。

• 有序比较（<、<=、>、>=）对源操作数执行通常算术转换，并产生有符号结果。结果值为 0 或 1。

• 支持使用 (.s64) 和 (.u64) 类型转换将表达式转换为有符号或无符号类型。

• 对于条件运算符（ ? : ），第一个操作数必须是整数，且第二个 和第三个操作数必须同为整数或同为浮点数。对第二个和第三个操作数执行通常算术 转换，结果类型与转换后的类型相同。

4.5.6. 常量表达式求值规则汇总#

表 5 汇总了常量表达式求值规则。