7. ABI 抽象#

PTX 不暴露某一特定调用约定、栈布局和应用二进制接口（ABI）的细节，而是提供稍高一层的抽象，以支持多种 ABI 实现。本章介绍 PTX 隐藏 ABI 所需的特性，包括函数定义、函数调用、参数传递和栈上内存分配（alloca）的语法。

有关如何生成符合 CUDA^® 架构应用二进制接口（ABI）的 PTX，请参阅 PTX Writers Guide to Interoperability。

7.1. 函数声明和定义#

在 PTX 中，使用 .func 指示符声明和定义函数。函数声明指定可选的返回参数列表、函数名和可选的输入参数列表；它们共同确定函数的接口或原型。函数定义除接口外还指定函数体。调用函数之前，必须先声明或定义它。

最简单的函数没有参数和返回值，在 PTX 中可写为：

.func foo
{
    ...
    ret;
}
...
call foo;
...
plaintext

此处执行 call 指令会将控制流转移到 foo，并隐式保存返回地址。在 foo 中执行 ret 指令时，控制流转移到调用指令之后的指令。

基本类型的标量和向量输入参数及返回参数可以直接用寄存器变量表示。调用时，实参可以是寄存器变量或常量，返回值可以直接放入寄存器变量。调用处的实参和返回变量在类型及大小上必须与被调函数对应的形参匹配。

示例

.func (.reg .u32 %res) inc_ptr ( .reg .u32 %ptr, .reg .u32 %inc )
{
    add.u32 %res, %ptr, %inc;
    ret;
}
...
call (%r1), inc_ptr, (%r1,4);
...
plaintext

使用 ABI 时，.reg 状态空间中的参数宽度必须至少为 32 位。源语言中小于字宽的标量对象应提升为 PTX 中的 32 位寄存器，或者使用下文所述 .param 状态空间中的字节数组。

C 结构体、联合体等对象在 PTX 中会展平为寄存器或字节数组，并使用 .param 空间的内存表示。例如，考虑按值传给函数的以下 C 结构体：

struct {
    double dbl;
    char   c[4];
};

在 PTX 中，该结构体会展平为字节数组。由于内存访问必须按访问大小的倍数对齐，本例中的结构体是一个按 8 字节对齐的 12 字节数组，以确保访问 .f64 字段时满足对齐要求。通过 .param 状态空间按值传递该结构体：

示例

.func (.reg .s32 out) bar (.reg .s32 x, .param .align 8 .b8 y[12])
{
    .reg .f64 f1;
    .reg .b32 c1, c2, c3, c4;
    ...
    ld.param.f64 f1, [y+0];
    ld.param.b8  c1, [y+8];
    ld.param.b8  c2, [y+9];
    ld.param.b8  c3, [y+10];
    ld.param.b8  c4, [y+11];
    ...
    ... // computation using x,f1,c1,c2,c3,c4;
}
{
.param .b8 .align 8 py[12];
…
st.param.b64 [py+ 0], %rd;
st.param.b8  [py+ 8], %rc1;
st.param.b8  [py+ 9], %rc2;
st.param.b8  [py+10], %rc1;
st.param.b8  [py+11], %rc2;
// scalar args in .reg space, byte array in .param space
call (%out), bar, (%x, py);
…

注意，本例以两种方式使用 .param 空间变量。首先，在函数 bar 的定义中，.param 变量 y 表示形参。其次，在调用函数的函数体中声明 .param 变量 py，用于准备传给 bar 的结构体。

可以从以下概念性角度理解设备函数如何使用 .param 状态空间。

对于调用方：

• 使用 .param 状态空间设置要传给被调函数的值，和／或接收被调函数的返回值。通常用 .param 字节数组汇集按值传递的结构体的各个字段。

对于被调方：

• 使用 .param 状态空间接收参数值，和／或将返回值传回调用方。

参数传递受到以下限制。

对于调用方：

• 实参可以是 .param 变量、.reg 变量或常量。

• 如果 .param 空间中的形参是字节数组，实参也必须是 .param 空间中的字节数组，且类型、大小和对齐方式均须匹配。.param 实参必须在调用方的局部作用域内声明。

• 如果 .param 空间中的形参是基本类型的标量或向量变量，对应实参可以是类型和大小匹配的 .param 或 .reg 空间变量，也可以是能由形参类型表示的常量。

• 如果形参位于 .reg 空间，对应实参可以是类型和大小匹配的 .param 或 .reg 空间变量，也可以是能由形参类型表示的常量。

• 如果形参位于 .reg 空间，该寄存器的宽度必须至少为 32 位。

• 用于传递函数调用实参的所有 st.param 指令，必须紧接在对应 call 指令之前；用于收集返回值的 ld.param 指令，必须紧接在 call 指令之后，期间不得改变控制流。传递实参所用的 st.param 和 ld.param 指令不能带谓词保护。这样既有利于编译器优化，也能确保 .param 变量不会在调用方栈帧中占用超出 ABI 所需的额外空间。.param 变量只是使调用点能够将可能分布在多个位置的数据（例如调用方在寄存器和内存中操作的结构体）映射为可传给被调方的参数或可接收的返回值。

对于被调方：

• 输入参数和返回参数可以是 .param 变量或 .reg 变量。

• .param 内存中的参数必须按 1、2、4、8 或 16 字节的倍数对齐。

• .reg 状态空间中的参数宽度必须至少为 32 位。

• .reg 状态空间可用于接收和返回基本类型的标量及向量值；在非 ABI 模式下编译时，也包括小于字宽的对象。继续支持 .reg 状态空间是为了兼容旧版本。

请注意，参数传递时选用 .reg 还是 .param 状态空间，不影响参数最终通过物理寄存器还是栈传递。参数到物理寄存器和栈位置的映射取决于 ABI 定义以及参数的顺序、大小和对齐方式。

7.1.1. 与 PTX ISA 1.x 相比的变化#

在 PTX ISA 1.x 版中，形参只能位于 .reg 状态空间，也不支持数组参数。C 结构体等对象会展平后通过多个寄存器传递或返回。为此，PTX ISA 1.x 支持多个返回值。

从 PTX ISA 2.0 版起，形参可以位于 .reg 或 .param 状态空间，.param 空间参数还支持数组。对于 sm_20 或更高的目标架构，PTX 将函数限制为单个返回值；无法放入一个寄存器的对象应通过 .param 字节数组返回。对于 sm_1x 目标，PTX 仍支持多个返回寄存器。

PTX ISA 3.0 之前的版本允许在模块作用域定义 .reg 和 .local 状态空间变量。从 PTX ISA 3.0 版起，使用 ABI 编译时禁止模块作用域的 .reg 和 .local 变量，并将其限制在函数作用域内。未使用 ABI 编译时，仍与以往一样支持模块作用域的 .reg 和 .local 变量。编译包含此类模块作用域变量的旧 PTX 代码（ISA 版本低于 3.0）时，编译器会静默禁用 ABI。

7.2. 可变参数函数#

PTX 6.0 版支持向函数传入未指定大小的数组参数，可用它实现可变参数函数。

详见Kernel 与函数指示符：.func。

7.3. Alloca#

PTX 提供 alloca 指令，用于在运行时从每线程的局部内存栈分配存储空间。使用 alloca 返回的指针，可通过 ld.local 和 st.local 指令访问所分配的栈内存。

为便于释放 alloca 分配的内存，PTX 还提供两条指令：stacksave 可将栈指针的值读取到局部变量，stackrestore 则可用保存的值恢复栈指针。

alloca、stacksave 和 stackrestore 指令见栈操作指令。