PTX编写者的互操作性指南

编写符合ABI规范的PTX指南。

1. 简介

本文档定义了生成PTX时CUDA^®架构的应用二进制接口(ABI)。通过遵循此ABI，外部开发者可以生成符合规范的PTX代码，并能与其他代码进行链接。

PTX是一种低级的并行线程执行虚拟机和指令集架构(ISA)。PTX可以由多种工具生成或由开发者直接编写。PTX的设计目标是实现GPU架构无关性，使得同一套代码能够在不同的GPU架构上重复使用。如需了解更多关于PTX的信息，请参阅最新版的PTX ISA参考文档。

CUDA包含多个架构系列，每个系列都有其专属的指令集架构（ISA）；例如SM 5.x属于麦克斯韦架构家族，SM 6.x属于帕斯卡架构家族。本文档描述了适用于所有架构的高级应用二进制接口（ABI）。符合特定ABI的程序应当在对应架构的GPU上执行，并可以假定该指令集架构中的指令均可使用。

2. 数据表示

2.1. 基础类型

下表展示了所支持的原生标量PTX类型。任何PTX生成器都必须使用这些大小和对齐方式，以确保其生成的PTX与其他生成器产生的PTX兼容。PTX还支持原生向量类型，相关讨论请参阅聚合与联合。

类型的大小由主机定义。例如，指针大小和长整型大小由主机的ABI决定。PTX有一个.address_size指令，用于指定整个PTX代码中使用的地址大小。指针的大小在32位主机上是32位，在64位主机上是64位。然而，本地内存空间和共享内存空间的地址大小始终为32位。

在单独编译过程中，我们会将主机平台信息存储在每个目标文件中。链接器将无法链接为不兼容主机平台生成的目标文件。

PTX 类型	大小 (字节)	对齐 (字节)	硬件表示
.b8	1	1	无类型字节
.b16	2	2	无类型半字
.b32	4	4	无类型字
.b64	8	8	无类型双字
.s8	1	1	有符号整数字节
.s16	2	2	有符号整型半字
.s32	4	4	有符号整数字
.s64	8	8	有符号双字整型
.u8	1	1	无符号整数字节
.u16	2	2	无符号半字整型
.u32	4	4	无符号整数字
.u64	8	8	无符号双字整型
.f16	2	2	IEEE半精度
.f32	4	4	IEEE单精度
.f64	8	8	IEEE双精度浮点数

2.2. 聚合与联合

除了标量类型外，PTX还支持这些标量类型的原生向量类型，包括向量语法和字节数组语法。对于大小不超过4字节的标量类型，存在1、2、3和4元素的向量类型；对于所有其他类型，仅存在1和2元素的向量类型。

所有聚合和联合操作都可以在PTX中通过其字节数组语法得到支持。

以下是所有聚合体和联合体的大小和对齐规则。

• 对于非原生向量类型，整个聚合体或联合体会按照其最严格对齐成员的对齐边界进行对齐。如果对齐方式由输入语言定义，则不遵循此规则。例如，在OpenCL中，内置向量数据类型的对齐方式设置为内置数据类型的字节大小。

• 对于原生向量类型——本节开头讨论过——对齐定义如下。（对于以下定义，原生向量包含n个元素，且元素类型为t。）

  • 对于元素个数为奇数的向量，其对齐方式与其成员相同：alignof(t)。

  • 对于一个元素数量为偶数的向量，其对齐方式设置为元素数量乘以其成员的对齐值：n*alignof(t)。

• 每个成员会被分配到最低可用偏移量处，并保持适当对齐。根据前一个成员的不同，可能需要进行内部填充。

• 如有必要，聚合体或联合体的大小会被增加，使其成为聚合体或联合体对齐大小的整数倍。根据最后一个成员的不同，这可能需要尾部填充。

2.3. 位字段

C语言的结构体和联合体定义可以包含位字段，用于指定整数对象所占用的位数。

