全局内存（Global Memory）：显存的基础，容量最大但延迟最高
全局内存，说白了就是咱们常说的显存。它是GPU里容量最大的存储区域，所有线程都能访问。但代价也很明显——延迟最高。
我刚开始接触CUDA时，总觉得全局内存随便用就行。直到有一次，一个简单的矩阵乘法，我直接拿全局内存硬怼，结果性能惨不忍睹。嗯，从那以后我才真正重视起这个“大块头”的脾气。
全局内存的本质
全局内存位于GPU芯片外部，通过高速总线与SM连接。你可以把它想象成一个巨大的仓库，所有线程都能从里面取货、放货。但每次取货都要走一段很长的路——这就是延迟的来源。
关键特性：
- 容量最大：从几GB到几十GB不等，取决于显卡型号
- 延迟最高：典型延迟在400-800个时钟周期
- 所有线程可访问：不分SM、不分block，全局共享
- 持久性：数据在kernel执行期间一直存在
为什么全局内存这么慢？
你想想看，GPU芯片内部跑在1-2GHz的频率上，而全局内存的DRAM颗粒跑在几GHz。但问题不在于频率，而在于物理距离和协议开销。
从SM发出一个内存请求，需要经过：
- SM内部的load/store单元
- L1缓存（如果miss）
- L2缓存（如果miss）
- 内存控制器
- DRAM颗粒
- 数据原路返回
这一来一回，几百个周期就没了。我在项目中遇到过，一个kernel如果频繁访问全局内存，性能瓶颈几乎100%在这里。
全局内存的访问模式
全局内存虽然慢，但有个重要的优化手段——合并访问。说白了，就是让相邻的线程访问相邻的地址。
为什么合并访问这么重要？因为GPU的内存控制器一次读取一个128字节的cache line。如果32个线程（一个warp）访问的地址是连续的，那一次内存事务就能搞定。否则，就要多次事务，性能直接打折扣。
我的经验：写kernel时，我习惯让线程索引和内存地址一一对应。比如thread 0访问地址0，thread 1访问地址1，以此类推。这样基本能保证合并访问。
代码示例：合并访问 vs 非合并访问
// 合并访问（推荐）
__global__ void coalesced_access(float *A, float *B, float *C, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
C[idx] = A[idx] + B[idx]; // 相邻线程访问相邻地址
}
}
// 非合并访问（避免）
__global__ void strided_access(float *A, float *B, float *C, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
C[idx] = A[idx * 2] + B[idx * 2]; // 跨步访问，破坏合并
}
}你看，第一个例子中，thread 0访问A[0]，thread 1访问A[1]，完全连续。第二个例子中，thread 0访问A[0]，thread 1访问A[2]，中间跳过了A[1]。这会导致内存控制器每次只用到一半的带宽。
全局内存的带宽
全局内存的带宽其实不低。以NVIDIA A100为例，HBM2e显存带宽超过2TB/s。但问题是，这个带宽是所有SM共享的。如果多个SM同时发起大量内存请求，带宽就成了瓶颈。
| GPU型号 | 显存类型 | 带宽 | 容量 |
|---|---|---|---|
| A100 | HBM2e | 2.0 TB/s | 40/80 GB |
| V100 | HBM2 | 900 GB/s | 16/32 GB |
| RTX 4090 | GDDR6X | 1.0 TB/s | 24 GB |
我曾经优化过一个图像处理kernel，一开始带宽利用率只有30%。后来通过调整数据布局和访问模式，把利用率提到了80%以上。性能直接翻了2倍多。
全局内存的避坑指南
我曾经踩过的坑：
- 频繁的原子操作：全局内存的原子操作非常慢，因为要保证所有线程的可见性。能用共享内存的原子操作就别用全局的。
- 未对齐的访问：如果地址不是32字节对齐，内存控制器会多读一次。我习惯用
cudaMalloc分配的内存，它默认对齐到256字节。 - 过度使用全局内存：能放共享内存或寄存器的数据，就别放全局内存。我见过有人把循环变量都放全局内存，那性能简直灾难。
全局内存的生命周期
全局内存的生命周期由程序员控制。用cudaMalloc分配，用cudaFree释放。数据在kernel执行期间一直存在，直到你显式释放或程序结束。
这里有个小技巧：我习惯在程序开始时一次性分配好所有需要的全局内存，而不是在循环里反复分配释放。因为cudaMalloc本身也有开销，频繁调用会影响性能。
全局内存与缓存的配合
全局内存不是孤立存在的。它和L1、L2缓存配合工作。如果一个数据被频繁访问，它会被缓存到L1或L2中，后续访问就快多了。
但要注意，缓存是透明的。你不能强制数据留在缓存里，也不能强制它被踢出。NVIDIA提供了__ldg()指令，可以绕过L1直接读全局内存，适合那些只读一次的大数组。
我的建议：对于只读且访问模式规律的数据，用const __restrict__修饰指针。这能告诉编译器数据不会别名，有助于生成更高效的代码。
全局内存的典型应用场景
- 输入数据：从CPU传过来的数据，先放全局内存
- 输出结果：kernel计算完的结果，写回全局内存
- 大数组：超过共享内存容量的数据，只能放全局内存
- 持久化数据：需要在多个kernel之间共享的数据
说白了，全局内存就是GPU的“主存”。它容量大、延迟高，但所有线程都能访问。用好它的关键是——减少访问次数，优化访问模式。
这张图展示了全局内存的访问路径。每次内存请求都要经过多层缓存和控制器。如果L1命中，延迟只有几十个周期。如果一路miss到全局内存，那就是几百个周期。所以，尽量让数据留在缓存里，这是性能优化的核心思路。
好了，全局内存就讲这么多。记住它的特点：容量大、延迟高、所有线程可访问。用好合并访问，避开原子操作，配合缓存策略，你就能把这块“大块头”用得得心应手。