第四章：CUDA内存模型深度解析

各位同学，今天我们来聊聊CUDA的内存模型。说实话，这是很多人在GPU编程中栽跟头的地方。我见过太多项目，算法设计得挺漂亮，结果一跑起来性能惨不忍睹——十有八九是内存访问模式出了问题。

CUDA的内存模型，说白了就是一张「数据存放地图」。你想想看，GPU里有几千个核心同时干活，数据放在哪儿、怎么搬、谁先谁后，这些搞不清楚，程序就跑不快。我个人习惯把CUDA内存分成六个层次：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存。咱们一个一个来看。

4.1 全局内存（Global Memory）

全局内存是GPU里最大的存储空间，所有线程都能访问。但大归大，速度嘛...嗯，比较慢。它的延迟通常在400-800个时钟周期左右。

关键特点：

  • 容量最大（通常几个GB到几十个GB）
  • 所有线程（包括CPU）都能读写
  • 延迟高，需要合并访问才能发挥性能
  • 生命周期由程序员控制（cudaMalloc/cudaFree）

我在项目中遇到过最典型的坑：新手把全局内存当普通数组用，随机访问。结果性能直接掉到谷底。为什么？因为全局内存的带宽利用率取决于「合并访问」——相邻线程访问相邻地址，才能一次搬一大块数据。

// 好的访问模式：合并访问
__global__ void goodAccess(float *data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[idx] = data[idx] * 2.0f;  // 相邻线程访问相邻地址
    }
}

// 差的访问模式：非合并访问
__global__ void badAccess(float *data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[threadIdx.x * N + blockIdx.x] = 1.0f;  // 跨步访问，性能极差
    }
}

避坑指南：我曾经调试过一个图像处理程序，全局内存访问模式没优化，带宽利用率只有15%。改成合并访问后，直接飙到85%。记住：让线程0访问地址0，线程1访问地址1，以此类推。

4.2 共享内存（Shared Memory）

共享内存是块内所有线程共享的「小仓库」。它就在芯片上，速度比全局内存快两个数量级。延迟只有几个时钟周期。

属性 值
容量 每个SM通常48KB-164KB（取决于架构）
延迟 约5-30个时钟周期
作用域 单个线程块内
生命周期 线程块生命周期内

共享内存最常用的场景就是「数据复用」。比如矩阵乘法，每个线程需要多次读取同一块数据。把它先搬到共享内存里，能省下大量全局内存访问。

__global__ void sharedMemExample(float *A, float *B, float *C, int N) {
    __shared__ float sA[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float sB[BLOCK_SIZE][BLOCK_SIZE];
    
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    float sum = 0.0f;
    for (int t = 0; t < N / BLOCK_SIZE; t++) {
        // 加载到共享内存
        sA[threadIdx.y][threadIdx.x] = A[row * N + t * BLOCK_SIZE + threadIdx.x];
        sB[threadIdx.y][threadIdx.x] = B[(t * BLOCK_SIZE + threadIdx.y) * N + col];
        __syncthreads();  // 同步！确保所有线程加载完成
        
        // 从共享内存计算
        for (int k = 0; k < BLOCK_SIZE; k++) {
            sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

注意：共享内存用不好反而会拖慢速度。我见过有人把共享内存当全局内存用，每个线程只读一次数据——那还不如直接用全局内存呢。共享内存的核心理念是「复用」和「规约」。

4.3 寄存器（Registers）

寄存器是GPU里最快、最小的存储。每个线程独享，延迟为0（一个时钟周期内完成）。但数量有限——每个线程最多255个寄存器（不同架构略有差异）。

你想想看，如果一个线程用了太多寄存器，编译器会把多余的数据「溢出」到本地内存。本地内存虽然名字叫「本地」，实际上走的还是全局内存的通道，速度一下子就掉下来了。

经验之谈：我调试性能时，第一件事就是看寄存器使用量。用nvcc的--ptxas-options=-v选项可以查看。如果每个线程用了超过64个寄存器，就要考虑优化了——要么减少局部变量，要么把一些数据放到共享内存里。

4.4 本地内存（Local Memory）

本地内存这个名字容易让人误解。它其实不是「本地」的，而是存在全局内存里，只是每个线程私有的。编译器会自动把寄存器放不下的数据、数组、或者一些大的结构体放到本地内存。

说白了，本地内存就是寄存器的「备胎」。能用寄存器尽量用寄存器，实在不行才用本地内存。

// 这个数组可能会被放到本地内存
__global__ void localMemExample() {
    int largeArray[100];  // 寄存器放不下，编译器会放到本地内存
    // ... 使用largeArray
}

建议：尽量避免在核函数里声明大数组。如果必须用，考虑用共享内存或者全局内存手动管理。我曾经优化过一个程序，把本地内存里的数组改成共享内存，性能提升了3倍。

4.5 常量内存（Constant Memory）

常量内存是全局内存的一个「只读缓存版本」。容量小（通常64KB），但访问速度快——如果所有线程访问同一个地址，延迟几乎和寄存器一样。

常量内存最适合存放那些不会改变、且所有线程都需要的「查表数据」。比如数学函数的系数表、物理常数等。

__constant__ float coefficients[256];

__global__ void constantMemKernel(float *input, float *output, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        // 所有线程访问同一个系数
        output[idx] = input[idx] * coefficients[blockIdx.x % 256];
    }
}

注意：常量内存的「广播」特性只在所有线程访问同一地址时有效。如果每个线程访问不同地址，性能反而比全局内存还差。我在项目中吃过这个亏，后来改成纹理内存才解决问题。

4.6 纹理内存（Texture Memory）

纹理内存最初是为图形学设计的，但后来发现它在某些计算场景下特别好用。它有硬件支持的「缓存」和「插值」功能，特别适合二维/三维数据的随机访问。

纹理内存有几个独特优势：

  • 硬件缓存：对空间局部性好的访问模式很友好
  • 边界处理：可以自动处理越界访问（返回0或边界值）
  • 插值：硬件支持线性插值，适合图像处理
  • 只读：不用担心数据竞争
// 声明纹理引用
texture<float, 2, cudaReadModeElementType> texRef;

__global__ void textureKernel(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) {
        // 使用纹理内存读取，自动处理缓存
        output[y * width + x] = tex2D(texRef, x + 0.5f, y + 0.5f);
    }
}

个人经验：纹理内存在处理「非合并访问」时特别有用。比如光线追踪、体渲染这些需要随机访问数据的场景。我做过一个医学图像重建的项目，用纹理内存比全局内存快了将近2倍。

4.7 内存模型总结

好了，咱们把六种内存捋一遍。我画个表格帮你对比：

内存类型 位置 速度 容量 作用域 生命周期
寄存器 芯片上 最快 极小（每个线程） 单个线程 核函数内
共享内存 芯片上 快 小（每个SM） 线程块内 线程块内
常量内存 芯片外+缓存 较快（广播时） 64KB 全局 应用程序内
纹理内存 芯片外+缓存 较快（有缓存） 大 全局 应用程序内
本地内存 芯片外 慢 大（每个线程） 单个线程 核函数内
全局内存 芯片外 最慢 最大 全局 应用程序内

最后说一句：选哪种内存，取决于你的数据访问模式。没有银弹。我见过有人不管三七二十一全用共享内存，结果因为容量限制导致occupancy下降，性能反而更差。记住一个原则：数据在哪里用得多，就放在离计算单元近的地方。

下一章咱们聊同步机制，那是另一个容易踩坑的地方。到时候我会讲讲__syncthreads()的坑，还有原子操作的性能陷阱。咱们下节课见。

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