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）.