1. CUDA内存模型全景：理解Host与Device内存架构

大家好，我是你们的老朋友。今天咱们聊聊CUDA内存模型。

说实话，我见过太多新手一上来就写cudaMalloc，然后数据拷贝完就开始算。结果呢？性能惨不忍睹。为什么？因为他们根本不了解GPU的内存长什么样。

你想想看，CPU和GPU之间隔着PCIe总线，这玩意儿带宽再高也就几十GB/s。而GPU自己的显存带宽，动不动就几百GB/s甚至上TB/s。这差距，说白了就是高速公路和乡间小道的区别。

所以，理解内存模型，是CUDA优化的第一步，也是最重要的一步。

1.1 Host与Device：两个世界

Host就是CPU端，Device就是GPU端。它们各自有独立的内存空间。

• Host内存：就是我们平时用的DDR内存，容量大（几十GB到几百GB），但带宽一般。
• Device内存：就是GPU上的显存，容量相对小（几GB到几十GB），但带宽极高。

数据要从Host搬到Device，或者反过来，都得经过PCIe总线。这个操作叫Host-Device数据传输。

我记得有个项目，客户把数据分成了1000个小块，每块都单独cudaMemcpy一次。结果光拷贝就占了90%的时间。我改成一次大块传输，性能直接翻了5倍。嗯，这就是血的教训。

1.2 全局内存：GPU的“主存”

全局内存（Global Memory）是GPU上最大的存储空间。所有线程都能访问它。但代价是——慢。

全局内存的延迟大概在400-800个时钟周期。什么概念？你算一次浮点加法只要几个周期，等一个数据却要几百个周期。这就像你去食堂打饭，排队的时间比吃饭的时间还长。

怎么优化？合并访问（Coalesced Access）是关键。

// 好的合并访问：相邻线程访问相邻地址
__global__ void good_kernel(float *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = data[idx] * 2.0f;  // 线程0访问data[0]，线程1访问data[1]...
}

// 差的非合并访问：跨步访问
__global__ void bad_kernel(float *data) {
    int idx = threadIdx.x * 32 + blockIdx.x * blockDim.x;
    data[idx] = data[idx] * 2.0f;  // 线程0访问data[0]，线程1访问data[32]...
}

1.3 共享内存：片上“缓存”

共享内存（Shared Memory）是GPU芯片上的SRAM，速度比全局内存快两个数量级。延迟只有几个时钟周期。

但容量很小——每个SM（流多处理器）通常只有几十KB到一百多KB。你得精打细算着用。

共享内存的作用是什么？线程间协作。同一个block里的线程可以通过共享内存交换数据，不用反复访问慢速的全局内存。

// 共享内存的典型用法：矩阵分块
__global__ void matmul_kernel(float *A, float *B, float *C, int N) {
    __shared__ float tileA[16][16];
    __shared__ float tileB[16][16];
    
    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;
    
    float sum = 0.0f;
    for (int t = 0; t < N / 16; t++) {
        // 从全局内存加载到共享内存
        tileA[threadIdx.y][threadIdx.x] = A[row * N + t * 16 + threadIdx.x];
        tileB[threadIdx.y][threadIdx.x] = B[(t * 16 + threadIdx.y) * N + col];
        __syncthreads();
        
        // 从共享内存计算
        for (int k = 0; k < 16; k++) {
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

1.4 寄存器：最快的存储

寄存器（Registers）是GPU上最快的存储，没有之一。延迟为0个周期。每个线程都有自己的私有寄存器。

但数量有限——每个线程通常只有几十到几百个寄存器。如果你声明太多局部变量，编译器会“溢出”（spill）到本地内存，速度瞬间掉下来。

1.5 本地内存：寄存器的“备胎”

本地内存（Local Memory）听起来像“本地”的，其实它存在全局内存里。只是每个线程私有的而已。

什么时候会用本地内存？寄存器不够用的时候。编译器会自动把多余的变量放到本地内存里。

本地内存的延迟和全局内存一样高。所以，尽量避免寄存器溢出。

1.6 常量内存：只读的“广播”

常量内存（Constant Memory）是只读的，容量64KB。它有个特殊能力——广播。当同一个warp里的所有线程访问同一个常量地址时，硬件只需要一次读取，然后广播给所有线程。

适合放什么？查找表、系数、配置参数等只读数据。

__constant__ float coeff[256];

__global__ void kernel(float *data) {
    int idx = threadIdx.x;
    // 所有线程访问同一个coeff[0]，硬件会广播
    data[idx] = data[idx] * coeff[0];
}

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

纹理内存（Texture Memory）也是只读的，但它有硬件缓存，特别适合二维或三维的随机访问模式。

比如图像处理、体渲染这些场景。纹理内存会自动处理边界条件，还支持插值。

我个人很少用纹理内存，因为大多数场景用共享内存+全局内存就能搞定。但如果你做图像处理，纹理内存是个好选择。

1.8 层级关系与带宽差异

好了，咱们把所有的内存类型串起来，看看它们的层级关系：

这张图很直观。从上到下，速度越来越快，容量越来越小。你想想看，优化说白了就是：把频繁访问的数据放到更快的存储里。

内存类型	位置	容量	延迟	作用域
寄存器	SM内部	每个线程最多255个	0周期	单个线程
共享内存	SM内部	48-164 KB/SM	~5周期	单个Block
常量内存	显存	64 KB	~1周期(广播)	所有线程(只读)
纹理内存	显存	取决于显存	取决于缓存	所有线程(只读)
全局内存	显存	几GB~几十GB	400-800周期	所有线程
本地内存	显存	取决于寄存器溢出	400-800周期	单个线程
Host内存	CPU端	几十GB~几百GB	~100ns	CPU线程

嗯，这就是CUDA内存模型的全景。说白了，就是一张地图。你知道了哪里快哪里慢，才能规划出最优的数据路径。

我曾经有个项目，优化前kernel跑200ms，优化后只用了15ms。怎么做到的？就是把热点数据从全局内存搬到了共享内存，然后做了合并访问。就这么简单。

但简单不代表容易。你得理解每一层内存的特性，才能用好它们。