第四章内存模型（下）：访问模式与带宽实战

好，咱们接着聊内存。上一章我把全局内存、共享内存、寄存器这些硬件单元讲了一遍，算是把「有什么」说清楚了。这一章咱们重点解决「怎么用」——说白了，就是怎么让数据在内存里跑得飞快。

我个人习惯把内存访问比作「仓库取货」。你想想看，如果仓库管理员每次只拿一个零件，来回跑一趟，效率肯定低。但如果他一次拿一整箱，甚至提前把相邻的货架都搬过来，那速度就完全不一样了。CUDA 的内存优化，本质上就是在做这件事。

4.1 合并访问：让数据「组团」来

先问一个问题：GPU 从全局内存读数据，一次能读多少？

答案是 32 字节、64 字节、128 字节，取决于数据类型和硬件。但关键是——这 32 字节必须来自连续的地址。如果线程 0 读地址 A，线程 1 读地址 A+4，线程 2 读地址 A+8……那 GPU 就能一次把这一整块搬过来，这叫「合并访问」。

我在项目中遇到过最典型的反例：

// 糟糕的访问模式：按列访问
__global__ void bad_access(float *data, int width) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    float val = data[row * width + col];  // 连续线程访问不连续地址
}

这段代码里，相邻线程访问的地址间隔了 width 个元素。如果 width=1024，那间隔就是 4096 字节。GPU 每次只能读 128 字节，结果就是每个线程都要单独发一次请求——效率直接掉到 1/32。

核心原则：让同一个 warp 里的 32 个线程，访问连续的内存地址。这样一次内存事务就能喂饱整个 warp。

正确的写法应该是：

// 好的访问模式：按行访问
__global__ void good_access(float *data, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[row * width + col];  // 连续线程访问连续地址
}

嗯，这里要注意：合并访问只对全局内存有效。共享内存、寄存器、常量内存各有各的规矩。

4.2 Bank Conflict：共享内存的「堵车」问题

共享内存比全局内存快得多，但它也有自己的脾气。共享内存被分成 32 个 bank（在较新架构上是 32 个），每个 bank 宽度 4 字节。如果同一个 warp 里的多个线程，同时访问同一个 bank 的不同地址，就会发生 bank conflict——说白了就是堵车了。

我曾经调试过一个矩阵转置的 kernel，性能死活上不去。后来用 nvcc 的 --ptxas-options=-v 一看，bank conflict 率高达 50%。

看个例子：

// 有 bank conflict 的访问
__shared__ float s_data[32][32];
int idx = threadIdx.x;
float val = s_data[idx][idx];  // 所有线程访问同一列，bank 冲突

这里每个线程访问的是 s_data[0][0]、s_data[1][1]……这些地址虽然不同，但它们落在同一个 bank 上（因为列索引相同）。32 个线程抢一个 bank，性能直接掉到 1/32。

小技巧：加一个 padding 就能解决。比如把数组声明成 __shared__ float s_data[32][33]，多出来的一列让地址错开，bank conflict 就消失了。

为什么会这样？因为 bank 的映射规则是：bank_index = (byte_address / 4) % 32。加 padding 后，同一列的地址除以 4 再取模，结果就不一样了。

4.3 内存带宽计算：理论 vs 实际

聊完访问模式，咱们来算算账。GPU 的内存带宽，理论值可以从规格书上查到。比如 RTX 3090 的显存带宽是 936 GB/s。但实际能跑多少？

我建议用这个公式估算：

实际带宽 = (读取字节数 + 写入字节数) / 执行时间

写个简单的 benchmark：

__global__ void bandwidth_test(float *in, float *out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = in[idx] * 2.0f;
    }
}

// 调用后测量时间
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
bandwidth_test<<<grid, block>>>(d_in, d_out, n);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms;
cudaEventElapsedTime(&ms, start, stop);

float bandwidth = (n * sizeof(float) * 2) / (ms / 1000.0) / 1e9;  // GB/s

注意：实际带宽通常只有理论值的 60%-80%。如果低于 50%，说明你的访问模式有问题，或者 kernel 太轻量，被访存延迟掩盖了。

我见过最夸张的一次，有人写了个 kernel，实际带宽只有理论值的 5%。查了半天，发现他把全局内存当共享内存用了——每个线程都去读全局内存，而且地址还不连续。嗯，这种坑踩过一次就记住了。

4.4 知识体系总览

下面这张图把本章的核心逻辑串起来了。你可以把它当作一个检查清单：

4.5 避坑指南

最后分享几个我踩过的坑：

别迷信「共享内存一定快」。如果 bank conflict 严重，共享内存可能比全局内存还慢。我见过有人用共享内存做矩阵乘法，结果 bank conflict 导致性能还不如直接用全局内存。
注意数据类型对齐。比如 float4 要求 16 字节对齐，如果地址不对齐，硬件会自动拆成多次访问，性能直接腰斩。
小心「伪共享」。多个线程写同一个 cache line 的不同位置，虽然地址不同，但 cache line 只有一份，会导致写冲突。解决办法是加 padding 让每个线程独占一个 cache line。

嗯，内存这块内容确实多，但核心就一句话：让数据流动起来，别让任何一级存储成为瓶颈。下一章咱们会把这些技巧用到实际算子优化里，到时候你就能感受到「优化前后性能翻倍」的快感了。

本章核心：合并访问 + 无 Bank Conflict + 带宽计算 = 高效内存访问。记住这三个关键词，面试和实战都能用上。

想要更多资料，联系可可

📚 课程目录

第四章 内存模型（下）：访问模式与带宽实战

4.1 合并访问：让数据「组团」来

4.2 Bank Conflict：共享内存的「堵车」问题

4.3 内存带宽计算：理论 vs 实际

4.4 知识体系总览

4.5 避坑指南

第四章内存模型（下）：访问模式与带宽实战