位域类型	宽度 w	范围
有符号字符	1 到 8	-2w-1 到 2w-1 - 1
无符号字符	1到8	0到2w - 1
有符号短整型	1 到 16	-2w-1 到 2w-1 - 1
无符号短整型	1到16	0到2w - 1
有符号整型	1到32位	-2w-1 至 2w-1 - 1
无符号整型	1到32	0到2w - 1
有符号长整型	1到64位	-2w-1 至 2w-1 - 1
无符号长整型	1到64	0到2w - 1

当前GPU仅支持小端内存模式，因此以下内容均基于小端字节序布局。

以下是适用于位字段的规则。

• 普通的位字段（既未指定有符号也未指定无符号）会被视为有符号类型。

• 当未提供类型时（例如指定了 signed : 6），默认类型为 int。

位域遵循与其他结构和联合成员相同的大小和对齐规则，但有以下修改。

• 对于小端序系统，位字段在内存中是从右向左（从最低位到最高位）分配的。

• 位字段必须完全位于适合其声明类型的存储单元内。位字段绝不应跨越其单元边界。

• 位字段可以与其它结构和联合成员共享一个存储单元，包括非位字段的成员，只要存储单元内有足够的空间。

• 未命名的位字段不会影响结构体或联合体的对齐方式。

• 零长度位字段会强制结构体的下一个成员对齐到与该位字段类型对应的下一个对齐边界。未命名的零长度位字段不会强制结构体的外部对齐到该边界。如果未命名的零长度位字段具有比外部对齐更严格的对齐要求，当结构体或联合体分配到内存时，不能保证会维持这种更严格的对齐。

以下图表展示了位字段的示例。图1显示了示例中使用的字节偏移量（上角）和位编号（下角）。其余图表展示了不同的位字段示例。

位编号

位域分配

边界对齐

存储单元共享

联合分配

未命名字段

2.4. 纹理、采样器和表面类型

纹理(texture)、采样器(sampler)和表面(surface)类型用于定义对纹理和表面内存的引用。CUDA架构提供了硬件和指令，可以高效地从纹理或表面内存(而非全局内存)读取数据。

纹理引用通过运行时函数绑定到设备的只读内存区域（称为纹理内存），然后才能被内核使用。纹理引用具有多个属性，例如归一化模式、寻址模式和纹理过滤等。采样器引用可用于在内核中读取时对纹理进行采样。表面引用用于从表面内存读取或写入数据。它还具有类似于纹理的各种属性。

在PTX级别，访问纹理或表面内存的对象被称为不透明对象。纹理由.texref或.samplerref类型表示，表面由.surfref类型表示。不透明对象的数据可以通过特定指令访问（TEX用于.texref/.samplerref，SULD/SUST用于.surfref）。不透明对象的属性通过在内存中分配一个由驱动程序填充的描述符来实现。PTX TXQ/SUQ指令会被转换为对描述符字段的内存读取。描述符的内部格式随架构而异，用户不应依赖它。如果在编译时已知纹理或表面引用，可以直接访问不透明对象的数据和属性，否则间接访问。如果在编译时未知引用，则读取数据和属性所需的所有信息都包含在一个称为句柄的.b64值中。该句柄可用于在函数之间传递和返回不透明对象引用，以及引用外部纹理、采样器和表面。

3. 函数调用序列

本节描述PTX级别的函数调用序列，包括寄存器使用、栈帧布局和参数传递。PTX级函数调用序列说明了在PTX中如何表示以实现函数调用。这一层级存在抽象概念，与函数调用序列相关的大部分细节都在SASS级别处理。

早于2.0版本的PTX不符合本文档定义的ABI规范，无法执行ABI兼容的函数调用。要使调用约定生效，必须使用PTX 2.0或更高版本。

3.1. 寄存器

在PTX级别，指定的寄存器是虚拟的。寄存器分配发生在PTX到SASS的转换过程中。PTX到SASS的转换还会将参数和返回值转换为物理寄存器或堆栈位置。

3.2. 栈帧

