4、GPU内存管理：全局内存、共享内存、寄存器、常量内存、纹理内存

做量化系统这几年，我踩过最大的坑，就是GPU内存管理。

你想想看，一张A100有80GB显存，听起来挺大对吧？但如果你不懂每种内存的特性，写出来的程序可能连10%的性能都发挥不出来。我刚开始做高频因子计算时，就吃过这个亏——程序跑得比CPU还慢，查了半天才发现是全局内存访问模式不对。

今天咱们就把GPU的五种内存掰开揉碎了讲清楚。嗯，这五种分别是：全局内存、共享内存、寄存器、常量内存、纹理内存。每种都有它的脾气，用对了是神器，用错了是累赘。

4.1 全局内存：最常用，但最容易被坑

全局内存是GPU里容量最大的存储空间。说白了，就是显存。所有线程都能读写，但代价是访问延迟很高——大概400到800个时钟周期。

我在做期权定价模型时，一开始直接把所有数据都扔在全局内存里。结果呢？一个网格计算要跑好几分钟。后来才发现，问题出在合并访问上。

// ❌ 错误示范：非合并访问
__global__ void bad_access(float *data) {
    int idx = threadIdx.x;
    float val = data[idx * 100];  // 跨度100，无法合并
}

// ✅ 正确示范：合并访问
__global__ void good_access(float *data) {
    int idx = threadIdx.x;
    float val = data[idx];  // 连续访问，完美合并
}

4.2 共享内存：线程间的"小会议室"

共享内存是每个Block私有的高速缓存。容量很小——通常只有几十KB，但速度极快，延迟大概30个周期左右。

我个人习惯把共享内存当作"手动管理的L1缓存"。什么意思呢？就是你把全局内存里会被反复访问的数据，先搬到共享内存里，然后所有线程从共享内存读。

举个例子。做矩阵转置时，如果直接从全局内存读，会有严重的bank conflict。我建议的做法是：

__global__ void transpose_v2(float *in, float *out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    // 先加载到共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 再从共享内存写回
    out[x * width + y] = tile[threadIdx.x][threadIdx.y];
}

4.3 寄存器：最快的存储，但别贪多

寄存器是每个线程私有的，访问零延迟。但数量有限——每个线程最多能用255个32位寄存器（具体看架构）。

你想想看，如果一个线程用了太多寄存器，会发生什么？编译器会把多余的数据"溢出"到本地内存（其实还是全局内存），速度一下子就掉下来了。

我在做高频因子计算时，就遇到过这个问题。一个核函数里定义了太多临时变量，结果寄存器压力太大，性能反而下降了30%。

4.4 常量内存：适合"广播"场景

常量内存是只读的，有专门的缓存。它的特点是：当所有线程访问同一个地址时，速度极快（跟寄存器差不多）。但如果每个线程访问不同的地址，那就惨了——会串行化。

在量化系统里，我常用常量内存来存一些"查表"数据。比如期权定价里的正态分布累积函数表，或者一些固定的参数数组。

__constant__ float lookup_table[256];

__global__ void use_constant() {
    int idx = threadIdx.x;
    // 所有线程访问同一个地址时，效率最高
    float val = lookup_table[0];  // 极快
    // 每个线程访问不同地址时，会变慢
    float val2 = lookup_table[idx];  // 可能串行化
}

4.5 纹理内存：被低估的"加速器"

纹理内存可能很多人不熟悉。它也是只读的，但有硬件缓存和插值功能。最妙的是，它对空间局部性的访问模式做了优化。

什么意思呢？就是如果你访问地址A，那么地址A附近的数也会被预取到缓存里。这在处理二维数据时特别有用。

我记得有一次做波动率曲面的插值计算，用纹理内存比用全局内存快了将近3倍。因为纹理缓存的硬件会自动处理边界条件和插值，省了我不少代码。

// 声明纹理引用
texture<float, 2, cudaReadModeElementType> tex;

// 绑定纹理
cudaBindTexture2D(NULL, tex, dev_data, &tex_channel_desc, width, height, pitch);

// 在核函数中使用
__global__ void texture_kernel(float *output, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    // 使用归一化坐标访问
    float u = (float)x / width;
    float v = (float)y / height;
    output[y * width + x] = tex2D(tex, u, v);
}

4.6 实战建议：如何选择内存类型？

说了这么多，到底该怎么选？我总结了一个简单的决策流程：

场景	推荐内存类型	原因
大数据量，一次性读取	全局内存	容量大，注意合并访问
Block内线程需要共享数据	共享内存	速度快，可编程控制
线程私有变量	寄存器	零延迟，但别用太多
所有线程读同一份数据	常量内存	广播效率极高
二维空间局部性访问	纹理内存	硬件缓存+插值

嗯，最后说一句。内存管理没有银弹。我见过有人为了用共享内存，把代码写得极其复杂，结果性能还不如直接用全局内存。所以我的建议是：先跑通，再优化，用profiler说话。

用nvprof或者Nsight Compute看看实际的带宽利用率、共享内存bank conflict、寄存器溢出情况。数据不会骗人。

---

公众号:蓝海资料掘金营，微信deep3321