大多數 CUDA 開發人員都熟悉cuModuleLoad API 及其對應的 API ，用于將包含設備代碼的模塊加載到 CUDA context 中。在大多數情況下，您希望在所有設備上加載相同的設備代碼。這需要將設備代碼顯式加載到每個 CUDA 上下文中。此外，不控制上下文創建和銷毀的庫和框架必須跟蹤它們，以顯式加載和卸載模塊。
本文討論了 CUDA 12.0 中引入的上下文無關加載，它解決了這些問題。
上下文相關加載
傳統上，模塊加載總是與 CUDA 上下文相關聯。下面的代碼示例顯示了將相同的設備代碼加載到兩個設備中，然后在它們上啟動內核的傳統方法。
// Device 0
cuDeviceGet(&device0, 0);
cuDevicePrimaryCtxRetain(&ctx0, device0);
cuModuleLoad(&module0, “myModule.cubin”);
// Device 1
cuDeviceGet(&device1, 1);
cuDevicePrimaryCtxRetain(&ctx1, device1);
cuModuleLoad(&module1, “myModule.cubin”);在每個設備上啟動內核需要檢索每個模塊CUfunction，如以下代碼示例所示：
// Device 0
cuModuleGetFuntion(&function0, module0, “myKernel”);
cuLaunchKernel(function0, …);
// Device 1
cuModuleGetFuntion(&function1, module1, “myKernel”);
cuLaunchKernel(function1, …);這增加了應用程序中的代碼復雜性，因為您必須檢索和跟蹤每個上下文和每個模塊類型。您還必須使用cuModuleUnload API 顯式卸載每個模塊。
當庫或框架主要使用 CUDA 驅動程序 API 來加載自己的模塊時，問題就更加嚴重了。他們可能無法完全控制應用程序擁有的上下文的生命周期。
// Application code
libraryInitialize();
cuDevicePrimaryCtxRetain(&ctx0, device0);
libraryFunc();
cuDevicePrimaryCtxRetain(&ctx0, device1);
libraryFunc();
libraryDeinitialize();
// Library code
libraryInitialize() {
map<CUcontext, CUmodule> moduleContextMap;
}
libraryFunc() {
cuCtxGetCurrent(&ctx);
if (!moduleContextMap.contains(ctx)){
cuModuleLoad(&module, “myModule.cubin”);
moduleContextMap[ctx] = module;
}
else {
module = moduleContextMap[ctx];
}
cuModuleGetFuntion(&function, module, “myKernel”);
cuLaunchKernel(function, …);
}
libraryDeinitialize() {
moduleContextMap.clear();
}在代碼示例中，庫必須檢查新上下文并顯式地將模塊加載到其中。它還必須保持狀態，以檢查模塊是否已加載到上下文中。
理想情況下，可以在上下文被破壞后釋放狀態。但是，如果庫無法控制上下文的生命周期，則這是不可能的。
這意味著資源的釋放必須延遲到庫的去初始化。這不僅增加了代碼的復雜性，而且還會導致庫占用資源的時間超過必須的時間，從而可能會阻止應用程序的另一部分使用該內存。
另一種選擇是庫和框架對用戶施加額外的約束，以確保他們對資源分配和清理有足夠的控制。
上下文無關加載
CUDA 12.0 引入了上下文無關的加載，并添加了cuLibrary*和cuKernel* API ，解決了這些問題。通過獨立于上下文的加載，當應用程序創建和銷毀上下文時， CUDA 驅動程序會自動將模塊加載和卸載到每個 CUDA 上下文中。
// Load library
cuLibraryLoadFromFile(&library,“myModule.cubin”, …);
cuLibraryGetKernel(&kernel, library, “myKernel”);
// Launch kernel on the primary context of device 0
cuDevicePrimaryCtxRetain(&ctx0, device0);
cuLaunchKernel((CUkernel)kernel, …);
// Launch kernel on the primary context of device 1
cuDevicePrimaryCtxRetain(&ctx1, device1);
cuLaunchKernel((CUkernel)kernel, …);
// Unload library
cuLibraryUnload(library);如代碼示例所示，cuLibraryLoadFromFile API 負責在創建或初始化上下文時加載模塊。在本例中，這是在cuDevicePrimaryCtxRetain期間完成的。
此外，您現在可以使用上下文無關句柄CUkernel啟動內核，而不必維護每個上下文CUfunction。cuLibraryGetKernel檢索設備函數myKernel的上下文無關句柄。然后，通過指定上下文無關句柄CUkernel，可以使用cuLaunchKernel啟動設備功能。 CUDA 驅動程序負責根據此時激活的上下文在適當的上下文中啟動設備功能。
庫和框架現在可以分別在初始化和去初始化期間簡單地加載和卸載模塊。
// Application code
libraryInitialize();
cuDevicePrimaryCtxRetain(&ctx0, device0);
libraryFunc();
cuDevicePrimaryCtxRetain(&ctx0, device1);
libraryFunc();
libraryDeinitialize();
// Library code
libraryInitialize() {
cuLibraryLoadFromFile(&library,“myModule.cubin”, …);
cuLibraryGetKernel(&kernel, library, “myKernel”);
}
libraryFunc() {
cuLaunchKernel((CUkernel)kernel, …);
}
libraryDeinitialize() {
cuLibraryUnload(library);
}庫不再需要維護和跟蹤每上下文狀態。上下文無關加載的設計使 CUDA 驅動程序能夠跟蹤模塊和上下文，并執行加載和卸載模塊的工作。
訪問__ managed __變量
托管變量可以從設備和主機代碼中引用。例如，可以查詢托管變量的地址，也可以直接從設備或主機函數讀取或寫入該地址。與__device__變量不同，__managed__變量具有創建 CUDA 上下文的生命周期，屬于模塊的cuLibraryGetManaged變量在所有 CUDA 上下文甚至設備中指向相同的內存。
在 CUDA 12.0 之前，無法通過驅動程序 API 檢索在 CUDA 上下文中唯一的托管變量的句柄。 CUDA 12.0 引入了一個新的驅動程序 API cuLibraryGetManaged，這使得在 CUDA 上下文中獲得唯一句柄成為可能。
開始獨立于上下文的加載
在這篇文章中，我們介紹了新的 CUDA 驅動程序 API ，它提供了獨立于 CUDA context 加載設備代碼的能力。我們還討論了啟動內核的上下文無關句柄。與傳統的加載機制相比，它們提供了在 GPU 上加載和執行代碼的更簡單的方式，降低了代碼復雜性，避免了維護每上下文狀態的需要。
要開始使用這些 API ，請下載 CUDA Driver and Toolkit version 12 or higher 。有關 cuLibrary*和cuKernel* API 的更多信息，請參閱 CUDA Driver API 文檔。
?