PTX层级没有软件栈的概念。栈的操作完全在SASS层级定义，并在PTX到SASS的转换过程中进行分配。

3.3. 参数传递

在PTX级别，设备函数中的所有参数和返回值都使用参数状态空间(.param)。下表包含了处理在源代码级别定义的参数和返回值的规则。针对每个源代码级别的类型，提供了应使用的对应PTX级别类型。

源类型	位大小	PTX类型
整数类型	8到32位(A)	.u32(无符号)或.s32(有符号)
整数类型	64	.u64 (无符号) 或 .s64 (有符号)
指针 (B)	32	.u32
指针 (B)	64	.u64
浮点类型 (C)	32	.f32
浮点类型 (C)	64	.f64
聚合或联合	任意尺寸	.align align .b8 name[size] 其中align表示整体聚合体或联合体的字节对齐值(D)，name是与聚合体或联合体关联的变量名，size表示聚合体或联合体的字节大小。
句柄 (E)	64	.b64 (从 .texref, .sampleref, .surfref 分配)

注意事项：

1. 小于32位的数值会根据其是否为有符号或无符号类型进行符号扩展或零扩展。

2. 除非在函数声明中指定了内存类型，否则在PTX级别传递的所有指针都必须使用通用地址。

3. 16位浮点类型仅用于存储。因此，它们不能用作参数或返回值。

4. 对齐必须是1、2、4、8、16、32、64或128字节。

5. PTX内置的不透明类型，如纹理(texture)、采样器(sampler)和表面(surface)类型，可以通过64位句柄作为参数传递给函数，并由函数返回。该句柄包含从纹理或表面内存访问实际数据所需的信息，以及存储在其类型描述符中的对象属性。有关句柄的更多信息，请参阅章节Texture, Sampler, and Surface Types。

4. 系统调用

系统调用是指调用驱动程序操作系统代码。在PTX中，它们看起来像常规调用，但没有给出函数定义。PTX文件中必须提供函数原型，但函数的实现由驱动程序提供。

vprintf系统调用的原型是：

.extern .func (.param .s32 status) vprintf (.param t1 format, .param t2 valist)

以下是vprintf函数的参数和返回值的定义。

• status : 由vprintf返回的状态值。

• format : 指向格式说明符输入的指针。对于32位地址，类型t1为.b32。对于64位地址，类型t1为.b64。

• valist : 指向valist输入的指针。对于32位地址，类型t2为.b32。对于64位地址，类型t2为.b64。

使用32位地址调用vprintf的示例如下：

cvta.global.b32    %r2, _fmt;
st.param.b32  [param0], %r2;
cvta.local.b32  %r3, _valist_array;
st.param.b32  [param1], %r3;
call.uni (_), vprintf, (param0, param1);

对于这段代码，_fmt是全局内存中的格式字符串，_valist_array是参数的valist。请注意，任何指针都必须转换为通用空间。vprintf系统调用作为"stdio.h"中定义的printf函数的一部分被发出。

malloc系统调用的原型是：

.extern .func (.param t1 ptr) malloc (.param t2 size)

以下是malloc参数和返回值的定义。

• ptr : 指向由malloc分配的内存的指针。对于32位地址，类型t1为.b32。对于64位地址，类型t1为.b64。

• size : 从malloc申请的内存大小。该大小由size_t类型定义。当size_t为32位时，类型t2为.b32。当size_t为64位时，类型t2为.b64。

free系统调用的原型是：

.extern .func free (.param t1 ptr)

以下是自由参数的定义。

• ptr : 指向应释放内存的指针。对于32位地址，类型t1为.b32。对于64位地址，类型t1为.b64。

malloc和free系统调用作为“malloc.h”中定义的malloc和free函数的一部分被发出。

为了支持断言功能，当断言表达式结果为假时，会调用PTX函数__assertfail。该系统调用的原型如下：

.extern .func __assertfail (.param t1 message, .param t1 file, .param .b32 line, .param t1 function, .param t2 charSize)

