最早進入 NVIDIA GPU CUDA 平臺的架構設計決策之一是支持 GPU 代碼的向后兼容性。這種設計意味著,新 GPU 應該能夠運行為之前的 GPU 編寫的程序,而無需進行修改。它由 CUDA 的兩個基本特性完成:
- NVIDIA Parallel Thread Execution (PTX) 虛擬指令集架構 (ISA)
- 即時 (JIT) 在運行時編譯 PTX 代碼的 NVIDIA 驅動程序
PTX 是面向 NVIDIA GPU 的虛擬 ISA。您可以將其想象成組裝代碼,但它不限于特定的物理芯片硬件架構,其設計足夠通用,可以與未來的 GPU 架構兼容。
自 NVIDIA 創建 CUDA 平臺使開發者能夠為 GPU 編寫通用程序以來,PTX 一直是 CUDA 不可或缺的一部分。為之前的 GPU 構建的 PTX 代碼可以由當前的驅動進行 JIT 編譯,并在當前的 GPU 上運行,無需修改。
舉個例子。這是一段簡單的代碼,可打印 GPU 名稱和計算能力,還可從 GPU 內核內部打印 hello。
#include <stdio.h> #include <iostream> __global__ void printfKernel() { printf ( ">>>>>>>>>>>>>>>>>>>>\n" ); printf ( "HELLO FROM THREAD %d\n" , threadIdx.x ); printf ( ">>>>>>>>>>>>>>>>>>>>\n" ); } int main( int argc, char ** argv) { // Query and display device properties cudaDeviceProp deviceProp; cudaGetDeviceProperties(&deviceProp, 0); std::cout << deviceProp.name << std::endl; std::cout << "Compute Capability: " << deviceProp.major << "." << deviceProp.minor << std::endl; printfKernel<<<1,1>>>(); cudaDeviceSynchronize(); std::cout << "End Program" << std::endl; return 0; } |
當我們使用 CUDA 12.8 編譯此代碼并在配備 NVIDIA RTX 4000 Ada 的系統上運行時,我們會得到以下結果:
$ nvcc -o x.device_info device_info.cu $ . /x .device_info NVIDIA RTX 4000 Ada Generation Compute Capability: 8.9 >>>>>>>>>>>>>>>>>>>> HELLO FROM THREAD 0 >>>>>>>>>>>>>>>>>>>> End Program |
由于我們沒有為 NVCC 指定任何編譯器標志,因此它使用此版本編譯器支持的最低 PTX 目標。您可以使用 cuobjdump
檢查可執行文件,以查看代碼中的 PTX 架構和 CUDA 二進制 (cubin) 架構 (為簡潔起見,會對輸出進行截圖):
$ cuobjdump x.device_info Fatbin elf code: ================ arch = sm_52 >>> snipped <<< Fatbin ptx code: ================ arch = sm_52 >>> snipped <<< |
您可以看到 ELF
(即 binary) 和 PTX
。當您看到這樣的輸出時,表示 cubin 和 PTX 都嵌入到目標文件中。架構為 sm_52
,即 Compute Capability (CC) 5.2。CC 由數字 X.Y 表示,其中 X 是主要修訂版本號,Y 是次要修訂版本號。
返回示例。GPU 為 CC 8.9,如運行代碼時的打印輸出所示,那么此代碼如何在此 GPU 上運行?
這就是 JIT 編譯發揮作用的地方。CUDA 驅動程序 JIT 編譯 PTX 以在 CC 8.9 GPU 上運行。只要您的代碼包含由等同于或更早于 GPU 架構的架構生成的 PTX,您的代碼就能正常運行。
您可以通過稍微更改編譯器標志來驗證這一點。添加參數 -gencode arch=compute_75,code=compute_75
。這將告知 NVCC 您希望它為您的應用程序構建版本為 compute_75
(計算能力 7.5) 的 PTX,然后將該 PTX 放入可執行文件并使用 cuobjdump
進行驗證。有關 NVCC 如何構建 PTX 和二進制代碼的更多信息,請參閱 理解 PTX (CUDA GPU 計算的匯編語言) 中的圖 1。
您可以看到它正常運行。
$ nvcc -gencode arch=compute_75,code=compute_75 -o x.device_info device_info.cu $ . /x .device_info NVIDIA RTX 4000 Ada Generation Compute Capability: 8.9 >>>>>>>>>>>>>>>>>>>> HELLO FROM THREAD 0 >>>>>>>>>>>>>>>>>>>> End Program |
現在,如果您將 code=compute_75
更改為 code=sm_75
,這將告知 NVCC 構建與之前 (arch=compute_75
) 相同的 PTX。但是,NVCC 應將 PTX 編譯為 SM_75
的 cubin,并將該 cubin 放入可執行文件中,而不是將其保留在可執行文件中用于 JIT 編譯。同樣,您可以使用 cuobjdump
進行驗證。結果如下:
$ nvcc -gencode arch=compute_75,code=sm_75 -o x.device_info device_info.cu $ . /x .device_info NVIDIA RTX 4000 Ada Generation Compute Capability: 8.9 End Program |
如果仔細觀察,您會發現 “HELLO FROM THREAD 0”
未打印。我們省略了所有錯誤檢查代碼,以使代碼示例更清晰。
如果我們像在真實代碼中一樣包含錯誤檢查,您將看到 GPU 內核未啟動,并且返回的錯誤消息是 “No kernel image is available for execution on the device”
。這意味著應用程序中沒有與此 CC 8.9 設備兼容的內核代碼,因此內核從未啟動。
直至 CC 8.9 (含 CC 8.9) 并支持 CUDA 的所有 GPU (Tegra 除外, 因為它們遵循不同的規則 ) 應遵循的經驗法則如下:
- PTX 兼容性 :具有某個 CC 的 PTX 的任何代碼都將在該 CC 的 GPU 以及具有后續 CC 的任何 GPU 上運行。
- “ Cubin 兼容性:具有特定 CC 的 cubin 的任何代碼都將在該 CC 的 GPU 以及具有相同主要功能的任何后續 GPU 上運行。 例如,使用 CC 8.6 的 GPU 可以運行為 CC 8.0 構建的 cubin。但事實并非如此。如果您為 CC 8.6 構建 cubin,則它僅在 CC 8.6 及更高版本上運行,而不是在 8.0 上運行。”
NVIDIA Hopper 中引入的架構特定功能集
從 NVIDIA Hopper 架構 (CC 9.0) 開始,NVIDIA 推出了一套高度專業化的小型功能集,這些功能被稱為 特定架構 ,只能保證在特定的目標架構上存在。其中大多數功能與 Tensor Cores 的使用有關。
要使用這些功能,您必須在應用中嵌入 PTX 或 cubin 代碼,使用 compute_90a
標志表示 PTX,或 sm_90a
標志表示編譯中的 cubin。使用 a
后綴構建特定于架構的目標時,PTX 或 cubin 代碼無法向前兼容任何未來的 GPU 架構。
例如,您使用以下 NVCC 行編譯 CUDA 核函數:
$ nvcc -gencode arch=compute_90a,code=sm_90a -c kernel.cu |
在本例中,您的代碼僅在 CC 9.0 的設備上加載和運行。在使用特定于架構的 a
后綴時,PTX 或 cubin 都無法向前兼容。
NVIDIA Blackwell 中引入的家族特定功能集
從 NVIDIA Blackwell 架構和 CUDA 12.9 開始,我們引入了一類新功能: 特定于系列的功能 。
“該系列的特定功能類似于架構特定的功能,不同之處在于它們由具有多個次要計算能力的設備提供支持。一個系列的所有設備共享相同的主要計算能力版本。該系列的特定功能可確保在同一 GPU 系列中提供,其中包括具有相同主要計算能力和較高次要計算能力的后續 GPU。”
特定于系列的編譯器目標類似于特定于架構的目標,但您可以使用 f
后綴,而不是使用 a
后綴的編譯器目標。
如需詳細了解同一系列中的 GPU,請參閱 編程指南 和 CUDA 計算能力 頁面。如需詳細了解該系列特定目標中包含的功能,請參閱 PTX ISA 中的表格。
例如,您使用以下 NVCC 行編譯 CUDA 核函數,該行可調用系列特定的代碼生成目標:
$ nvcc -gencode arch=compute_100f,code=sm_100 -c kernel.cu |
在本例中,您為 sm_100f
系列生成特定于架構的 cubin 代碼,并且您的代碼將僅在具有 compute capability 10.x 的設備上運行。
目前,這是具有 10.0 和 10.3 計算能力的 GPU。如果引入具有 10.x 計算能力的新 GPU,這些 GPU 上的代碼也將兼容,因為它們屬于 sm_100f
系列。在本例中,code=sm_100
和 code=sm_100f
是彼此的別名,并將生成將在 sm_100f 系列設備上運行的相同 cubin。
在 NVCC 中思考這些不同的特征集的方式如下:
- 無后綴:您的 PTX 或 cubin 兼容性一如既往。
- 后綴 f: 無論您是停留在 PTX 還是從該代碼生成 cubin,該代碼都兼容在具有相同主要計算能力版本以及具有相同或更高次要計算能力版本的 GPU 設備上運行。
-
a
后綴: 代碼僅在該特定 CC 的 GPU 上運行,不得在其他 GPU 上運行。
開發者指南
現在,我們已經解釋了如何使用 NVCC 構建架構和特定系列的代碼目標,我們希望為您在構建應用程序時應該執行的操作提供建議。
通常,您應構建能夠在盡可能多的架構上運行的代碼。只要您沒有使用架構或系列特定功能,就不必在應用中包含架構或系列特定目標,而且您可以繼續像往常一樣構建代碼。即使您使用的是使用架構或系列特定功能的庫,只要這些庫以二進制形式分發,它們也能正常運行。
那么,您何時需要使用 family 或特定于 architecture 的編譯器目標?
如前所述,在使用主要與 Tensor Cores 相關的功能時,特別是通過 PTX 對 Tensor Cores 進行編程時,會使用這些目標。如果您直接編寫 PTX 并使用系列或架構特定的功能,則必須分別使用 f
或 a
標志構建代碼,具體取決于您使用的 PTX 指令是否在 f
特征集中,或者它們是否僅在 a
特征集中。
如果您希望在不同 CC 的 GPU 之間實現可移植性,則必須在代碼中包含適當的防護措施,以確保在不具備這些功能的不同 GPU 架構上運行時,有可用的備用代碼路徑。根據您正在使用的系列和架構特定功能,使用以下宏控制代碼路徑:
__CUDA_ARCH_FAMILY_SPECIFIC__
__CUDA_ARCH_SPECIFIC__
這些宏的定義類似于 __CUDA_ARCH__
。有關更多信息,請參閱《 CUDA Programming Guide 》。
例如,如果您正在構建應用程序并使用 CUTLASS 等頭文件庫或任何包含 CUTLASS 的庫(例如 cuBLASDx ),并且您在 CC 9.0 (NVIDIA Hopper)或更高版本上運行應用程序,則應針對將運行代碼的 GPU 設備構建特定于架構的目標。
CUTLASS 專為實現高性能而設計,具有特殊的代碼路徑,可使用特定于架構的功能來更大限度地提高性能。這些庫在內部已經有后備路徑,可以與其他架構完全兼容。
換言之,如果您使用的是庫,則無需擔心使用宏的后備路徑。
付諸實踐
現在,我們已經討論了架構和系列特定目標,以及何時使用它們,我們將把所有內容放在一起。
一般情況
首先要確定的是,您的代碼使用的是架構還是系列特定功能。您可能知道自己是否在使用這些功能,因為您要直接編寫 PTX,或者包含像 CUTLASS 這樣的頭文件庫。如果不是這樣,對于大多數開發者來說都是如此,那么構建應用程序就像以前一樣。
為了提供最佳性能和未來兼容性,通常的指導是為您知道代碼將運行的每個架構構建 binary code。這可提供最佳性能。
您還應嵌入適用于最新架構的 PTX,以提供出色的未來兼容性。例如,您可能知道自己的代碼將在 CC 8.0、9.0 和 10.0 的設備上運行。以下代碼示例展示了如何為這些架構編譯二進制文件,以及如何為將來的兼容性編譯 CC 10.0 PTX。
$ nvcc -gencode arch=compute_80,code=sm_80 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_100,code=sm_100 -gencode arch=compute_100,code=compute_100 -c kernels.cu |
系列特定功能
如果您選擇使用在不同架構中不可移植的特定功能來優化代碼,則應首先確定這些功能是否屬于該系列的特定功能集。
如果是,那么您可以使用后綴 f
構建目標,并且您將與該系列兼容。如果您希望可移植到系列以外的 GPU,則必須為任何使用系列特定功能的代碼添加備用代碼路徑。
通常,這是通過應用中的條件宏保護特定系列的代碼來實現的。在擴展先前的示例并為 CC 10.0 添加系列特定功能后,您對 NVCC 的使用可能類似于以下代碼示例:
$ nvcc -gencode arch=compute_80,code=sm_80 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_100f,code=sm_100 -gencode arch=compute_100,code=compute_100 -c kernels.cu |
這讓您的代碼能夠在 CC 8.0、9.0 和 10.0 的設備上運行,并具有 10.0 的系列特定功能。通過嵌入式 PTX,您的代碼也將在未來的設備上運行。
使用特定于系列的功能的另一種可能情況是,您知道您的應用必須利用這些功能,而應用僅設計為在該系列的設備上運行。例如,如果您將代碼設計為僅使用 100f
系列的功能,并且只想在此系列的設備上運行,則應用程序構建類似于以下代碼示例:
$ nvcc -gencode arch=compute_100f,code=sm_100 -c kernels.cu |
在這種情況下,您的代碼只能在此系列的設備中進行移植。
架構特定功能
如果您已確定系列特定功能不足以滿足您的應用需求,并且必須使用架構特定功能集中的功能,則必須使用 a
標志進行構建。
與使用 f
進行構建的情況類似,您必須使用應用內的條件宏保護代碼,從而確定必須在應用中構建何種代碼可移植性。為了獲得與上一個示例相同的可移植性,請按照以下代碼示例構建代碼:
$ nvcc -gencode arch=compute_80,code=sm_80 -gencode arch=compute_90,code=sm_90 -gencode arch=compute_100a,code=sm_100a -gencode arch=compute_100,code=compute_100 -c kernels.cu |
您的代碼將具有相同的兼容性,可在 CC 8.0、CC 9.0、CC 10.0 及更高版本上運行,
與僅為特定系列設計應用類似,您也可以選擇針對特定架構進行設計和優化。如果您設計和編寫的應用程序使用特定于架構的功能,并且知道它不必在任何其他 GPU 上運行,則可以構建類似于以下代碼示例的應用程序:
$ nvcc -gencode arch=compute_100a,code=sm_100a -c kernels.cu |
您的應用程序將僅在 CC 10.0 上運行,不兼容任何其他 GPU。
總結
總結一下,以下簡單流程說明了您應如何考慮構建您的代碼:
- 您是直接編寫 PTX,還是直接調用 CUTLASS 等庫?否則,您不需要添加
f
或a
標志。像往常一樣構建代碼。 - 如果您要編寫 PTX,或包含僅包含頭文件的庫,則需要確定該庫是否使用架構或系列特定特征集中的特征,如果是,則確定要使用的
f
或a
標志。您需要參閱該庫的文檔,以確定為您的架構構建的最佳方式。例如, CUTLASS 構建說明指定在構建 CC 9.0 和 10.0 的設備時使用a
標志
本文介紹了很多內容,向您展示了如何構建可以使用架構和系列特定功能的代碼。我們想明確的是,在 NVCC 編譯行中使用 a
和 f
后綴并不是一種神奇的優化技術。要具體使用這些功能,您必須直接編寫 PTX 或調用具有此功能的庫。
立即下載 CUDA 12.9 ,開始在 Blackwell 上的代碼中使用家族和架構特定的功能。
致謝
感謝以下 NVIDIA 貢獻者:Cody Addison、Vyas Venkataraman、Rob Armstrong、Girish Bharambe 和 Mridula Prakash。
?