4.8. 延迟加载

4.8.3.1. CUDA 运行时版本要求

从 CUDA 运行时版本 11.7 开始提供延迟加载。由于 CUDA 运行时通常以静态方式链接到程序和库中，因此只有来自 CUDA 11.7 及更高版本工具包或使用这些工具包编译的程序和库才能受益于延迟加载。使用较旧 CUDA 运行时版本编译的库会立即加载所有模块。

4.8.3.2. CUDA 驱动程序版本要求

延迟加载要求驱动程序版本为 515 或更高。对于低于 515 的驱动程序版本，即使使用 CUDA 工具包 11.7 或更高版本，也无法使用延迟加载。

4.8.3.3. 编译器要求

延迟加载不需要任何编译器支持。使用 11.7 之前版本的编译器编译的 SASS 和 PTX 都可以在启用延迟加载的情况下加载，并且能够充分受益于该功能。不过，如上所述，仍然需要版本为 11.7 或更高的 CUDA 运行时。

4.8.3.4. 核函数要求

延迟加载不会影响包含托管变量的模块，这些模块仍会立即加载。

4.8.4.1. 启用与禁用

将CUDA_MODULE_LOADING环境变量设置为LAZY即可启用延迟加载。将CUDA_MODULE_LOADING环境变量设置为EAGER可以禁用延迟加载。从 CUDA 12.3 起，所有平台都默认启用延迟加载。

4.8.4.2. 在运行时检查是否启用了延迟加载

CUDA 驱动程序 API 中的cuModuleGetLoadingModeAPI 可用于确定是否启用了延迟加载。注意，运行此函数之前必须初始化 CUDA。下面的代码片段展示了用法示例。

#include "<cuda.h>"
#include "<assert.h>"
#include "<iostream>"

int main() {
        CUmoduleLoadingMode mode;

        assert(CUDA_SUCCESS == cuInit(0));
        assert(CUDA_SUCCESS == cuModuleGetLoadingMode(&mode));

        std::cout << "CUDA Module Loading Mode is " << ((mode == CU_MODULE_LAZY_LOADING) ? "lazy" : "eager") << std::endl;

        return 0;
}

4.8.4.3. 在运行时强制立即加载模块

核函数和变量会自动加载，无需进行任何显式加载。即使不执行核函数，也可以通过以下方式显式加载它们：

• cuModuleGetFunction()函数会使模块加载到设备内存中

• cudaFuncGetAttributes()函数会使核函数加载到设备内存中

4.8.5.1. 对核函数并发执行的影响

一些程序错误地假定核函数并发执行是有保证的。如果需要跨核函数同步，但核函数执行已被串行化，则可能发生死锁。要尽量减小延迟加载对核函数并发执行的影响，请执行以下操作：

• 在启动希望并发执行的所有核函数之前，预加载这些核函数；或者

• 使用CUDA_MODULE_LOADING = EAGER运行应用程序，以强制立即加载数据，而不强制立即加载每个函数

4.8.5.2. 大规模内存分配

延迟加载将 CUDA 模块的内存分配从程序初始化时推迟到更接近执行的时间。如果应用程序在启动时分配了全部显存，CUDA 可能在运行时无法为模块分配内存。可能的解决方案：

• 使用cudaMallocAsync()，而不是使用在启动时分配全部显存的分配器

• 留出一些余量，以应对核函数的延迟加载

• 在尝试初始化分配器之前，预加载程序中将使用的所有核函数

4.8.5.3. 对性能测量的影响

延迟加载可能将 CUDA 模块初始化移入测量的执行时间窗口，从而使性能测量产生偏差。为避免这种情况：

• 在测量之前至少进行一次预热迭代

• 在启动接受基准测试的核函数之前，先预加载该核函数