以下是__assertfail参数的定义。

• message : 指向应输出字符串的指针。对于32位地址，类型t1为.b32。对于64位地址，类型t1为.b64。

• file : 指向与断言关联的文件名字符串的指针。对于32位地址，类型t1为.b32。对于64位地址，类型t1为.b64。

• line : 与该断言关联的行号。

• function : 指向与断言相关联的函数名称字符串的指针。对于32位地址，类型t1为.b32。对于64位地址，类型t1为.b64。

• charSize：__assertfail参数字符串中包含的字符的字节大小。唯一支持的字符大小为1。字符大小由类型size_t定义。当size_t为32位时，类型t2为.b32。当size_t为64位时，类型t2为.b64。

__assertfail 系统调用是作为“assert.h”中定义的 assert 宏的一部分发出的。

5. 原子操作应用二进制接口

编程语言的原子操作到PTX ISA的映射需要以一致的方式在所有可能并发访问共享内存的编程语言中实现。 针对CUDA架构的C++11原子操作映射在A Formal Analysis of the NVIDIA PTX Memory Consistency Model中被证明是正确的。 PTX ISA为获取、释放、获取-释放以及宽松的C++内存排序语义提供了原子内存操作和栅栏。 以下是C++顺序一致原子操作的PTX ABI：

C 或 C++ 或 CUDA C++ API	PTX ABI ISA 映射
atomic_thread_fence(memory_order_seq_cst, thread_scope_)	fence.sc.;
atomic_load(memory_order_seq_cst, thread_scope_)	fence.sc.; ld.acquire.;
atomic_store(memory_order_seq_cst, thread_scope_)	fence.sc.; st.release.;
atomic_ op>(memory_order_seq_cst, thread_scope_)	fence.sc.; atom.acq_rel.. op>;

6. 调试信息

调试信息以DWARF（任意记录格式调试）格式编码。

6.1. 调试信息生成

生成调试信息的职责由PTX生产者和PTX到SASS的后端共同承担。PTX生产者负责使用PTX中的.section和.b8-.b16-.b32-and-.b64指令将二进制DWARF信息输出到PTX文件中。这应包含.debug_info和.debug_abbrev部分，以及可选的.debug_pubnames和.debug_aranges部分。这些部分是标准的DWARF2部分，它们引用PTX中的标签和寄存器。

PTX-to-SASS后端负责从PTX文件中的.file和.loc指令生成.debug_line段。该段将源代码行映射到SASS地址。后端还会生成.debug_frame段。

6.2. CUDA专用DWARF定义

为了支持多内存段的调试，定义了地址类别代码以反映变量的内存空间。所有变量和参数的调试信息条目(Debugging Information Entries)都会将地址类别值作为DW_AT_address_class属性输出。地址类别代码定义如下表所示。

代码	值	描述
ADDR_code_space	1	代码存储空间
ADDR_reg_space	2	寄存器存储空间
ADDR_sreg_space	3	特殊寄存器存储空间
ADDR_const_space	4	常量存储空间
ADDR_global_space	5	全局存储空间
ADDR_local_space	6	本地存储
ADDR_param_space	7	参数存储空间
ADDR_shared_space	8	共享存储空间
ADDR_surf_space	9	表面存储空间
ADDR_tex_space	10	纹理存储
ADDR_tex_sampler_space	11	纹理采样器存储空间
ADDR_generic_space	12	通用地址存储空间

7. 示例

以下是带有调试信息的示例PTX代码，用于实现以下调用程序的程序：

__device__ __noinline__ int foo (int i, int j)
{
  return i+j;
}

__global__ void test (int *p)
{
  *p = foo(1, 2);
}

生成的PTX代码大致如下：

.version 4.2
.target sm_20, debug
.address_size 64

 .file   1 "call_example.cu"

