過去，我們使用 nvcc 等離線工具將 GPU 設備代碼與應用程序一起編譯。在這種情況下，GPU 設備代碼在 CUDA 運行時內部進行管理。然后，您可以使用 <<<>>> 啟動內核，并且 CUDA 運行時可確保啟動所調用的內核。
但是，在某些情況下，需要動態編譯和加載 GPU 設備代碼。本文介紹了使用 CUDA 運行時實現此目標的方法，同時還展示了在 CUDA 驅動程序和 CUDA 運行時核句柄之間實現互操作性的方法。
在 CUDA 12.0 中，NVIDIA 通過 CUDA 驅動引入了 cuLibraryLoad APIs。這些 APIs 使您能夠以與上下文無關的方式動態選擇和加載 GPU 設備代碼。有關更多信息，請參閱 CUDA Context-Independent Module Loading 。
現在，我們將此功能擴展為通過 CUDA 運行時加載動態 GPU 設備代碼，并使用一組新的庫管理 API 來擴展 CUDA 驅動程序 API，這與其他 CUDA 運行時 API 類似。
動態 GPU 設備代碼加載的優勢
啟用動態 GPU 設備代碼加載具有以下優勢：
- 顯式控制正在加載的 GPU 設備代碼，以防該代碼與加載編譯單元分開進行修改。
- 通過加載 API 選項來控制 GPU 設備代碼的加載時間以及如何加載的選項。
- 使用 NVRTC 等其他 CUDA 工具包組件進行動態編譯，以生成 GPU 設備代碼模組。
- 使用 nvJitLink 等其他 CUDA 工具包組件進行動態選擇性 GPU 設備代碼鏈接，以實現鏈路時間優化。
- 使用 nvcc 編譯且必須執行動態 GPU 設備代碼加載的僅包含報文頭的庫可以通過這些更改關聯到 CUDA 運行時。
我們將在本帖中詳細討論每項好處。
CUDA 運行時中的靜態加載
CUDA 運行時會維護有關初始化期間加載的 GPU 設備代碼的狀態。GPU 設備代碼模組由編譯內容以及與編譯工具 (如 nvcc) 關聯的內容決定。在初始化期間，CUDA 運行時會加載這些 GPU 設備代碼模組，您可以隱式地與它們交互，如下例所示：
main.cu:#include <stdio.h>__global__ void helloWorld() { printf(“Hello from the GPU!\n”); }int main(int argc, char *argv[]) { cudaSetDevice(0); helloWorld<<<1,1,1>>>(); return cudaDeviceSynchronize();} |
此簡化示例使用 nvcc 編譯時，可使用適當的 GPU 設備代碼模塊創建可執行文件，使 CUDA 運行時能夠在 GPU 上運行和執行 helloWorld 核函數。
CUDA 驅動程序中的動態加載
CUDA 驅動程序要求您動態加載要執行的 GPU 設備代碼，并管理更多狀態，例如 CUDA 運行時為您自動管理的 CUDA 上下文。我們會將一個類似的示例分解為兩個具有單獨編譯軌跡的文件。有關各種編譯軌跡的更多信息，請參閱 NVIDIA CUDA Compiler Driver NVCC 。
GPU 的代碼將使用 nvcc 編譯到獨立的 GPU 設備代碼模塊中，例如 .fatbin、.cubin 或獨立的 PTX 文件 (在本示例中即為 device.fatbin)。
然后，您將擁有一個主源文件，用于使用和管理此 .fatbin 文件，其中包含已編譯并鏈接的 GPU 設備代碼模塊。部分主源文件如下所示，未進行錯誤檢查，便于閱讀：
main.c:#include <cuda.h>int main(int argc, char *argv[]) { … cuDeviceGet(&dev, 0); cuDevicePrimaryCtxRetain(&ctx, dev); cuCtxPushCurrent(ctx); cuLibraryLoadFromFile(&library, “device.fatbin”, NULL, NULL, 0, NULL, NULL, 0); cuLibraryGetKernel(&kernel, library, “helloWorld”); cuLaunchKernel((CUfunction)kernel, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, NULL); cuCtxSynchronize(); cuLibraryUnload(library); cuDevicePrimaryCtxRelease(dev); return 0;} |
前面列出的動態加載的優勢已擴展到 CUDA 運行時，并在 用例啟用部分 中進行了進一步描述。
CUDA 運行時中的動態加載
通過更改 CUDA 以支持 CUDA 運行時中的動態加載，我們為 CUDA 運行時提供了動態加載 GPU 設備代碼的靈活性。這意味著前面的示例可以壓縮為以下代碼。這消除了驅動示例所需的顯式 CUDA 上下文管理開銷。此處顯示了更新后主源文件的一部分，未進行錯誤檢查，便于閱讀：
main.cu:#include <cuda_runtime_api.h>int main(int argc, char *argv[]) { … cudaLibraryLoadFromFile(&library, “device.fatbin”, NULL, NULL, 0, NULL, NULL, 0); cudaLibraryGetKernel(&kernel, library, “helloWorld”); cudaLaunchKernel((const void*)kernel, 1, 1, NULL, 0, NULL); cudaDeviceSynchronize(); cudaLibraryUnload(library); return 0;} |
用例支持?
這實現了哪些用例？以下是一些以前不可能實現的示例：
- 純 CUDA 運行時 API 使用情況
- CUDA 驅動程序和 CUDA 運行時之間類型的可互換性
- 處理 CUDA 運行時實例之間的共享
純 CUDA 運行時 API 使用情況?
到目前為止，加載的所有動態 GPU 設備代碼模組都需要驅動 API。如果其他庫或應用可以使用 NVRTC 進行編譯，或者使用 nvJitLink 動態關聯 GPU 設備代碼，則需要驅動加載生成的輸出。
借助新的 CUDA 運行時動態加載 API，這些動態輸出的加載、管理和使用完全可以通過 CUDA 運行時完成。
以下是根據前面提到的 NVRTC 文檔修改的示例：已更新的 NVRTC SAXPY 示例，可使用新的 CUDA 運行時 API。
Current NVRTC SAXPY Example Snippet// Load the generated PTX and get a handle to the SAXPY kernel.CUdevice cuDevice;CUcontext context;CUmodule module;CUfunction kernel;CUDA_SAFE_CALL(cuInit(0));CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, ptx, 0, 0, 0));CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, “saxpy”));…//Execute SAXPYvoid *args[] = {&a, &dX, &dY, &dOut, &n};CUDA_SAFE_CALL(cuLaunchKernel(kernel, NUM_BLOCKS, 1, 1, // grid dim NUM_THREADS, 1, 1, // block dim 0, // shmem NULL, // stream args, 0)); // arguments |
Updated NVRTC SAXPY Example Snippet// Load the generated PTX and get a handle to the SAXPY kernel.cudaLibrary_t library;cudaKernel_t kernel;CUDART_SAFE_CALL(cudaLibraryLoadData(&library, ptx, 0,0,0,0,0,0));CUDART_SAFE_CALL(cudaLibraryGetKernel(&kernel, library, “saxpy”));…//Execute SAXPYvoid *args[] = {&a, &dX, &dY, &dOut, &n};CUDART_SAFE_CALL(cudaLaunchKernel((void*)kernel, NUM_BLOCKS, // grid dim NUM_THREADS, // block dim args, // arguments 0, // shmem NULL)); // stream |
另一個好處是，在此之前，使用 nvcc 編譯且必須進行動態 GPU 設備代碼加載的純頭庫會增加用戶在編譯時鏈接到 CUDA 驅動程序的要求。現在，通過使用與 nvcc 關聯的 CUDA 運行時，這些僅包含標頭的庫可以不需要顯式鏈接到 CUDA 驅動程序。
這也意味著可以使用 CUDA 運行時將兩套不同的代碼（一套用于加載 CUDA 運行時 GPU 設備代碼的歷史靜態方法，另一套用于加載 CUDA 驅動程序的動態 GPU 設備代碼）融合為一組代碼。
CUDA 驅動程序和 CUDA 運行時之間類型的可互換性
以前，與許多其他句柄（例如 CUDA 流和 CUDA 事件）一樣，內核句柄無法在 CUDA 運行時和 CUDA 驅動程序之間互換。
以前，您無法在 CUDA 運行時環境和 CUDA 驅動程序之間交換內核句柄，而現在 cudaKernel_t 和 CUkernel (以及 cudaLibrary_t 和 CUlibrary) 可互換。要使用 CUDA 運行時 API 進行加載，但使用 CUDA 驅動程序 API 啟動或設置內核屬性，您可以在這些類型之間轉換。
現在，要執行動態 GPU 設備代碼加載，您不必僅使用 CUDA 驅動程序 API。您可以使用一組 API，并且僅在 CUDA 驅動程序和 CUDA 運行時類型之間轉換。
處理 CUDA 運行時實例之間的共享
假設有兩個理論庫，即庫 A 和庫 B，每個庫都關聯到各自的靜態 CUDA 運行時。
歷史 CUDA 運行時加載的隱式特性無法在多個 CUDA 運行時實例之間共享 CUDA 內核句柄。在這種情況下，無法共享每個庫的內核句柄。
現在，借助 CUDA 運行時 API cudaGetKernel，您可以獲得任何內核的句柄，并將其傳遞給另一個 CUDA 運行時實例。如果需要在兩個庫之間共享 CUDA 內核，則庫 A 可以調用 cudaGetKernel 并將句柄傳遞給庫 B。這樣做的潛在好處是增加庫之間的代碼共享量，并減少每個庫包含自己的內核實現的需求。
在以下代碼示例中，libmatrix_mul.cu 使用 CUDA 運行時 API 中的新動態加載，libvector_add.cu 使用 CUDA 運行時中的傳統隱式加載，但利用新的 cudaGetKernel API 獲取可共享 CUDA 核函數的句柄。
在這兩種情況下，您都可以將句柄傳遞給 cudaKernel_t 第三個獨立庫 libcommon，以啟動并使用 cudaKernel_t，即使它們關聯到自己的靜態 CUDA 運行時實例，也可以實現這一點。
// matrix_mul.cu - using dynamic shared handlesvoid matrix_mul() { cudaLibrary_t lib; cudaKernel_t kern; cudaLibraryLoadData(&lib, ptx, …); // ptx from nvrtc cudaLibraryGetKernel(&kern, lib, “matrixMul”); libcommon.foo(kern);}// vector_add.cu - using implicit shared handles__global__ void vectorAdd() { … }void vector_add() { cudaGetKernel(&kern, vectorAdd); libcommon.foo(kern);}// libcommon.cu - takes a shareable kernel handlevoid foo(cudaKernel_t kern) { cudaLaunchKernel(kern, ...);} |
此示例意義重大，但它展示了通過在彼此所需的內核之間進行重復數據刪除，庫可以節省主機和 GPU 顯存空間。
開始使用 CUDA 運行時動態加載
在本文中，我們介紹了新的 CUDA 運行時 API，這些 API 能夠加載 GPU 設備代碼。當僅需要 CUDA 運行時 API 時，這是一種在 GPU 上加載和執行設備代碼的更簡單方法。
要開始使用這些 API，請從 CUDA 工具包 下載 CUDA 工具包版本 12.8 或更高版本。有關 cudaLibrary* 和 cudaKernel* API 的更多信息，請參閱 CUDA 運行時 API 文檔 。
?
