CUDA内存优化：全局内存、共享内存、常量内存、纹理内存与访问模式优化

CUDA程序性能的瓶颈往往不在计算单元，而在内存系统。GPU拥有庞大的并行线程规模，如果内存访问模式低效，数千个线程会同时陷入等待，导致计算单元空转。理解各类内存的物理特性、作用域与适用场景，是编写高性能CUDA内核的必修课。本章从内存层次结构出发，逐一剖析全局内存、共享内存、常量内存和纹理内存，并重点讲解内存访问模式对带宽利用率的影响。

一、CUDA内存层次结构概览

CUDA设备内存模型可视为一个金字塔结构，从寄存器到全局内存，容量递增、延迟递增、带宽递减。每个线程拥有私有寄存器和局部内存；每个线程块共享一块共享内存；所有线程共享全局内存、常量内存和纹理内存。下表总结了各类内存的关键属性：

二、全局内存与合并访问

全局内存是GPU中容量最大的存储空间，但也是延迟最高的。其性能关键不在于单次访问的延迟，而在于内存事务的利用率。GPU以32个线程为一组（warp）执行指令，当warp中的线程访问全局内存时，硬件会将这些访问合并为尽可能少的内存事务（通常为32字节、64字节或128字节的缓存行）。

合并访问（Coalesced Access）是指warp内所有线程访问的地址落在同一个或少数几个连续的内存段中。理想情况下，32个线程访问连续的32个浮点数（128字节），硬件只需一次事务即可完成。反之，如果线程访问地址分散（如按列访问二维数组），则可能产生32次独立事务，带宽利用率下降数倍甚至数十倍。

// 合并访问示例：连续索引访问
__global__ void coalesced_copy(const float* in, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = in[idx];  // 相邻线程访问相邻地址 → 合并
    }
}

// 非合并访问示例：按列访问（假设行优先存储）
__global__ void strided_copy(const float* in, float* out, int rows, int cols) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int col = blockIdx.y * blockDim.y + threadIdx.y;
    if (row < rows && col < cols) {
        out[row * cols + col] = in[col * rows + row];  // 转置访问 → 跨步
    }
}

三、共享内存：块内协作的利器

共享内存是片上存储，延迟远低于全局内存，且带宽极高。它允许同一线程块内的线程互相交换数据，是实现数据复用和块内归约的关键。共享内存需要显式声明和同步管理，使用不当会导致数据竞争或bank冲突。

共享内存被划分为32个bank，每个bank宽度为4字节。当多个线程同时访问不同bank时，可以并行执行；若多个线程访问同一bank的不同地址，则发生bank冲突，访问被串行化。理想情况是warp内32个线程访问32个不同bank。

// 共享内存矩阵转置示例（使用padding避免bank冲突）
#define TILE_SIZE 32
__global__ void transpose_shared(const float* in, float* out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE + 1];  // +1 padding
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    if (x < width && y < width) {
        tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    }
    __syncthreads();
    
    int x_out = blockIdx.y * TILE_SIZE + threadIdx.x;
    int y_out = blockIdx.x * TILE_SIZE + threadIdx.y;
    
    if (x_out < width && y_out < width) {
        out[y_out * width + x_out] = tile[threadIdx.x][threadIdx.y];
    }
}

四、常量内存与纹理内存

常量内存用于存储只读数据，具有专用的常量缓存。当warp内所有线程访问同一地址时，常量缓存效率极高（广播机制，一次事务服务所有线程）。但如果线程访问不同地址，则会被串行化，性能急剧下降。因此常量内存适合存储系数、参数表等所有线程都会用到的数据。

纹理内存最初为图形渲染设计，具有以下特性：支持硬件插值（如双线性过滤）、处理边界条件（clamp/wrap）、对二维空间局部性有优化。纹理缓存针对空间局部性设计，适合图像处理、随机访问模式不规律但具有空间相关性的数据。在计算场景中，纹理内存常用于图像卷积、光线追踪等。

// 常量内存声明与使用
__constant__ float coeff[256];

__global__ void apply_coeff(const float* in, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = in[idx] * coeff[idx % 256];  // 所有线程访问相同coeff地址时效率最高
    }
}

// 纹理内存声明（需要纹理对象）
texture<float, 2, cudaReadModeElementType> texRef;

__global__ void texture_sample(float* out, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        out[y * width + x] = tex2D(texRef, x, y);  // 硬件处理边界与插值
    }
}

五、内存访问模式优化实战

以矩阵乘法为例，展示如何通过共享内存和合并访问将性能提升数倍。朴素实现中，每个线程直接读取全局内存中的A和B矩阵元素，存在大量重复读取。使用分块（tiling）策略后，每个线程块将子矩阵加载到共享内存，块内线程复用这些数据，大幅减少全局内存访问次数。

// 分块矩阵乘法（每个块计算TILE_SIZE×TILE_SIZE子矩阵）
#define TILE_SIZE 16
__global__ void matmul_tiled(const float* A, const float* B, float* C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int t = 0; t < (N + TILE_SIZE - 1) / TILE_SIZE; ++t) {
        // 协作加载A和B的子块到共享内存
        if (row < N && t * TILE_SIZE + threadIdx.x < N)
            As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;
            
        if (col < N && t * TILE_SIZE + threadIdx.y < N)
            Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
            
        __syncthreads();
        
        for (int k = 0; k < TILE_SIZE; ++k) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 防止下一轮覆盖数据
    }
    
    if (row < N && col < N) {
        C[row * N + col] = sum;
    }
}

六、误区与复盘

• 误区1：所有数据都放共享内存。共享内存容量有限（通常几十KB/SM），且过度使用降低占用率。应只缓存被多次复用的数据。
• 误区2：忽略L2缓存的作用。现代GPU的L2缓存（可达40MB）能自动缓存全局内存访问，对于有一定局部性的访问模式，即使未显式使用共享内存，也能获得较好性能。
• 误区3：常量内存随意使用。如果warp内线程访问不同常量地址，性能反而比全局内存更差。常量内存只适合广播式访问。
• 误区4：纹理内存只用于图形。纹理缓存的二维局部性优化对许多科学计算（如有限差分、图像处理）同样有效，值得尝试。
• 误区5：忽略内存对齐。结构体数组（AoS）往往导致访问跨步，应优先使用数组结构体（SoA）布局，确保每个字段连续存储。

内存优化是一个反复测量与调整的过程。建议使用NVIDIA Nsight Compute分析内存吞吐量、合并率、bank冲突等指标，定位瓶颈后再针对性优化。记住：先保证合并访问，再考虑共享内存复用，最后才考虑常量/纹理等特殊内存，这是大多数场景下的最优策略顺序。