CUDA内存管理：全局内存、共享内存、常量内存、纹理内存的使用与优化

核心原理：CUDA内存层次结构与带宽差异

CUDA设备内存并非单一存储池，而是由多个具有不同容量、延迟和带宽的层级构成。理解这一层次结构是编写高性能内核的前提。从程序员视角看，最常打交道的四类内存为：全局内存（Global Memory）、共享内存（Shared Memory）、常量内存（Constant Memory）和纹理内存（Texture Memory）。它们各自服务于不同的访问模式，性能差异可达两个数量级。

• 全局内存：容量最大（通常数GB至数十GB），但延迟高达400~600周期。所有线程均可读写，是CPU与GPU交换数据的主通道。其性能高度依赖访问的合并度（Coalescing）。
• 共享内存：位于SM内部，容量仅几十KB（可配置），延迟约20~30周期，带宽极高。它用于线程块内数据复用与协作，是优化最核心的武器。
• 常量内存：只读，容量64KB，具有专用的常量缓存。当所有线程同时访问同一地址时，广播机制使其速度接近寄存器。
• 纹理内存：只读，通过纹理缓存（Texture Cache）服务，针对二维/三维空间局部性访问做了硬件优化，适合图像、插值、边界处理等场景。

方法流程：内存选择与优化决策路径

面对一个具体的内核，选择哪种内存并非凭直觉，而应遵循系统化的决策流程。以下步骤可帮助快速定位内存优化方向：

1. 分析数据访问模式：统计每个数据被线程复用的次数。若复用次数≥2，优先考虑共享内存；若所有线程访问同一地址且数据只读，考虑常量内存。
2. 检查全局内存合并度：使用nvprof --metrics gld_efficiency或Nsight Compute查看当前内核的全局内存加载效率。若低于80%，需调整索引映射或使用__ldg()只读缓存。
3. 评估共享内存容量与Bank冲突：计算所需共享内存大小（每块线程），确认不超过SM上限（如48KB/块）。使用填充（Padding）技术避免Bank冲突。
4. 判断是否适合纹理/常量：若数据具有二维空间局部性（如图像卷积），或需要硬件双线性插值，纹理内存是首选。若数据为固定系数且所有线程同址访问，常量内存更优。
5. 迭代测量：每次修改后重新测量有效带宽（Effective Bandwidth = (bytes_read + bytes_written) / time），而非仅看内核耗时。

案例/场景：矩阵乘法与图像卷积的实战对比

以经典的矩阵乘法（C = A × B）为例，展示共享内存的优化效果。朴素实现中，每个线程直接读取全局内存中的A和B元素，未合并访问导致带宽浪费。使用分块（Tiling）策略后，每个线程块协作将子矩阵加载到共享内存：

// 共享内存分块矩阵乘法核心（块大小16x16）
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];
    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;
    float sum = 0.0f;

    for (int k = 0; k < N; k += 16) {
        As[threadIdx.y][threadIdx.x] = A[row * N + k + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * N + col];
        __syncthreads();  // 确保加载完成
        for (int i = 0; i < 16; i++)
            sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        __syncthreads();  // 防止覆盖
    }
    C[row * N + col] = sum;
}

该实现将全局内存访问次数从2×N³降至2×N³/16，在N=1024时性能提升约8~12倍。而图像卷积场景则更适合纹理内存：

• 纹理内存的硬件缓存天然利用像素的空间局部性，相邻线程访问相邻像素时命中率极高。
• 使用tex2D()获取像素值时，硬件自动处理边界条件（如Clamp或Wrap），省去手动判断。
• 对于浮点纹理，硬件支持双线性插值，可免费实现图像缩放。

误区/复盘：常见性能陷阱与修正策略

在实际开发中，即使理解了内存层次，仍容易陷入以下误区。复盘这些典型错误，有助于建立正确的优化直觉。

误区3：常量内存使用不当。常量内存的广播优势仅在所有线程访问同一地址时成立。若线程访问不同地址（如coef[threadIdx.x]），常量缓存会退化为串行访问，性能远低于全局内存。修正：将此类数据放入全局内存并使用__ldg()。

误区4：忽视纹理内存的容量限制。纹理内存虽方便，但最大绑定尺寸受限于cudaDeviceProp::maxTexture1D等参数。对于超大数组，需分块绑定或改用全局内存只读路径。此外，纹理内存的写入需通过cudaMemcpyToArray或cudaMemcpy2DToArray，拷贝开销需纳入考量。