高效内存访问模式：合并访问、共享内存与只读缓存

各位同学，今天我们来聊聊GPU计算里最核心、也最容易被忽视的一个话题——内存访问模式。说实话，我见过太多人把核函数写得天花乱坠，结果一跑性能惨不忍睹。问题出在哪？十有八九是内存访问没搞好。

GPU的算力其实很猛，但它的内存系统是个「娇气」的东西。你访问得对，它给你满带宽；你访问得不对，它直接给你打三折。嗯，今天我就把这三个关键点掰开揉碎了讲清楚。

合并访问（Coalesced Access）原理

先说说合并访问。这个概念说白了就是：让同一个warp里的32个线程，访问连续的内存地址。

为什么这么重要？因为GPU的内存控制器一次只能处理一个128字节的缓存行（cache line）。如果32个线程各自访问不同的128字节区域，那就要发起32次内存事务。但如果它们访问的是同一块连续区域，一次事务就搞定了。

我举个例子。假设我们要处理一个float数组，每个线程处理一个元素：

// 合并访问 —— 推荐写法
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];  // 线程0访问data[0]，线程1访问data[1]...

// 非合并访问 —— 千万别这么写
int idx = threadIdx.x * gridDim.x + blockIdx.x;
float val = data[idx];  // 同一个warp里的线程访问地址间隔很大

你看，第一种写法里，线程0访问data[0]，线程1访问data[1]，地址是连续的。第二种写法呢？同一个warp里的线程，blockIdx.x不同，地址直接跳了gridDim.x个元素。这就不连续了。

我在项目中遇到过一个问题：一个图像滤波的核函数，性能死活上不去。后来一查，发现是线程到像素的映射写反了。改了一行代码，性能直接翻了三倍。你想想看，就一行代码的事。

共享内存（Shared Memory）与Bank Conflict

共享内存是GPU里的一块「宝贝」——它比全局内存快两个数量级。但它有个脾气：被分成了32个bank。

每个bank每个时钟周期只能响应一个请求。如果同一个warp里的多个线程访问同一个bank的不同地址，就会发生bank conflict，请求被串行化。

我画了一张图，帮你理解bank conflict是怎么发生的：

怎么避免bank conflict？我个人的习惯是：在共享内存声明时加一个padding。

// 有bank conflict的写法
__shared__ float data[32][32];  // 每行32个float，正好占满32个bank

// 无bank conflict的写法
__shared__ float data[32][33];  // 每行33个float，多了一个padding

为什么加一个padding就能解决问题？因为原本每行32个元素，每个元素对应一个bank。加了padding后，每行33个元素，下一行的起始地址偏移了33个bank位置，就不会和上一行冲突了。

只读缓存（ReadOnly Cache）与常量内存

最后聊聊只读缓存和常量内存。这两个东西，说白了都是「加速只读数据访问」的机制，但适用场景完全不同。

只读缓存（ReadOnly Cache）

只读缓存是Kepler架构引入的，专门用来缓存那些只读且访问模式不规则的数据。比如，你要查一个查找表，每个线程查的位置都不一样——这种场景下，只读缓存比L1缓存更高效。

怎么用？有两种方式：

1. 用 __ldg() 内建函数：显式告诉编译器走只读缓存路径
2. 用 const __restrict__ 修饰指针：让编译器自动推断

// 方式1：显式使用只读缓存
__global__ void kernel(const float* __restrict__ data) {
    int idx = threadIdx.x;
    float val = __ldg(&data[idx]);  // 强制走只读缓存
}

// 方式2：让编译器自动推断
__global__ void kernel(const float* __restrict__ data) {
    int idx = threadIdx.x;
    float val = data[idx];  // 编译器可能自动使用只读缓存
}

我个人更推荐方式1，因为显式声明不会让编译器「猜」你的意图。我曾经在一个项目里遇到过，同样的代码在不同架构的GPU上表现差异很大，后来发现是编译器自动优化策略不同导致的。

常量内存（Constant Memory）

常量内存是另一种只读内存，但它有个特点：同一个warp里的所有线程必须访问同一个地址，才能发挥最大效率。

为什么？因为常量内存只有8KB，但它有专门的广播机制。如果32个线程都访问同一个常量地址，硬件会一次广播给所有线程。但如果每个线程访问不同的地址，那就退化成串行访问了。

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

// 主机端赋值
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 核函数中使用
__global__ void apply_filter(float* output, const float* input) {
    int idx = threadIdx.x;
    float val = input[idx];
    // 所有线程都访问同一个filter数组
    output[idx] = val * filter[0] + val * filter[1];
}

嗯，这里要注意：常量内存的更新开销很大。每次调用 cudaMemcpyToSymbol 都会同步所有正在执行的核函数。所以，尽量在核函数启动前一次性把常量数据准备好。

三种内存访问模式的对比

特性	合并访问	共享内存	只读缓存/常量内存
适用场景	全局内存的连续访问	线程间数据共享、重用	只读数据、查找表
延迟	高（~400周期）	低（~20周期）	中（~100周期）
容量	大（显存大小）	小（48KB~164KB）	极小（8KB~48KB）
主要问题	非合并访问导致带宽浪费	Bank conflict导致串行化	常量内存广播限制

最后总结一下我的经验：写GPU代码时，先想清楚数据访问模式，再动手写代码。合并访问是基础，共享内存是加速器，只读缓存和常量内存是锦上添花。把这三点搞明白了，你的CUDA算子性能至少能提升一个数量级。

想要更多资料，联系可可