了解 CUDA GPU 计算的汇编语言 PTX

并行线程执行（PTX）是一种虚拟机指令集架构，从一开始就是 CUDA 的一部分。您可以将 PTX 视为 NVIDIA CUDA GPU 计算平台的汇编语言。

在本文中，我们将解释这意味着什么，PTX 的用途，以及您需要了解哪些内容才能在您的应用中充分利用 CUDA。我们将首先介绍 CUDA 如何生成、存储和加载最终在 GPU 上运行的代码。然后，我们将展示 PTX 如何实现前向兼容性，以及如何使用 PTX 让特定领域的编程语言和其他编程语言面向 CUDA。

指令集架构

指令集架构（ISA）是对处理器可以执行的指令、其格式、这些指令的行为以及二进制编码的规范。每个处理器都有 ISA。例如，x86_64 是 CPU ISA。ARM64 是另一类。GPU 也具有 ISA。对于 NVIDIA GPU，对于不同世代的 GPU，甚至是一代内不同产品线的 GPU，ISA 可能会有所不同。

虚拟机 ISA 是对一组适用于虚拟处理器的受支持指令、格式和行为的规范。也就是说，它只是抽象处理器的 ISA，而不是实际生成的处理器。虚拟机 ISA 可能不会为指令指定二进制编码，因为只有在物理处理器上运行时才需要这种编码。

PTX 在 CUDA 平台中的作用

为说明 PTX 如何融入 CUDA 平台，以下示例展示了如何编译简单的 CUDA 文件。此源文件包含单个核函数 (即并行添加两个向量的经典示例) ，以及应用程序主函数和辅助函数的骨架。

__global__ void vecAdd(float* a, float* b, float* c, int n)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < n)
    {
        c[index] = a[index] + b[index];
    }
}


void vecAddLauncher(float* a, float* b, float* c, int n)
{
   … // Utility function for launching the kernel
}


int main()
{
 … // Main function of the application
}

使用 NVIDIA CUDA 编译器 NVCC 编译此文件时，系统会将源代码分为 GPU 代码和 CPU 代码。GPU 代码被发送到 GPU 编译器，CPU 代码被发送到主机编译器。主机编译器不是 NVCC 的一部分。NVCC 会调用在命令行中传入的主机编译器，或者该编译器是系统上的默认编译器。

对于将在 GPU 上运行的函数和核函数，GPU 编译器会为 CUDA 平台生成汇编语言：PTX。这段程序集 (PTX) 然后通过一个名为 ptxas 的汇编程序运行，该程序会为 GPU 生成可执行的二进制代码。GPU 二进制文件称为 cubin ，是 CUDA 二进制文件的简称。

编译 GPU 代码包括两个阶段：首先，将高级语言代码 (C++) 编译为 PTX。然后，将 PTX 编译成一个 cubin。请求二进制输出时，NVCC 会自动调用 ptxas。

此编译路径与热门编译器类似 clang 操作。Clang 首先将代码编译到名为 LLVM IR 的虚拟机 ISA 中。LLVM 代表低级虚拟机，IR 代表中间表示。第二阶段或后端编译器称为 LLVM，然后将虚拟机表示形式 LLVM IR 编译为特定处理器的可执行代码。此结构的一个优势是，对于 LLVM 后端编译器支持的任何硬件架构，可以将程序的 LLVM IR 编译为二进制文件。

PTX 与 LLVM IR 类似，因为程序的 PTX 表示可以编译到各种 NVIDIA GPU 中。重要的是，针对特定 GPU 的 PTX 编译可以在应用程序运行时即时完成 (JIT) 。如图 1 所示，应用程序的可执行文件可以嵌入 GPU 二进制文件 (cubins) 和 PTX 代码。通过在可执行文件中嵌入 PTX，CUDA 可在应用程序运行时通过 JIT 将 PTX 编译到相应的 cubin。适用于 PTX 的 JIT 编译器是 NVIDIA GPU 驱动 的一部分。

在应用程序中嵌入 PTX 可在编译应用程序时运行编译的第一阶段——高级语言到 PTX。当应用程序运行时，可以延迟编译的第二阶段——PTX 到 cubin。如下所示，这样做可以让应用在更广泛的 GPU 上运行，包括构建应用后很久才发布的 GPU。

以下是上述示例中 vecAdd 核函数的 PTX 代码。那些看过任何平台的汇编语言的人应该会发现 PTX 的语法和格式都很熟悉。无需了解代码的细节。而是为了让大家一窥 PTX，并进一步阐明 PTX 是什么：CUDA 平台的汇编语言。