.visible .func  (.param .b32 func_retval0) // return value
_Z3fooii(
        .param .b32 _Z3fooii_param_0, // parameter "i"
        .param .b32 _Z3fooii_param_1) // parameter "j"
{
        .reg .s32       %r<4>;
        .loc 1 1 1      // following instructions are for line 1

func_begin0:
        ld.param.u32    %r1, [_Z3fooii_param_0]; // load 1st param
        ld.param.u32    %r2, [_Z3fooii_param_1]; // load 2nd param
        .loc    1 3 1   // following instructions are for line 3
        add.s32         %r3, %r1, %r2;
        st.param.b32    [func_retval0+0], %r3; // store return value
        ret;
func_end0:
}

.visible .entry _Z4testPi(
        .param .u64 _Z4testPi_param_0) // parameter *p
{
        .reg .s32       %r<4>;
        .reg .s64       %rd<2>;
        .loc 1 6 1

func_begin1:
        ld.param.u64    %rd1, [_Z4testPi_param_0]; // load *p
        mov.u32         %r1, 1;
        mov.u32         %r2, 2;
        .loc    1 8 9
        .param .b32 param0;
        st.param.b32    [param0+0], %r1; // store 1
        .param .b32 param1;
        st.param.b32    [param1+0], %r2; // store 2
        .param .b32 retval0;
        call.uni (retval0), _Z3fooii, ( param0, param1); // call foo
        ld.param.b32    %r3, [retval0+0]; // get return value
        st.u32  [%rd1], %r3;              // *p = return value
        .loc    1 9 2
        ret;
func_end1:
}

.section .debug_info {
 .b32 262
 .b8 2, 0
 .b32 .debug_abbrev
 .b8 8, 1, 108, 103, 101, 110, 102, 101, 58, 32, 69, 68, 71, 32, 52, 46, 57
 .b8 0, 4, 99, 97, 108, 108, 49, 46, 99, 117, 0
 .b64 0
 .b32 .debug_line // the .debug_line section will be created by ptxas from the .loc
 .b8 47, 104, 111, 109, 101, 47, 109, 109, 117, 114, 112, 104, 121, 47, 116
 .b8 101, 115, 116, 0, 2, 95, 90, 51, 102, 111, 111, 105, 105, 0, 95, 90
 .b8 51, 102, 111, 111, 105, 105, 0
 .b32 1, 1, 164
 .b8 1
 .b64 func_begin0 // start and end location of foo
 .b64 func_end0
 .b8 1, 156, 3, 105, 0
 .b32 1, 1, 164
 .b8 5, 144, 177, 228, 149, 1, 2, 3, 106, 0
 .b32 1, 1, 164
 .b8 5, 144, 178, 228, 149, 1, 2, 0, 4, 105, 110, 116, 0, 5
 .b32 4
 .b8 2, 95, 90, 52, 116, 101, 115, 116, 80, 105, 0, 95, 90, 52, 116, 101
 .b8 115, 116, 80, 105, 0
 .b32 1, 6, 253
 .b8 1
 .b64 func_begin1 // start and end location of test
 .b64 func_end1
 .b8 1, 156, 3, 112, 0
 .b32 1, 6, 259
 .b8 9, 3
 .b64 _Z4testPi_param_0
 .b8 7, 0, 5, 118, 111, 105, 100, 0, 6
 .b32 164
 .b8 12, 0
}
.section .debug_abbrev {
 .b8 1, 17, 1, 37, 8, 19, 11, 3, 8, 17, 1, 16, 6, 27, 8, 0, 0, 2, 46, 1, 135
 .b8 64, 8, 3, 8, 58, 6, 59, 6, 73, 19, 63, 12, 17, 1, 18, 1, 64, 10, 0, 0
 .b8 3, 5, 0, 3, 8, 58, 6, 59, 6, 73, 19, 2, 10, 51, 11, 0, 0, 4, 36, 0, 3
 .b8 8, 62, 11, 11, 6, 0, 0, 5, 59, 0, 3, 8, 0, 0, 6, 15, 0, 73, 19, 51, 11
 .b8 0, 0, 0
}
.section .debug_pubnames {
 .b32 41
 .b8 2, 0
 .b32 .debug_info
 .b32 262, 69
 .b8 95, 90, 51, 102, 111, 111, 105, 105, 0
 .b32 174
 .b8 95, 90, 52, 116, 101, 115, 116, 80, 105, 0
 .b32 0
}

