Parallel Thread Execution（PTX/NVPTX）是一种较底层的并行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。
寄存器
PTX 使用任意大小的处理器寄存器集；编译器的输出几乎是纯粹的静态单赋值形式（SSA），连续的行通常引用连续的寄存器。程序以如下形式的声明开始：
.reg .u32 %r; // 声明 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分别表示线程索引、块维度、块索引和网格维度。
状态空间
加载 ( ld ) 和存储 ( st ) 命令主要操作如下几个不同的状态空间（内存库），例如ld.param 。共有八个状态空间：
; .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 模块
参见
- (SPIR)
参考
外部链接
- [https://docs.nvidia.com/cuda/parallel-thread-execution/ NVIDIA开发者专区的PTX ISA页面]
评论 (0)