.visible .entry _Z6vecAddPfS_S_j(
.param .u64 _Z6vecAddPfS_S_j_param_0,
.param .u64 _Z6vecAddPfS_S_j_param_1,
.param .u64 _Z6vecAddPfS_S_j_param_2,
.param .u32 _Z6vecAddPfS_S_j_param_3
)
{
.reg .pred %p<2>;
.reg .f32 %f<4>;
.reg .b32 %r<6>;
.reg .b64 %rd<11>;


ld.param.u64 %rd1, [_Z6vecAddPfS_S_j_param_0];
ld.param.u64 %rd2, [_Z6vecAddPfS_S_j_param_1];
ld.param.u64 %rd3, [_Z6vecAddPfS_S_j_param_2];
ld.param.u32 %r2, [_Z6vecAddPfS_S_j_param_3];
mov.u32 %r3, %tid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %ctaid.x;
mad.lo.s32 %r1, %r5, %r4, %r3;
setp.ge.u32 %p1, %r1, %r2;
@%p1 bra $L__BB0_2;


cvta.to.global.u64 %rd4, %rd1;
mul.wide.u32 %rd5, %r1, 4;
add.s64 %rd6, %rd4, %rd5;
cvta.to.global.u64 %rd7, %rd2;
add.s64 %rd8, %rd7, %rd5;
ld.global.f32 %f1, [%rd8];
ld.global.f32 %f2, [%rd6];
add.f32 %f3, %f2, %f1;
cvta.to.global.u64 %rd9, %rd3;
add.s64 %rd10, %rd9, %rd5;
st.global.f32 [%rd10], %f3;


$L__BB0_2:
ret;
}

计算能力和 NVIDIA GPU 硬件 ISAs

所有 NVIDIA GPU 都有一个版本标识符 计算能力 或 CC 编号。每种计算能力都有一个主要版本号和一个次要版本号。例如，计算能力 8.6 的主要版本为 8，次要版本为 6。

与其他处理器一样，NVIDIA GPU 也有特定的 ISA。不同代的 GPU 具有不同的 ISA。这些 ISA 通过与 GPU 计算能力相对应的版本号进行识别。当编译二进制文件 (cubin) 时，会针对特定的计算能力对其进行编译。

例如，GeForce 和 RTX GPU NVIDIA Ampere generation 的计算能力为 8.6，其 cubin 版本为 sm_86。所有 cubin 版本的格式为 sm_XY，其中 X 和 Y 对应计算能力的主要和次要数字。

NVIDIA 不同代的 GPU，甚至同一代产品中的不同产品，都可能具有不同的 ISA。这是采用 PTX 的部分原因。

PTX：一种版本化的虚拟机汇编语言

PTX 是虚拟机 ISA。如前所述，虚拟机 ISA 是一组由假设的处理器 (而非任何特定的真实处理器) 支持的指令集。由于虚拟机 ISA 比实际硬件 ISA 略抽象，因此 CUDA 汇编程序 ptxas 的行为更像是编译器，而不是传统汇编程序。它将 PTX 程序编译为 GPU 二进制文件，或 cubins。

自首次推出 CUDA 以来，GPU 功能不断发展。随着新一代 GPU 硬件的推出，GPU 将新增功能。PTX 描述的虚拟机也进行了扩展以匹配。对 PTX 规范的更改通常涉及添加新指令。

因此，有不同版本的 PTX 支持不同的指令集。PTX 版本号表示虚拟架构中可用的指令。与 cubin 版本一样，这些版本号对应 GPU 计算能力。

例如， NVIDIA Ampere generation GA100 GPU 的计算能力为 8.0。名为 compute_80 的 PTX 版本具有 GA100 支持的所有指令。PTX 版本称为 compute_XY，其中 X 和 Y 对应计算能力的主要和次要数量。

GPU 代码兼容性

CUDA 为不同 GPU 之间的代码兼容性提供了两种机制：二进制兼容性和 PTX JIT 兼容性。

二进制兼容性

NVIDIA GPUs 在主要计算能力版本中与二进制兼容，前提是次要版本相同或更高。这意味着为 sm_86 编译的 cubin 可以加载到 x 大于或等于 6 的任何 sm_8x 上。

例如，为 sm_86 (例如 NVIDIA RTX A4000) 编译的 cubin 也可以在 sm_89 (例如 NVIDIA RTX 4000 Ada Generation) 上加载和运行。但是，它不会在计算能力为 8.0 的设备上加载，因为该 GPU 计算能力的次要版本低于 cubin 的次要版本。

在主要计算能力版本中，NVIDIA GPUs 并不二进制兼容。为 sm_86 编译的 cubin 将不会在 9.0 ( NVIDIA Hopper 架构 ) 或更高版本的 GPU 上加载和运行。

PTX JIT 兼容性

在可执行文件中嵌入 PTX 提供了一种机制，可在单个二进制文件中跨不同计算能力的 GPU (包括不同的主要版本) 实现兼容性。如图 1 中的可执行文件所示，PTX 和 cubin 均可存储在最终应用程序可执行文件中。PTX 和 cubin 也可以存储在库中。