8. C++

C++设备函数的实现遵循Itanium C++ ABI规范。但并非所有C++特性都受支持。特别需要注意的是，以下特性在设备代码中不被支持。

• 异常与try/catch代码块

• RTTI

• STL库

• 全局构造函数和析构函数

• 主机和设备之间的虚函数和类（即，vtable不能跨主机和设备使用）

还有一些目前不支持的C语言特性：

• 除printf外的标准输入输出

9. 公告

9.1. 注意事项

本文档仅供信息参考之用，不应视为对产品功能、状态或质量的保证。NVIDIA公司（“NVIDIA”）对本文件所含信息的准确性或完整性不作任何明示或暗示的陈述或保证，并对其中可能存在的错误不承担任何责任。NVIDIA对于因使用此类信息而产生的后果、或因使用该信息导致的第三方专利或其他权利侵权概不负责。本文件不构成对开发、发布或交付任何材料（定义见下文）、代码或功能的承诺。

NVIDIA保留随时对本文件进行更正、修改、增强、改进以及任何其他变更的权利，恕不另行通知。

客户在下单前应获取最新的相关信息，并确认这些信息是最新且完整的。

除非NVIDIA与客户授权代表签署的单独销售协议中另有约定，否则NVIDIA产品的销售均以订单确认时提供的NVIDIA标准销售条款和条件为准（以下简称"销售条款"）。NVIDIA特此明确反对将任何客户通用条款适用于本文件所述NVIDIA产品的采购。本文件不直接或间接构成任何合同义务。

NVIDIA产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于那些可以合理预期NVIDIA产品故障或失灵会导致人身伤害、死亡、财产或环境损害的应用场景。NVIDIA对于在此类设备或应用中使用和/或包含NVIDIA产品不承担任何责任，因此客户需自行承担相关风险。

NVIDIA不声明或保证基于本文档的产品适用于任何特定用途。NVIDIA未必会对每个产品的所有参数进行测试。客户应全权负责评估和确定本文档所含信息的适用性，确保产品适合并满足客户计划的应用需求，并执行必要的应用测试以避免应用或产品出现故障。客户产品设计中的缺陷可能会影响NVIDIA产品的质量和可靠性，并可能导致超出本文档范围的其他或不同的条件和/或要求。对于任何因以下原因导致的故障、损坏、成本或问题，NVIDIA不承担任何责任：(i) 以违反本文档的任何方式使用NVIDIA产品或(ii) 客户产品设计。

本文档不授予任何NVIDIA专利权、版权或其他NVIDIA知识产权的明示或暗示许可。NVIDIA发布的关于第三方产品或服务的信息，不构成NVIDIA对这些产品或服务的使用许可或担保认可。使用此类信息可能需要获得第三方基于其专利或其他知识产权的许可，或需要获得NVIDIA基于其专利或其他知识产权的许可。

本文件中的信息仅可在获得NVIDIA事先书面批准、未经改动完整复制且完全符合所有适用的出口法律法规，并附带所有相关条件、限制和声明的情况下进行复制。

本文件及所有NVIDIA设计规格、参考板、文件、图纸、诊断工具、清单和其他文档（统称及单独称为"材料"）均以"现状"提供。NVIDIA不对材料作出任何明示或默示的保证，包括但不限于对不侵权、适销性和特定用途适用性的默示保证免责。在法律允许的最大范围内，NVIDIA不就因使用本文件导致的任何损害承担责任，包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论损害成因如何，也无论责任理论为何，即使NVIDIA已被告知发生此类损害的可能性。不论客户因任何原因可能遭受的任何损害，NVIDIA对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。

9.2. OpenCL

OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。

9.3. 商标

NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。