Parallel Thread Execution
Parallel Thread Execution（PTX/NVPTX[1]）是一种较底层的并行VM和指令集架构，服务于Nvidia的CUDA编程环境。Nvidia CUDA 编译器 nvcc 将 OpenCL C 语言代码或 CUDA C/C++ 语言代码翻译成 PTX 指令（中间语言）。Nvidia显卡驱动包含对应编译器，可以将PTX指令翻译为可执行程序，将程序运行在Nvidia图形处理器上。
clang 和 GNU GCC均能够根据 CUDA、OpenCL C/C++、SYCL C++ 或OpenACC或OpenMP指令生成 PTX。 [1][2][3]
寄存器
[编辑]PTX 使用任意大小的处理器寄存器集；编译器的输出几乎是纯粹的静态单赋值形式（SSA），连续的行通常引用连续的寄存器。程序以如下形式的声明开始：
.reg .u32 %r<335>; // 声明 335 个寄存器 %r0, %r1, ..., %r334，类型为无符号 32 位整数
PTX是一种三参数汇编语言，几乎所有指令都会明确列出其操作的数据类型（包括符号和宽度）。寄存器名称前会加上一个 % 字符，所有常量都是字面常量，例如：
shr.u64 %rd14, %rd12, 32; // shift right an unsigned 64-bit integer from %rd12 by 32 positions, result in %rd14
cvt.u64.u32 %rd142, %r112; // convert an unsigned 32-bit integer to 64-bit
PTX存在谓词寄存器，但着色器模型 1.0 中的编译代码仅在配合分支指令时使用这些寄存器，例如：
@%p14 bra $label; // 跳转至 $label
setp.cc.type指令将谓词寄存器设置为比较两个寄存器的结果，例如：
set.le.u32.u64 %r101, %rd12, %rd28 会将32位寄存器%r101设为0xffffffff（若 %rd12小于%rd28）或0x00000000。
有一些预定义的标识符用于表示伪寄存器。例如， %tid, %ntid, %ctaid和%nctaid分别表示线程索引、块维度、块索引和网格维度。 [4]
状态空间
[编辑]加载 ( ld ) 和存储 ( st ) 命令主要操作如下几个不同的状态空间（内存库），例如ld.param 。共有八个状态空间： [4] [5]
.reg- 寄存器
.sreg- 只读的专用寄存器
.const- 共享只读内存
.global- 全局内存，由所有线程共享
.local- 每个线程私有的本地内存
.param- 传递给Kernel的参数
.shared- 内存块中线程间共享的内存
.tex- 全局纹理内存（已弃用）
共享内存在 PTX 文件头声明：
.shared .align 8 .b8 pbatch_cache[15744]; // 定义 15,744 字节，按 8 字节边界对齐
使用 PTX 编写Kernel需要通过 CUDA 驱动程序 API 显式注册 PTX 模块，这通常比使用 CUDA 运行时 API 和 Nvidia 的 CUDA 编译器 nvcc 更为繁琐。GPU Ocelot 项目(已停止维护)提供了一个 API，用于在 CUDA 运行时 API Kernel中注册 PTX 模块 [6]
参见
[编辑]- 标准便携式中间表示(SPIR)
参考
[编辑]- ^ 1.0 1.1 User Guide for NVPTX Back-end – LLVM 7 documentation. llvm.org. [2026-06-27]. （原始内容存档于2026-05-28）.
- ^ Compiling CUDA with clang — LLVM 22.0.0git documentation. llvm.org. [2026-06-27]. （原始内容存档于2026-07-20）.
- ^ nvptx. GCC Wiki. [2026-06-27]. （原始内容存档于2023-09-30）.
- ^ 4.0 4.1 PTX ISA Version 2.3 (PDF).
- ^ 1. Introduction — PTX ISA 9.0 documentation. docs.nvidia.com. [2025-09-22]. （原始内容存档于2019-12-01）.
- ^ GPUOCelot: A dynamic compilation framework for PTX. github.com. 7 November 2022. （原始内容存档于2023-09-24）.