当 PTX 代码存储在应用程序或库二进制文件中时，可以为加载该代码的 GPU 进行 JIT 编译。例如，如果应用程序或库包含面向 compute_70 的 PTX，则可以针对任何计算能力 7.0 或更高版本的 GPU (包括计算能力 8.x、9.x、10.x 和 12.x) 对该 PTX 进行 JIT 编译。

如果计算能力低于 PTX 版本，则无法对 PTX 进行 JIT 编译。例如，针对 compute_70 的 PTX 无法针对 5.x 或 6.xGPU 的计算能力进行 JIT 编译。

Fatbins

在构建 CUDA 应用程序或库时，它们有一个名为 fatbin 的容器。fatbin 可以包含多个 cubins 和 PTX 版本的 GPU 代码。例如，图 2 所示的可执行文件中的 fatbin 包含 compute_70 的 PTX 以及 sm_70、sm_80 和 sm_86 的 cubin。这意味着应用程序已有用于计算能力 7.0、8.0 和 8.6 的 GPU 二进制代码。如果应用程序在计算能力 8.9 的 GPU 上运行，也可以加载 sm_86 cubin。

compute_70 PTX 可用于任何计算能力为 7.0 或更高的 GPU 的 JIT 编译，因此此应用程序可以在更新的 GPU 上运行，而不是可用 cubin 的目标。例如，此应用程序可以在计算能力为 9.0、10.0 或 12.0 的 GPU 上运行，而无需重建应用程序。表 1 显示了每个嵌入式 cubin 和 PTX 如何在此特定示例中实现兼容性。

当应用程序启动或代码首次在 GPU 上使用时，驱动会编译 PTX。有关控制 JIT 编译发生时间的详细信息，请参阅《CUDA 编程指南》中的 Lazy Loading 部分。

PTX 的优势

CUDA 平台将 PTX 用作中间代码格式，使开发者能够构建将在尚未创建的 GPU 上运行的应用程序二进制文件。2018 年为 NVIDIA Turing 架构 (CC 7.5) 编译的应用程序可以在 2025 年的 NVIDIA Blackwell (CC 12.0) GPU 上运行，也可以在未来发布的 GPU 上运行。

通过在可执行文件或库中嵌入 GPU 代码的 PTX 表示，CUDA 驱动程序可以在运行时 JIT 编译 PTX 代码，以用于编译应用程序时甚至没有构思的架构。对于分发其应用程序或库的二进制版本的开发者而言，这允许应用程序或库在未来的 GPU 架构上运行，而无需更新二进制文件。

作为 CUDA GPU 计算平台的汇编语言，PTX 还提供了一种可供任何语言的编译器使用的表示形式。例如，域特定语言 (DSL) 编译器可以生成 PTX 代码，然后在 NVIDIA GPU 上运行。 OpenAI Triton 是一个用于生成 PTX 的 DSL 示例。

有兴趣制作 DSL 的开发者应探索 NVVM IR 以及 libNVVM 因为这些可能比实现定制的 PTX 生成更可取。

手写 PTX

NVIDIA 记录了 PTX 虚拟机 ISA 。您可以手写 PTX 代码。与其他汇编语言一样，这通常不是大型软件项目的好选择。与直接处理汇编语言或虚拟汇编语言相比，更高级别的语言可提高开发者的工作效率。

手动编写 PTX 应视为与手动编写 CPU 组件代码类似：它允许专家深入了解处理器将如何执行代码，从而对最终可执行文件中的指令进行精细控制。虽然这样做可以提高性能，但对于大多数开发者来说，这通常不是必要的，也不可取。

也就是说，一些开发者确实选择通过直接为将运行数十亿次或更多次的内部循环代码编写 PTX 来优化代码。在此类代码中，即使是微小的性能提升也会乘以大量的行程数量，因此仔细和手动优化工作也很有价值。我们计划在后续文章中提供此类优化的一些示例。

CUDA 工具包 中包含的 libcu++ 提供了一个 cuda::ptx 命名空间，可提供直接映射到 PTX 指令的函数。这有助于在 C++ 应用中轻松使用特定的 PTX 指令。有关 cuda::ptx 命名空间的更多信息，请参阅 libcu++ 文档 。此外，NVIDIA 还提供了有关 在 C++ 代码中直接内联 PTX 的文档 。

总结

PTX 是一种虚拟机 ISA，可视为 CUDA GPU 计算平台的汇编语言。PTX 是 CUDA GPU 计算平台的重要组成部分。高级语言编译为 PTX，然后 PTX 在编译时或运行时编译为二进制代码。

通过在其二进制文件中嵌入 PTX 代码，应用和库可以在单个二进制文件中实现跨代兼容性。此外，适用于其他编程或领域特定语言的编译器可以编译为 PTX，然后使用 ptxas 或 JIT 编译来生成能够在 NVIDIA GPU 上运行的二进制代码。

作为开发者，您可以在应用或库中加入 PTX，从而更大限度地提高 GPU 代码的兼容性。

致谢

感谢 Rob Armstrong 和 Jake Hemstad 为此博文做出的贡献。