CUDA软件架构中的模块、函数与内存管理

### CUDA 软件架构中的模块、函数与内存管理 #### 1. 上下文状态与模块 在 CUDA 编程中，上下文状态的配置十分关键。`cuCtxSetLimit()` 和 `cuCtxGetLimit()` 函数可用于配置与类 CPU 功能相关的限制，如内核中的 `malloc()` 和 `printf()`。而 `cuCtxSetCacheConfig()` 则用于指定启动内核时首选的缓存配置，例如是为共享内存和 L1 缓存分配 16K 还是 48K。不过这只是一个提示，若内核使用的共享内存超过 16K，则需要将配置设置为 48K 共享内存。此外，上下文状态可被特定内核的状态（如 `cuFuncSetCacheConfig()`）覆盖。这些状态具有上下文范围，因为更改它们的成本较高。 模块是一起加载的代码和相关数据的集合，类似于 Windows 上的 DLL 或 Linux 上的 DSO。CUDA 运行时并不显式支持模块，模块仅在 CUDA 驱动 API 中可用。CUDA 没有类似于目标文件的中间结构可合成 CUDA 模块，而是由 `nvcc` 直接生成可作为 CUDA 模块加载的文件： | NVCC 参数 | 描述 | | ---- | ---- | | -cubin | 编译为特定的 GPU 架构 | | -ptx | 作为驱动即时编译的中间表示 | | -fatbin | 结合了 cubin 和 PTX。若可用则加载合适的 cubin，否则将 PTX 编译到 GPU 上，仅适用于 CUDART | 加载 CUDA 模块后，应用程序可以查询其中包含的资源，如全局变量、函数（内核）和纹理引用。需要注意的是，这些资源在模块加载时就已创建，因此查询函数不会因资源不足而失败。 与上下文类似，CUDA 运行时隐藏了模块的存在和管理。所有模块在 CUDART 初始化时同时加载。对于包含大量 GPU 代码的应用程序，使用驱动 API 而非 CUDA 运行时的主要原因之一是能够通过加载和卸载模块来显式管理驻留性。 cubin 已编译为特定的 GPU 架构，无需“即时”编译，加载速度更快，但不具备向后和向前兼容性。因此，只有事先了解目标 GPU 架构的应用程序才能在不嵌入相同模块的 PTX 版本作为备份的情况下使用 cubin。PTX 是驱动即时编译的中间语言，编译可能需要大量时间，驱动会保存编译后的模块，只要硬件和驱动未更改，就会重复使用。而 fatbin 使 CUDA 运行时能够自动处理合适 cubin 的使用，若不可用则编译 PTX。 ```mermaid graph LR A[nvcc 编译] --> B[-cubin 文件] A --> C[-ptx 文件] A --> D[-fatbin 文件] B --> E[特定 GPU 架构加载] C --> F[驱动即时编译] D --> G{是否有合适 cubin?} G -->|是| E G -->|否| F ``` #### 2. 内核（函数） 在内核编程中，`.cu` 文件里的内核由 `__global__` 关键字标识。使用 CUDA 运行时，可通过 `<<< >>>` 语法内联调用内核。模块的 GPU 可执行代码以内核形式存在，可通过 CUDA 运行时的语言集成特性（`<<< >>>` 语法）或驱动 API 中的 `cuLaunchKernel()` 函数调用。目前，CUDA 不会对 CUDA 模块中的可执行代码进行动态驻留管理，模块加载时，所有内核都会加载到设备内存中。 加载模块后，可使用 `cuModuleGetFunction()` 查询内核，使用 `cuFuncGetAttribute()` 查询内核的属性，使用 `cuLaunchKernel()` 启动内核。`cuLaunchKernel()` 使许多旧的 API 入口点过时，因为旧 API 设置内核启动状态的效率较低，而块大小等参数最好与启动内核的请求一起原子性地指定。 `cuFuncGetAttribute()` 函数可用于查询函数的特定属性，例如： - 每个块的最大线程数 - 静态分配的共享内存量 - 用户分配的常量内存大小 - 每个函数使用的本地内存量 - 函数每个线程使用的寄存器数量 - 函数编译所针对的虚拟（PTX）和二进制架构版本 使用驱动 API 时，建议使用 `extern "C"` 抑制 C++ 的默认名称修饰行为，否则需向 `cuModuleGetFunction()` 指定修饰后的名称。 CUDA 运行时在加载使用其构建的可执行文件时，会在主机内存中创建全局数据结构，描述创建 CUDA 设备时要分配的 CUDA 资源。初始化 CUDA 设备后，这些全局变量用于一次性创建 CUDA 资源。由于这些全局变量在整个进程中由 CUDA 运行时共享，因此使用 CUDA 运行时无法逐步加载和卸载 CUDA 模块。并且由于 CUDA 运行时与 C++ 语言的集成方式，应通过名称（而非字符串字面量）向 API 函数（如 `cudaFuncGetAttributes()` 和 `cudaMemcpyToSymbol()`）指定内核和符号。 #### 3. 设备内存 设备内存（或线性设备内存）位于 CUDA 地址空间，CUDA 内核可通过普通的 C/C++ 指针和数组解引用操作访问。大多数 GPU 都有一个直接连接到 GPU 并由集成内存控制器访问的专用设备内存池。CUDA 硬件不支持按需分页，所有内存分配都由实际物理内存支持。与 CPU 应用程序不同，当物理内存耗尽时，CUDA 的内存分配功能会失败。 ##### 3.1 CUDA 运行时 CUDA 运行时应用程序可通过调用 `cudaGetDeviceProperties()` 并检查 `cudaDeviceProp::totalGlobalMem` 来查询给定设备上可用的设备内存总量。`cudaMalloc()` 和 `cudaFree()` 分别用于分配和释放设备内存，`cudaMallocPitch()` 用于分配带间距的内存，`cudaMalloc3D()