3、内存访问优化：全局内存合并访问、共享内存bank冲突、常量内存与纹理内存使用

内存访问，说白了就是GPU性能的命门。我见过太多项目，核函数写得花里胡哨，结果一跑profile，发现大部分时间都卡在内存读写上。今天咱们就把这块硬骨头啃下来。

3.1 全局内存合并访问

全局内存是GPU里最大的存储空间，但也是访问最慢的。不过，如果你能让它「合并访问」，速度能提升一个数量级。

什么是合并访问？

当一个warp（32个线程）同时访问全局内存时，硬件会把这些访问合并成尽可能少的内存事务。如果相邻线程访问相邻的地址，那就是一次完美的合并访问。反之，如果线程访问的是分散的地址，那就得拆成多个事务，性能直接崩掉。

举个例子，假设我们要处理一个二维数组 data[N][M]：

// ❌ 糟糕的访问模式：按列访问
float val = data[threadIdx.x][col];  // 线程0访问data[0][col]，线程1访问data[1][col]... 地址不连续

// ✅ 合并访问：按行访问
float val = data[row][threadIdx.x];  // 线程0访问data[row][0]，线程1访问data[row][1]... 地址连续

我在项目中遇到过一位同事，他写矩阵乘法时用了列优先遍历，结果性能只有理论值的5%。改成行优先后，直接快了20倍。嗯，这就是合并访问的威力。

3.2 共享内存bank冲突

共享内存是片上的SRAM，速度比全局内存快两个数量级。但它有个坑——bank冲突。

什么是bank？

共享内存被划分成32个bank（对于计算能力3.5+的设备）。每个bank宽度4字节。同一时钟周期内，不同线程可以同时访问不同bank。但如果两个线程访问同一个bank的不同地址，就会发生bank冲突，访问会被串行化。

访问模式	性能表现	说明
无冲突	最佳	每个线程访问不同bank
广播	良好	所有线程访问同一bank的同一地址
n路冲突	降为1/n	n个线程访问同一bank的不同地址

举个例子，共享内存声明为 __shared__ float sdata[32][32]：

// ❌ 2路bank冲突：线程0和线程16都访问bank0
float a = sdata[threadIdx.x][0];  // 线程0→bank0, 线程16→bank0

// ✅ 无冲突：每个线程访问不同bank
float a = sdata[threadIdx.x][threadIdx.x];  // 线程0→bank0, 线程1→bank1...

为什么会这样？因为共享内存的地址到bank的映射是 bank = (byte_address / 4) % 32。当步长是32的倍数时，就会踩到同一个bank上。

3.3 常量内存与纹理内存

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

常量内存

常量内存只有64KB，但它的读取速度非常快——前提是warp内所有线程读取的是同一个地址。硬件会把这个值广播给所有线程，零开销。

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

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

// 核函数中使用
__global__ void conv_kernel(float* input, float* output) {
    float val = filter[blockIdx.x];  // 所有线程读同一个地址 → 广播
}

我个人习惯把那些所有线程都用到的参数（比如滤波器系数、查找表）放到常量内存里。64KB虽然小，但装个几百个浮点数绰绰有余。

纹理内存

纹理内存最初是为图形学设计的，但它的一些特性在通用计算中也很有用：

• 硬件插值：支持线性插值、双线性插值，适合图像处理
• 边界处理：支持clamp、wrap等寻址模式，省去边界判断
• 缓存优化：有独立的纹理缓存，适合空间局部性强的访问模式

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

// 绑定纹理
cudaArray* cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);
cudaMemcpyToArray(cuArray, 0, 0, host_data, size, cudaMemcpyHostToDevice);
cudaBindTextureToArray(tex, cuArray);

// 核函数中使用
__global__ void texture_kernel(float* output) {
    float val = tex2D(tex, x, y);  // 硬件自动处理边界
}

3.4 知识体系总览

下面这张图总结了本章的核心逻辑：

说白了，内存优化的本质就是「让数据靠近计算，让访问模式匹配硬件」。全局内存要合并，共享内存要避免冲突，常量内存和纹理内存要用对场景。你想想看，这些优化做下来，性能翻个几倍是很常见的事。

嗯，这一章的内容就到这里。记住，纸上得来终觉浅，绝知此事要躬行。建议你拿自己的项目跑一下profile，看看内存访问模式到底长什么样。