全局内存访问：全局内存特性、合并访问原则、对齐与跨步访问、内存事务优化

大家好，我是你们的老朋友。今天我们来聊聊CUDA编程里一个绕不开的话题——全局内存访问。说实话，我见过太多新手写的核函数，计算逻辑没问题，但跑起来就是慢。为什么？十有八九是内存访问模式出了问题。

全局内存，说白了就是GPU里那块最大的显存。所有线程都能读写它，容量大，但延迟也高。我刚开始做CUDA优化时，总觉得这玩意儿就是个黑盒子，直到有一次调优一个矩阵转置的核函数，性能死活上不去，折腾了两天才发现是内存访问没对齐。嗯，从那以后我再也不敢小看全局内存了。

全局内存的特性

先说说全局内存的几个关键特性：

• 容量大：通常几个GB到几十个GB，能装下大模型、大矩阵
• 延迟高：几百个时钟周期，比共享内存慢两个数量级
• 带宽大：现代GPU的全局内存带宽能到几百GB/s，但前提是你得会用
• 所有线程可见：这是最常用的内存空间，也是性能瓶颈的重灾区

我个人习惯把全局内存想象成一个巨大的仓库。你从仓库里搬东西，一次搬得越多、搬得越整齐，效率就越高。如果东一榔头西一棒子，那效率肯定上不去。

合并访问原则

合并访问，这是全局内存优化的核心原则。什么意思呢？

GPU在执行核函数时，会把32个线程组成一个warp。这32个线程同时访问全局内存时，硬件会尽量把这些访问合并成一次或少数几次内存事务。如果访问模式是连续的、对齐的，那一次事务就能搞定。如果访问是随机的、跨步的，那就得拆成很多次事务，带宽利用率直线下降。

举个例子，假设我们有一个float数组，每个线程读一个元素：

// 好的访问模式：连续访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];  // 线程0读data[0]，线程1读data[1]...

// 坏的访问模式：跨步访问
int idx = threadIdx.x * 2;
float val = data[idx];  // 线程0读data[0]，线程1读data[2]...

我在项目中遇到过这样一个案例：一个图像处理核函数，原本用第二种方式写，带宽利用率只有20%左右。改成连续访问后，性能直接翻了3倍。你想想看，只是改了一下索引计算方式，效果就这么明显。

对齐与跨步访问

对齐，指的是内存访问的起始地址要满足一定的对齐要求。对于全局内存，我个人建议尽量保证访问地址是32字节对齐的。为什么呢？因为GPU的内存事务是以32字节、64字节或128字节为单位的。

内存事务大小	适用场景	对齐要求
32字节	1字节或2字节数据类型	32字节对齐
64字节	4字节数据类型（如float）	64字节对齐
128字节	16字节数据类型（如float4）	128字节对齐

跨步访问，说白了就是线程访问的地址之间有间隔。比如：

// 跨步访问示例
int stride = 4;
int idx = threadIdx.x * stride;
float val = data[idx];

这种情况下，warp内的32个线程访问的地址是0, 4, 8, 12... 它们之间隔着3个元素。硬件没法合并这些访问，只能拆成多次事务。我曾经优化过一个三维卷积的核函数，就是因为跨步访问导致性能惨不忍睹。后来通过重新组织数据布局，把跨步访问变成了连续访问，性能提升了5倍。

内存事务优化

内存事务，就是GPU从全局内存读取或写入数据的最小操作单元。优化内存事务，说白了就是让每次事务都物尽其用，别浪费带宽。

几个实用的优化技巧：

• 使用向量化数据类型：比如用float4代替float，一次事务能读4个float
• 保证对齐：用cudaMalloc分配的内存默认是256字节对齐的，但你自己做偏移时要小心
• 避免非对齐访问：非对齐访问会导致一次事务被拆成两次，带宽减半
• 减少跨步：尽量让warp内的线程访问连续地址

举个例子，用float4优化内存访问：

// 未优化版本
float val = data[idx];

// 优化版本：使用float4
float4 val4 = ((float4*)data)[idx / 4];
float val = val4.x;  // 或者val4.y, val4.z, val4.w

为什么这样能优化？因为float4是16字节对齐的，一次128字节的事务能读8个float4，也就是32个float。刚好一个warp的32个线程一次搞定。我做过测试，在带宽受限的核函数里，用float4优化后性能能提升20%-50%。

最后，我画了一张图来总结全局内存访问的核心逻辑：

这张图把整个逻辑串起来了。你从线程访问模式出发，判断是否连续且对齐。如果是，恭喜你，一次事务搞定，带宽利用率拉满。如果不是，那就得拆成多次事务，性能自然就下来了。

好了，关于全局内存访问的内容就聊到这里。记住一句话：让warp内的线程访问连续、对齐的内存地址。这是CUDA优化的基本功，也是性价比最高的优化手段之一。

想要更多资料，联系可可