内存访问优化：全局内存合并访问、共享内存bank conflict、常量内存与纹理内存

各位同学，今天我们来聊聊CUDA编程里最核心、也最让人头疼的一块——内存访问优化。

说实话，我见过太多人把核函数写得飞快，结果一跑，性能惨不忍睹。问题出在哪？十有八九是内存访问没搞好。GPU的计算能力再强，如果数据喂不进去，那也是白搭。这就像你开着一辆法拉利，结果堵在早高峰的二环上，再好的引擎也跑不起来。

所以，这一章我们专门来啃这块硬骨头。我会结合我这些年踩过的坑，把全局内存、共享内存、常量内存和纹理内存这几个关键点讲透。

1. 全局内存合并访问

先说说全局内存。这是GPU里最大的存储空间，但也是访问延迟最高的。为什么？因为它离计算单元最远。

但好在，GPU有它的补救措施——合并访问。说白了，就是让相邻的线程访问相邻的内存地址。这样，硬件可以把多次内存请求合并成一次大的传输，效率直接拉满。

我举个例子。假设你有一个float数组，每个线程要读一个元素。如果线程0读data[0]，线程1读data[1]，线程2读data[2]……这就是完美的合并访问。一次内存事务就能搞定。

但如果你让线程0读data[0]，线程1读data[32]，线程2读data[64]……那就完蛋了。每个线程都要单独发一次内存请求，性能直接掉一个数量级。

我在项目中遇到过这样一个案例：一个矩阵转置的核函数，一开始写得比较随意，结果带宽利用率只有10%左右。后来我调整了数据布局，让访问模式对齐到warp的访问粒度，带宽直接飙到了80%。嗯，这就是合并访问的威力。

// 好的访问模式：合并访问
__global__ void good_access(float *A, float *B, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        B[idx] = A[idx] * 2.0f;  // 连续访问
    }
}

// 坏的访问模式：非合并访问
__global__ void bad_access(float *A, float *B, int N) {
    int idx = blockIdx.x + threadIdx.x * gridDim.x;
    if (idx < N) {
        B[idx] = A[idx] * 2.0f;  // 跨步访问，步长为gridDim.x
    }
}

2. 共享内存bank conflict

共享内存是GPU里的“高速缓存”，由程序员手动管理。它的速度比全局内存快得多，但容量有限，而且有个坑——bank conflict。

共享内存被划分成32个bank，每个bank的宽度是4字节。如果同一个warp里的多个线程同时访问同一个bank的不同地址，就会发生bank conflict，导致访问串行化。

你想想看，本来可以并行访问的，结果因为冲突，变成了串行，性能自然就下来了。

为什么会这样？因为共享内存的bank映射是基于地址的。地址对32取模，结果相同的，就落在同一个bank上。如果你的访问模式恰好让多个线程访问同一个bank，那就冲突了。

常见的解决方法是填充（padding）。比如，你有一个32x32的共享内存数组，声明成__shared__ float s[32][32 + 1]。多出来的那一列，就是为了打破bank对齐，避免冲突。

// 有bank conflict的版本
__shared__ float s[32][32];
int tx = threadIdx.x;
int ty = threadIdx.y;
float val = s[tx][ty];  // 如果tx相同，就会冲突

// 无bank conflict的版本（加padding）
__shared__ float s[32][33];  // 注意第二维多了1
int tx = threadIdx.x;
int ty = threadIdx.y;
float val = s[tx][ty];  // 现在不会冲突了

3. 常量内存与纹理内存

除了全局内存和共享内存，CUDA还提供了两种特殊的内存类型——常量内存和纹理内存。它们各有各的用武之地。

3.1 常量内存

常量内存，顾名思义，存的是不会变的数据。它只有64KB，但有个好处：硬件会广播。如果一个warp里的所有线程都访问同一个常量地址，那只需要一次读取，然后广播给所有线程。

这玩意儿特别适合存查找表、滤波器系数、数学常数这类数据。

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

// 在主机端初始化
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);

// 在核函数中使用
__global__ void apply_filter(float *in, float *out, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        out[idx] = in[idx] * filter[idx % 256];  // 所有线程访问同一个filter值
    }
}

3.2 纹理内存

纹理内存，最初是为图形处理设计的。但它用在通用计算里，也有奇效。

纹理内存有两大特点：

• 空间局部性优化：它有自己的缓存，特别适合二维或三维的随机访问模式。
• 硬件插值：可以直接做线性插值、最近邻插值，省去手动计算的麻烦。

我记得有一次做图像处理，需要频繁地访问像素的邻域。用全局内存直接读，缓存命中率很低。后来换成纹理内存，性能直接翻了一倍。说白了，纹理内存就是为这种“不规则但局部”的访问模式量身定做的。

// 声明纹理对象
cudaTextureObject_t texObj;
cudaResourceDesc resDesc;
cudaTextureDesc texDesc;

// 绑定到二维数组
cudaArray *cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
cudaMemcpyToArray(cuArray, 0, 0, hostData, size, cudaMemcpyHostToDevice);

resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = cuArray;

texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModeLinear;  // 线性插值
texDesc.readMode = cudaReadModeElementType;

cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

// 在核函数中使用
__global__ void texture_kernel(cudaTextureObject_t tex, float *out, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        out[y * width + x] = tex2D<float>(tex, x + 0.5f, y + 0.5f);  // 带插值的访问
    }
}

知识体系总览

为了让你更直观地理解这几种内存的关系和适用场景，我画了一张图。你可以把它当作本章的“导航地图”。

这张图把四种内存的特点、优化要点和注意事项都列出来了。你可以把它当作一个速查表，写代码的时候对照着看。

好了，这一章的内容就到这里。内存访问优化是个细活，需要你在实践中不断调试和体会。别怕踩坑，我当年也是从坑里爬出来的。记住一句话：数据怎么放，决定了你的核函数能跑多快。

想要更多资料，联系可可