过去，我们使用 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 SAXPY
void *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 SAXPY
void *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 handles
void 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 handle
void 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 文档 。
