了解 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 文件。此源文件包含單個核函數 (即并行添加兩個向量的經典示例) ，以及應用程序主函數和輔助函數的骨架。

使用 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 平臺的匯編語言。

計算能力和 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 為此博文做出的貢獻。

?

?