浅谈NVIDIA GPU 的PTX

本文通过一个简单的整数加法例子，浅谈 NVIDIA GPU 的 PTX（Parallel Thread Execution）。

为便于理解，下面只介绍 PTX 的基本概念、编译流程和作用，不展开 Tensor Core、Warp 调度、内存模型等复杂细节。

What

先把最容易混淆的一点说清楚：

1
2
3

PTX 代码
    ≠
GPU 最终执行的机器指令

PTX 是 NVIDIA CUDA 平台的一种虚拟指令集架构（Virtual ISA），可以把它理解为面向 NVIDIA GPU 的一种中间汇编语言。 它位于 CUDA C++ 等高级语言与具体 GPU 的机器指令之间。

可以简单理解为：

CUDA C++：告诉 GPU 要完成什么计算

PTX：用接近汇编的指令描述计算过程

GPU 机器指令：针对具体 GPU 架构生成的可执行代码

例如，CUDA C++ 中的整数加法：

1	int c = a + b;

可以在 PTX 中表示为：

1	add.u32 %r3, %r1, %r2;

这里的 add.u32 表示执行 32 位无符号整数加法，%r1、%r2 和 %r3 是 PTX 中的虚拟寄存器。

这只是一个独立的指令示例，并不是完整的 CUDA Kernel。

所以：

PTX 描述 GPU 应该执行的指令，但最终如何映射到具体 GPU 的机器指令，由后续编译阶段决定。

Why

为什么 NVIDIA 不直接让 CUDA C++ 编译成某一代 GPU 的机器指令？

因为 GPU 架构会持续演进。不同代 GPU 的底层机器指令并不完全相同；如果程序只包含面向某一代 GPU 的二进制代码，换到不兼容的 GPU 时，就可能需要重新编译。

PTX 在高级语言与具体硬件之间提供了一个相对稳定的中间层：

CUDA C++ / 其他编程语言
            ↓
           PTX
            ↓
    针对目标 GPU 编译
            ↓
       GPU 机器代码

这样，编译器不必让每种上层语言都直接适配每一代 GPU；其他编译器也可以生成 PTX，再借助 NVIDIA 的工具链生成可执行代码。

因此，PTX 的目的可以概括为：

在高级编程语言与具体 GPU 架构之间提供一个可编译、可移植的中间表示。

基础 Example

假设我们编写了一个 CUDA Kernel，用来计算两个数组的逐元素相加：

__global__ void add(const float *a,
                    const float *b,
                    float *c)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

每个 GPU 线程负责计算一个位置的结果。

从 CUDA C++ 到 PTX

编译器可以将其中的浮点加法转换成类似这样的 PTX 指令：

ld.global.f32 %f1, [%rd1];
ld.global.f32 %f2, [%rd2];
add.f32       %f3, %f1, %f2;
st.global.f32 [%rd3], %f3;

这是为了说明思路的示意片段，省略了地址计算、参数准备等代码；真实生成的 PTX 会随编译器和编译选项变化。

可以这样理解：

ld.global.f32：从 GPU 全局内存读取一个 float

add.f32：      将两个 float 相加

st.global.f32：把结果写回 GPU 全局内存

PTX 还不是最后执行的代码

PTX 通常还会继续编译，生成针对特定 GPU 架构的二进制代码，通常称为 cubin。其底层机器指令常被称为 SASS。

CUDA C++ Kernel
      ↓
PTX 虚拟指令
      ↓
面向目标 GPU 编译
      ↓
cubin / GPU 机器指令
      ↓
GPU 执行

因此，看到 PTX 中有一条 add.f32，并不代表 GPU 最终一定只执行一条与它一一对应的硬件指令。后端编译器还会针对目标 GPU 进行优化和指令选择。

How

只看三个关键点。

1. 编译阶段生成 PTX

使用 nvcc 编译 CUDA 程序时，设备端代码可以先转换为 PTX；之后，编译工具链可以把 PTX 编译为特定 GPU 架构的二进制代码。

1	CUDA C++ → PTX → cubin

开发者也可以显式生成 PTX，或在需要时直接编写 PTX。但对于大多数 CUDA 程序，通常不需要手写 PTX。

2. 程序可以同时携带 cubin 和 PTX

CUDA 程序中的 GPU 代码通常可以打包在 fatbin 中。它可以包含面向不同 GPU 架构的 cubin，也可以包含 PTX。

应用程序
    └── GPU 代码（fatbin）
          ├── 针对某些 GPU 的 cubin
          └── PTX

运行时，驱动会选择适合当前 GPU 的可执行代码。如果有可用且兼容的 cubin，就可以直接使用；如果需要从 PTX 生成目标 GPU 的代码，驱动可以执行 JIT（Just-In-Time，即时编译）。

3. PTX 有助于支持更新的 GPU

假设应用编译时还没有针对某一款新 GPU 生成 cubin，但应用中包含兼容的 PTX，那么新 GPU 的驱动就可能将 PTX 即时编译为适合该 GPU 的机器代码。

应用中保留 PTX
       ↓
在新 GPU 上运行
       ↓
驱动将 PTX JIT 编译为目标代码
       ↓
执行 Kernel

这为未来 GPU 提供了前向兼容的途径，不必一定为每一代新 GPU 重新发布一份应用。实际能否运行仍取决于 PTX 版本、目标计算能力、使用的指令特性以及驱动支持情况。

一个简单的取舍

PTX 带来了灵活性，但也有需要注意的地方。

包含可用的 cubin：
    通常可以直接加载目标机器码，启动时开销较小

需要 JIT 编译 PTX：
    可以为目标 GPU 生成代码，但首次加载时可能增加编译开销

驱动通常会缓存 JIT 生成的二进制代码，以减少后续运行重复编译的开销。

另外，PTX 是虚拟 ISA，不是性能保证。同一段 PTX 在不同 GPU 上可能被编译成不同的机器指令；最终性能还取决于编译器优化、目标架构和程序本身。

记忆图

          CUDA C++ / 其他语言
                   │
                   ▼
              PTX 虚拟 ISA
                   │
       ┌───────────┴───────────┐
       ▼                       ▼
离线编译为 cubin          运行时 JIT 编译
       │                       │
       └───────────┬───────────┘
                   ▼
          目标 GPU 的机器代码
                   │
                   ▼
              GPU 执行

可以简单记成：

PTX：面向 NVIDIA GPU 的虚拟指令集

cubin / SASS：更接近具体 GPU 的可执行代码

JIT：运行时将 PTX 编译成目标 GPU 的代码

一句话概括：

PTX 是 CUDA 生态中的中间汇编层，让程序可以先用相对稳定的虚拟指令描述 GPU 计算，再针对实际运行的 GPU 生成机器代码。

参考资料：

文章目录