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

说到CUDA编程，内存管理绝对是绕不开的核心话题。我刚开始接触CUDA时，总觉得把数据扔到GPU上就能自动跑得快，结果被现实狠狠教育了一顿。说白了，不理解内存模型，你写的CUDA代码可能比CPU还慢。

今天咱们就聊聊CUDA里的四种关键内存：全局内存、共享内存、常量内存和纹理内存。每种内存都有它的脾气，用对了是神器，用错了是累赘。

3.1 全局内存：最常用，但最慢

全局内存是GPU上容量最大的存储空间，所有线程都能访问。我习惯叫它"大仓库"——什么都能放，但取东西特别慢。

它的延迟有多夸张？大概400-800个时钟周期。你想想看，一个线程等一次全局内存访问，够ALU做几百次加法了。

关键优化点：合并访问

全局内存访问必须满足"合并访问"条件才能达到最佳带宽。说白了，相邻的线程要访问相邻的内存地址。

我在项目中遇到过这样一个坑：一个图像处理任务，我按行处理数据，结果带宽利用率只有20%。后来改成按列处理，带宽直接拉满。为什么？因为线程束里的32个线程访问的内存地址不连续，导致多次内存事务。

// ❌ 糟糕的访问模式：线程间地址不连续
__global__ void bad_access(float *data, int width) {
    int idx = threadIdx.x;
    for (int i = 0; i < 100; i++) {
        data[idx + i * width] *= 2.0f;  // 跨度大，不连续
    }
}

// ✅ 优化的访问模式：线程间地址连续
__global__ void good_access(float *data, int width) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] *= 2.0f;  // 连续访问
}

我的小技巧：分配全局内存时，尽量使用cudaMallocPitch()而不是cudaMalloc()。它能自动对齐内存，保证每行的起始地址满足对齐要求，这对合并访问很有帮助。

3.2 共享内存：片上高速缓存

共享内存是GPU上的"高速缓存"，位于芯片内部，延迟只有几个时钟周期。每个线程块独享一份共享内存，块内线程可以互相通信。

它的容量很小——老一点的卡只有16KB，新的卡最多也就48KB或96KB。所以你得精打细算。

我记得有一次做矩阵乘法优化，用共享内存把数据分块加载，性能提升了将近10倍。当时我盯着性能分析器看了半天，确认不是幻觉。

// 共享内存实现矩阵分块乘法（核心片段）
__global__ void matmul_shared(float *A, 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 tile = 0; tile < N / TILE_SIZE; tile++) {
        // 协作加载数据到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * TILE_SIZE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * TILE_SIZE + threadIdx.y) * N + col];
        __syncthreads();  // 同步，确保所有线程加载完成
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 防止下一轮覆盖数据
    }
    C[row * N + col] = sum;
}

注意：共享内存使用不当会导致bank conflict。当多个线程访问同一个bank的不同地址时，访问会被串行化。我曾经因为这个原因，优化了半天性能反而下降了。解决办法是调整数据布局，或者使用padding技巧。

3.3 常量内存：只读且广播

常量内存是GPU上一块只读的存储区域，容量只有64KB。但它有个绝活：当线程束内所有线程访问同一个常量地址时，硬件会把这个值广播给所有线程，只产生一次内存事务。

这玩意儿最适合放什么？查找表、滤波器系数、变换矩阵——就是那些所有线程都要用到的只读数据。

我做过一个图像滤波的项目，滤波器系数放在全局内存里，性能一般。后来改成常量内存，速度提升了30%。原因就是线程束里的线程都在读同一个系数，广播机制发挥了作用。

// 声明常量内存
__constant__ float filter[256];

// 主机端复制数据到常量内存
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 内核中使用
__global__ void apply_filter(float *input, float *output, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) {
        float val = input[idx];
        // 所有线程访问同一个filter地址，硬件广播
        output[idx] = val * filter[0] + val * filter[1];
    }
}

避坑指南：我曾经把一个大数组声明为常量内存，结果编译通过但运行时出错。后来才想起来常量内存只有64KB。记住，常量内存不是用来存大数据的，它只适合存小尺寸的只读参数。

3.4 纹理内存：为空间局部性而生

纹理内存最初是为图形学设计的，但它对某些计算任务特别有用。它利用空间局部性——如果线程访问了某个地址，附近的地址会被缓存起来。

纹理内存还提供硬件级别的边界处理、插值等功能。做图像处理时，用纹理内存处理边界像素特别方便，不用手动写边界检查代码。

我记得有个做图像缩放的同事，用纹理内存的线性插值功能，几行代码就实现了双线性插值。如果用全局内存自己实现，至少要多写几十行。

// 绑定纹理
texture<float, 2, cudaReadModeElementType> tex;
cudaArray *cuArray;
// ... 创建并填充cuArray ...
cudaBindTextureToArray(tex, cuArray);

// 内核中使用纹理内存
__global__ void texture_kernel(float *output, int width, int height) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    if (x < width && y < height) {
        // 纹理坐标归一化到[0,1)
        float u = (float)x / width;
        float v = (float)y / height;
        output[y * width + x] = tex2D(tex, u, v);
    }
}

3.5 如何选择？一张表说清楚

内存类型 容量 延迟 作用域 最佳使用场景
全局内存 GB级别 高（~400 cycles） 所有线程 大数据、需要随机访问
共享内存 KB级别 低（~5 cycles） 线程块内 数据重用、线程间通信
常量内存 64KB 低（广播时） 所有线程 只读参数、查找表
纹理内存 GB级别 中（有缓存） 所有线程 空间局部性强的访问

3.6 实战建议

说了这么多，到底怎么用？我总结几条经验：

  • 先保证全局内存的合并访问——这是最基础的优化，做不好其他都白搭
  • 数据重用超过2次，考虑用共享内存——加载一次，多次使用，稳赚不赔
  • 小尺寸只读参数，用常量内存——64KB够用就别浪费
  • 图像处理、有空间局部性的任务，试试纹理内存——硬件缓存不是吃素的

嗯，内存管理这块内容确实不少，但掌握了这些，你的CUDA程序性能至少能翻一番。下一章咱们聊聊流和事件，看看怎么让GPU干活更高效。

← 上一章 📖 返回目录 下一章 →