大多数 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 文档。
