4. 全局内存访问优化：合并访问

全局内存，说白了就是GPU里那块最大的显存。所有线程都能访问它，但速度嘛...嗯，得看你怎么用。我刚开始写CUDA程序时，总觉得全局内存就是个大仓库，随便怎么搬都行。结果一跑性能测试，直接傻眼了——带宽利用率连10%都不到。

为什么会这样？因为全局内存有个脾气：它喜欢连续的数据访问。你让线程们东一榔头西一棒子地去读数据，它就会罢工给你看。今天我们就来聊聊怎么哄好它。

4.1 合并访问原理

合并访问（Coalesced Access）这个概念，我理解起来其实很简单：相邻的线程访问相邻的内存地址。你想想看，GPU里有成百上千个线程同时跑，如果它们都去读内存，那总得有个高效的调度方式吧？

硬件层面是这样的：当warp（32个线程）执行一次内存加载指令时，内存控制器会把这32个请求合并成尽可能少的内存事务。如果访问是连续的、对齐的，一个事务就能搞定。如果访问是散乱的，那就得拆成好几个事务，效率自然就下来了。

我在项目中遇到过这样一个场景：一个简单的向量加法，我一开始用了一维数组，线程索引直接映射到数组下标。结果性能还不错。后来换了数据结构，把数组变成了结构体数组（AoS），性能直接掉了40%。排查了半天，才发现是内存访问模式变了。

4.2 对齐与跨步访问

对齐访问，说白了就是内存地址要规整。比如你读一个128字节的段，起始地址最好是128的倍数。CUDA里，cudaMalloc默认返回的地址就是256字节对齐的，这点倒不用太担心。

真正坑人的是跨步访问（Strided Access）。来看个例子：

// 跨步访问 - 性能差
__global__ void strideAccess(float* data, int stride) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val = data[idx * stride];  // 每个线程跳着读
}

// 连续访问 - 性能好
__global__ void coalescedAccess(float* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val = data[idx];  // 相邻线程读相邻地址
}

跨步访问有多伤？我做过测试：步长为1时，带宽利用率能到80%以上。步长变成2，直接掉到40%。步长再大点，10%都悬。所以啊，设计数据结构时，一定要把跨步访问这个坑考虑进去。

4.3 使用cudaMallocPitch与cudaMalloc3D

说到二维数组，就不得不提cudaMallocPitch。这个函数我特别喜欢，因为它自动帮你处理了对齐问题。

二维数组在显存里怎么存？如果你直接用cudaMalloc分配width * height * sizeof(float)，那每一行的起始地址可能不对齐。而cudaMallocPitch会返回一个pitch值，保证每行都对齐到512字节（具体值取决于设备）。

float* d_data;
size_t pitch;
cudaMallocPitch(&d_data, &pitch, width * sizeof(float), height);

// 访问时用pitch计算偏移
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float val = d_data[row * pitch / sizeof(float) + col];

注意这里pitch的单位是字节，所以要除以sizeof(float)才能得到元素偏移。这个细节我一开始也搞错过，调试了半天才发现是pitch用错了。

对于三维数据，cudaMalloc3D就更方便了。它返回一个cudaPitchedPtr结构体，包含了指针、pitch和深度信息：

cudaExtent extent = make_cudaExtent(width * sizeof(float), height, depth);
cudaPitchedPtr d_data;
cudaMalloc3D(&d_data, extent);

// 访问三维数组
float* ptr = (float*)((char*)d_data.ptr + 
                      z * d_data.pitch * height + 
                      y * d_data.pitch + 
                      x * sizeof(float));

知识体系总览

下面这张图，是我梳理的全局内存访问优化的核心逻辑。你看一眼就能明白各个知识点之间的关系：

性能对比一览

为了让你更直观地感受不同访问模式的性能差异，我整理了一张表：

访问模式	带宽利用率	典型场景	优化建议
连续对齐访问	80% - 95%	向量加法、数组运算	保持默认对齐即可
跨步访问（步长2）	40% - 50%	矩阵列访问、图像通道分离	改用SoA布局或共享内存
随机访问	5% - 15%	稀疏矩阵、哈希表	考虑纹理内存或重新组织数据
使用cudaMallocPitch	75% - 90%	二维图像处理	注意pitch的单位是字节

想要更多资料，联系可可