CUDA内存管理全解析

# CUDA 内存管理全解析 在 CUDA 编程中，内存管理是至关重要的一环，它直接影响着程序的性能和稳定性。本文将深入探讨 CUDA 中各种类型的内存，包括 ECC 内存、常量内存、局部内存、纹理内存、共享内存以及内存复制操作，帮助大家更好地理解和运用 CUDA 内存管理机制。 ## 1. ECC 内存 ECC（Error-Correcting Code）内存具有以下特点： - **内存减少**：会使可用内存减少 12.5%。例如，在 Amazon EC2 的 cg1.4xlarge 实例中，内存会从 3071MB 减少到 2687MB。 - **上下文同步成本增加**：使上下文同步的开销变大。 - **非合并内存事务成本增加**：启用 ECC 时，非合并内存事务的成本会更高。 可以使用 `nvidia-smi` 命令行工具（具体描述可参考相关文档）或 NVML（NVIDIA Management Library）来启用和禁用 ECC。当检测到不可纠正的 ECC 错误时，同步错误报告机制会返回 `cudaErrorECCUncorrectable`（针对 CUDA 运行时）和 `CUDA_ERROR_ECC_UNCORRECTABLE`（针对驱动 API）。 ## 2. 常量内存 常量内存针对向多个线程进行只读广播进行了优化。编译器使用常量内存来存储那些难以计算或无法直接编译到机器代码中的常量。常量内存位于设备内存中，但通过特殊的“常量缓存”进行访问。编译器有 64K 的常量内存可供使用，开发者还可以使用 `__constant__` 关键字声明另外 64K 的内存。这些限制是按模块（对于驱动 API 应用程序）或按文件（对于 CUDA 运行时应用程序）来计算的。 ### 2.1 主机和设备的 `__constant__` 内存 Mark Harris 介绍了一种使用预定义宏 `__CUDA_ARCH__` 来维护 `__constant__` 内存的主机和设备副本的方法，方便 CPU 和 GPU 访问。示例代码如下： ```c __constant__ double dc_vals[2] = { 0.0, 1000.0 }; const double hc_vals[2] = { 0.0, 1000.0 }; __device__ __host__ double f(size_t i) { #ifdef __CUDA_ARCH__ return dc_vals[i]; #else return hc_vals[i]; #endif } ``` ### 2.2 访问 `__constant__` 内存 除了 C/C++ 运算符隐式访问常量内存外，开发者还可以进行常量内存的复制操作，甚至查询常量内存分配的指针。 #### 2.2.1 CUDA 运行时 CUDA 运行时应用程序可以使用 `cudaMemcpyToSymbol()` 和 `cudaMemcpyFromSymbol()` 分别进行常量内存的复制操作。可以使用 `cudaGetSymbolAddress()` 查询常量内存的指针。 ```c cudaError_t cudaGetSymbolAddress( void **devPtr, char *symbol ); ``` 需要注意的是，开发者在使用该指针通过内核写入常量内存时，要避免在另一个内核正在读取该常量内存时进行写入操作。 #### 2.2.2 驱动 API 驱动 API 应用程序可以使用 `cuModuleGetGlobal()` 查询常量内存的设备指针。由于驱动 API 没有像 CUDA 运行时那样的语言集成，因此没有专门的内存复制函数。应用程序需要先使用 `cuModuleGetGlobal()` 查询地址，然后调用 `cuMemcpyHtoD()` 或 `cuMemcpyDtoH()` 进行内存复制。 可以使用 `cuFuncGetAttribute(CU_FUNC_ATTRIBUTE_CONSTANT_SIZE_BYTES)` 查询内核使用的常量内存量。 ## 3. 局部内存 局部内存包含 CUDA 内核中每个线程的栈，其用途如下： - **实现应用程序二进制接口（ABI）**：即调用约定。 - **寄存器溢出数据**：当寄存器空间不足时，将数据溢出到局部内存。 - **存储编译器无法解析索引的数组**：对于那些索引无法在编译时确定的数组，会存储在局部内存中。 在早期的 CUDA 硬件实现中，使用局部内存会严重影响性能，开发者通常会采取各种措施避免使用局部内存。但随着 Fermi 架构中 L1 缓存的出现，只要局部内存流量限制在 L1 内，这些性能问题就不再那么紧迫。 开发者可以使用 `nvcc` 选项 `-Xptxas –v,abi=no` 让编译器报告给定内核所需的局部内存量。在运行时，可以使用 `cuFuncGetAttribute(CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES)` 查询内核使用的局部内存量。 ### 3.1 性能优化策略 寄存器溢出会带来两个成本：指令数量增加和内存流量增加。可以使用 L1 和 L2 性能计数器来确定内存流量是否影响性能。以下是一些提高性能的策略： - **增加寄存器数量**：在编译时，通过 `–maxregcount` 指定更高的寄存器数量限制。增加线程可用的寄存器数量可以减少指令数量和内存流量。在使用 PTXAS 在线编译内核时，可以使用 `__launch_bounds__` 指令调整该参数。 - **使用非缓存加载**：对于全局内存，使用非缓存加载，例如 `nvcc –Xptxas -dlcm=cg`。 - **增加 L1 缓存大小**：将 L1 缓存大小增加到 48K，可以调用 `cudaFuncSetCacheConfig()` 或 `cudaDeviceSetCacheconfig()`。 ### 3.2 示例代码 下面是一个 `GlobalCopy` 内核的示例代码，它展示了寄存器溢出时的“性能悬崖”： ```c template<class T, const int n> __global__ void GlobalCopy( T *out, const T *in, size_t N ) { T temp[n]; size_t i; for ( i = n*blockIdx.x*blockDim.x+threadIdx.x; i < N-n*blockDim.x*gridDim.x; i += n*blockDim.x*gridDim.x ) { for ( int j = 0; j < n; j++ ) { size_t index = i+j*blockDim.x; temp[j] = in[index]; } for ( int j = 0; j < n; j++ ) { size_t index = i+j*blockDim.x; out[index] = temp[j]; } } // to avoid the (index<N) conditional in the inner loop, // we left off some work at the end for ( int j = 0; j < n; j++ ) { for ( int j = 0; j < n; j++ ) { size_t index = i+j*blockDim.x; if ( index<N ) temp[j] = in[index]; } for ( int j = 0; j < n; j++ ) { size_t index = i+j*blockDim.x; if ( index<N ) out[index] = temp[j]; } } } ``` ### 3.3 性能分析 以下是 `globalCopy.cu` 在 GK104 GPU 上的部分输出，仅显示 64 位操作数的复制性能： | Unroll | 32 | 64 | 128 | 256 | 512 | maxBW | maxThreads | | ---- | ---- | ---- | ---- | ---- | ---- | ---- | ---- | | 1 | 75.57 | 102.57 | 116.03 | 124.51 | 126.21 | 126.21 | 512 | | 2 | 105.73 | 117.09 | 121.84 | 123.07 | 124.00 | 124.00 | 512 | | 3 | 112.49 | 120.88 | 121.56 | 123.09 | 123.44 | 123.44 | 512 | | 4 | 115.54 | 122.89 | 122.38 | 122.15 | 121.22 | 122.89 | 64 | | 5 | 113.81 | 121.29 | 120.11 | 119.69 | 116.02 | 121.29 | 64 | | 6 | 114.84 | 119.49 | 120.56 | 118.09 | 117.88 | 120.56 | 128 | | 7 | 117.53 | 122.94 | 118.74 | 116.52 | 110.99 | 122.94 | 64 | | 8 | 116.89 | 121.68 | 119.00 | 113.49 | 105.69 | 121.68 | 64 | | 9 | 116.10 | 120.73 | 115.96 | 109.48 | 99.60 | 120.73 | 64 | | 10 | 115.02 | 116.70 | 115.30 | 106.31 | 93.56 | 116.70 | 64 | | 11 | 113.67 | 117.36 | 111.48 | 102.84 | 88.31 | 117.36 | 64 | | 12 | 88.16 | 86.91 | 83.68 | 73.78 | 58.55 | 88.16 | 32 | | 13 | 85.27 | 85.58 | 80.09 | 68.51 | 52.66 | 85.58 | 64 | | 14 | 78.60 | 76.30 | 69.50 | 56.59 | 41.29 | 78.60 | 32 | | 15 | 69.00 | 65.78 | 59.82 | 48.41 | 34.65 | 69.00 | 32 | | 16 | 65.68 | 62.16 | 54.71 | 43.02 | 29.92 | 65.68 | 32 | 从输出结果可以看出，当循环展开为 12 时，由于寄存器溢出，性能开始明显下降，带宽从 117GB/s 降至不到 90GB/s，随着循环展开增加到 16，性能进一步下降到 30GB/s 以下。 下面是 `globalCopy` 内核对应展开循环的寄存器和局部内存使用情况总结： | UNROLL FACTOR | REGISTERS | LOCAL MEMORY (BYTES) | | ---- | ---- | ---- | | 1 | 20 | None | | 2 | 19 | None | | 3 | 26 | None | | 4 | 33 | None | | 5 |