CUDA编程模型：全局内存

CH-18-02-003·约 15 分钟·AI硬件与算力·更新 2026-08-11

想象你管理着一座位于城市边缘的巨大仓库。仓库里货架一眼望不到头，容量惊人，但每次提货都要开车往返；这就是 GPU 的全局内存（global memory）。GPU 里有成千上万个“订单拣货员”——线程。线程们的活很简单：从仓库某个货位取一件货物，搬回工位。程序员的任务不是教他们怎么搬，而是设计提货单上的货位顺序。如果 32 个拣货员恰好站在同一排、货位连续，搬运车一次就能把一托盘的货拉到他们面前，这就是合并访问；如果他们东一个西一个，搬运车就得来回跑好几趟，这就是非合并访问的性能惩罚。本篇文章将带你认识这座大仓库的货架结构、搬运规则和摆放技巧，让 GPU 的全局内存访问又快又省。

岗位关联

学懂全局内存的访存特性与合并访问，是写出高性能 CUDA kernel 的基本功；也是 AI 推理引擎、训练框架在 GPU 上做数据搬运时绕不开的优化关口。

可视化演示
1. 存储层次定位：城外的大仓库
2. 访存事务原理：搬运车一次能搬多少
3. 合并访问模式：让一趟车装完一整组货
4. 非合并代价：四处散落的提货单
5. 数据布局优化：重新摆放货架
6. 性能评估方法：给仓库门口装个计数器
7. 优化实践清单：贴在墙上的提货规则
AI 应用场景
Python 演示
常见误区
面试题

可视化演示

展示GPU全局内存的合并访问与非合并访问模式、带宽利用率对比及数据布局优化，交互切换查看内存访问效率关键因素。

本节配有交互动效（进入视口后自动加载，约 3KB）。

合并访问：连续线程访问连续地址 Warp中32线程访问连续地址触发单次内存事务非对齐跨步访问导致多次事务带宽浪费 Coalesced = 1 transaction/Warp 合并访问是GPU性能基石

1. 存储层次定位：城外的大仓库

把 GPU 想象成一个忙得不可开交的仓库园区。寄存器是工人手上的口袋，共享内存是工位旁的小货架，而全局内存则是园区外面那座一眼望不到边的巨型仓库。它容量最大，动辄几个 GB 甚至几十 GB；但离计算单元最远，取一次货往往要几百个时钟周期。正因为地方远，工人们不会把每件货都往那里搬，只把真正需要共享的大块数据放在那里。

在 CUDA 的存储体系里，全局内存在设备端，所有线程、所有线程块都能读能写。也就是说，任何一个“拣货员”都可以拿着提货单去城外仓库取货。它不像共享内存那样由程序员显式管理，也不像寄存器那样属于某个线程私有；它更像是公共货场。正因为它“又大又公共”，访问规则就变得格外重要：去得越频繁、去法越乱，开销就越大。

现代 GPU 会给全局内存配 L2 缓存，但第一次访问仍然要跨过漫长的内存总线。你可以把缓存想象成仓库门口的几个临时周转箱，能装下最近一批货，却装不下整座仓库。因此，理解“提货单怎么开”，比盲目依赖缓存更重要。

2. 访存事务原理：搬运车一次能搬多少

仓库里最值钱的设备不是货架，而是那辆记忆体搬运车。它每次出车，不会只为你捎一个盒子，而是至少搬一个托盘。CUDA 硬件读取全局内存时，也不是一个字节一个字节地搬，而是按“内存事务”批量搬运。一个事务通常覆盖 32B 的扇区，连续 4 个扇区组成 128B 的缓存行。换句话说，只要你的地址落在同一个托盘上，这趟车就能顺便把周围的东西一起带回来。

这个机制带来了一个关键结论：决定性能的不是“你请求了多少字节”，而是“你的请求分布在多少个事务里”。如果 32 个线程的地址分散在 8 个 128B 段里，搬运车就得跑 8 趟；哪怕每趟只用到其中 4 个字节，整车货物也算被搬运了。

用公式表示搬运效率：

所以，想让搬运车跑得快，要么让每趟车都尽量装满，要么减少出车次数。这正是合并访问要做的事。

\displaystyle B=\frac{D}{T}

其中 $D$ 是实际搬运的有效字节数，$T$ 是耗时；这个值就是全局内存访问的有效带宽。

有效字节数 $D$真正被程序用到的数据量。比如 32 个线程各读 1 个 float，$D=32\times4=128$ 字节。
$\displaystyle D = N_{\text{threads}} \times \text{bytes per thread}$
耗时 $T$搬运车从出发到把货放到工位前的总时间，包括排队、传输和返回。
$\displaystyle T = T_{\text{latency}} + T_{\text{transfer}}$
带宽 $B$单位时间能送达的有效数据量，反应仓库这条路的真实吞吐能力。
$\displaystyle B = \frac{D}{T}$

3. 合并访问模式：让一趟车装完一整组货

现在你给 32 个拣货员发提货单。最好的情况是什么？他们按编号排好队，第 0 号取 0 号货位，第 1 号取 1 号货位，第 2 号取 2 号货位……这样 32 个货位紧紧挨在一起，搬运车只需开进那一排，把一整托盘货放下，所有人一次拿完。这就是合并访问（coalesced access）。

在 CUDA 里，一个 warp 有 32 个线程。当线程 i 访问的地址是：

如果每个元素大小 $s$ 与线程步长一致，且基址对齐，这 32 个地址就会落在同一个对齐的 128B 区间里。

更严格地说，合并访问要求：

其中 $S$ 是事务段大小，$C$ 是同一个段编号。

还有一种特殊情况：如果所有线程读同一个地址，硬件也支持“广播”，一次就把同一份数据发给所有人。这就像 32 个工人同时要 0 号货位上的那张说明书，搬运车把说明书复印一遍就行。

\displaystyle A_i = \text{base} + i \times s

\displaystyle \left\lfloor \frac{A_i}{S} \right\rfloor = C, \quad \forall i \in \{0,1,\dots,31\}

取整符号 $\lfloor \cdot \rfloor$向下取整，表示“看这个地址属于哪一段托盘”。
$\displaystyle \text{segment} = \left\lfloor \frac{\text{address}}{\text{segment size}} \right\rfloor$
段大小 $S$硬件搬运的基本单位，常取 128 字节。地址越界到另一个段，就要多开一趟车。
$\displaystyle S = 128\ \text{bytes}$
同一个常数 $C$所有线程的段号完全相同，说明它们的目标货物装在同一辆车上。
$\displaystyle C = \text{same segment id for all threads}$
对所有线程成立只要有一个线程掉到别的段，整车货物就可能装不下。
$\displaystyle \forall i \in \{0,1,\dots,31\}$

4. 非合并代价：四处散落的提货单

如果提货单上的货位是“第 0、8、16、24 号……”这样跳着写，拣货员们就会分散在仓库各个角落。搬运车不可能同时服务所有人，只能按片区一趟趟跑：先送 0 号，再送 8 号，再送 16 号……每一趟都只装了一小部分有用的货，其余位置全部空着。这就是非合并访问。

在 CUDA 中，非合并没有语法错误，程序结果仍然正确，但性能会明显下降。原因很直接：访问区间被切成更多个 128B 段，事务数量增加。假设 32 个线程每隔 8 个 float 取一个数据，地址范围就从 128B 膨胀到约 1KB，可能触发 8 个甚至更多事务。硬件实际搬运的数据量，可能比你真正需要的多出好几倍。

这种代价经常出现在“跨步访问”里，比如结构体数组只取其中一个字段：线程 i 读第 i 个结构体的 x 字段，相邻地址之间隔了 12 或 16 字节，货位不再连续。于是每个线程都需要单独一趟车，带宽被大量浪费，延迟也随之上升。更麻烦的是，这种问题用眼睛看代码未必看得出来，只有用 profiler 数事务数才能发现。

一句话总结：地址越散，车次越多；车次越多，有效带宽越低。仓库依然正确，但运费全花在空跑上了。

5. 数据布局优化：重新摆放货架

既然散货位会让搬运车空跑，那仓库管理员能做什么？答案是重新安排货物的摆放方式。CUDA 程序员最常用的一招，是把“结构体数组”（Array of Structures，AoS）改成“数组结构体”（Structure of Arrays，SoA）。

举个例子。你要记录一堆粒子的位置，每个粒子有 x, y, z 三个坐标。如果用 AoS，内存里是这样排的：

如果某个 kernel 只想读所有粒子的 x，线程 0 读地址 0，线程 1 读地址 12，线程 2 读地址 24……货位隔得很远，合并访问彻底失效。

如果改成 SoA，内存变成：

所有 x 紧紧挨在一起，线程按编号连续读，就变成合并访问。

SoA 的地址公式可以写成：

而 AoS 的地址公式是：

两者的差别，就是“按人的订单放货”还是“按货的种类放货”。

除了改布局，还要注意对齐和填充。如果结构体大小不是 4、8、16 的倍数，可以在尾部加 padding，让每个结构体的起始地址对齐到合适的边界。说白了，货架上的箱子不能被台阶卡着，否则搬运车没法一托托盘整整齐齐地叉起来。

\displaystyle A_i^{\text{SoA}} = \text{base}_x + i \times \text{sizeof(float)}

\displaystyle A_i^{\text{AoS}} = \text{base} + i \times \text{sizeof(Particle)}

基址 $\text{base}_x$x 数组的起始地址。SoA 把同名字段放在连续空间里。
$\displaystyle \text{base}_x = \text{address of first x}$
步长 $\text{sizeof(float)}$相邻两个 x 之间的距离是 4 字节，正好让 32 个线程凑满一个 128B 段。
$\displaystyle \text{stride}_{\text{SoA}} = 4\ \text{bytes}$
AoS 步长 $\text{sizeof(Particle)}$相邻结构体之间隔着 12 或 16 字节，导致线程步长大于元素大小。
$\displaystyle \text{stride}_{\text{AoS}} = 12\ \text{bytes}$

text1 行

x0 y0 z0  x1 y1 z1  x2 y2 z2 ...

text1 行

x0 x1 x2 ...  y0 y1 y2 ...  z0 z1 z2 ...

6. 性能评估方法：给仓库门口装个计数器

优化做完了，怎么知道有没有效果？你要在仓库门口装一个计数器，分别记录“真正有用的货”和“搬运车实际搬的货”。对 GPU 程序来说，这就是有效带宽（achieved bandwidth）。

有效带宽的计算很简单：

然后和设备的理论峰值对比：

在实际工作中，NVIDIA Nsight Compute 等 profiler 会直接给出“memory throughput“和”excessive sector loads”之类的指标。看到 sectors per request 远大于理论最小值，就说明搬运车多跑了不少冤枉路。用数据说话，而不是靠猜，是仓库管理员最基本的素养。

\displaystyle B_{\text{achieved}} = \frac{\text{useful bytes}}{\text{time}}

\displaystyle E = \frac{B_{\text{achieved}}}{B_{\text{peak}}} \times 100\%

有效字节 useful bytes程序真正读写的字节数，比如 32 个线程读 32 个 float，有效字节就是 128。
$\displaystyle \text{useful bytes} = N_{\text{threads}} \times \text{element size}$
耗时 timekernel 执行时间，通常用 profiler 或 CUDA event 测量。
$\displaystyle \text{time} = t_{\text{end}} - t_{\text{start}}$
理论峰值 $B_{\text{peak}}$硬件手册给出的最大带宽，相当于传送带最高设计速度。
$\displaystyle B_{\text{peak}} = \text{memory clock} \times \text{bus width}$
效率 $E$把“实际吞吐”和“巅峰吞吐”做比值，用来判断还有多少优化空间。
$\displaystyle E = \frac{B_{\text{achieved}}}{B_{\text{peak}}} \times 100\%$

7. 优化实践清单：贴在墙上的提货规则

最后，把前面所有规则变成一张可以贴墙上的检查清单。

让线程编号与内存地址线性对应：addr = base + threadIdx.x，这是最朴素也最可靠的合并访问写法。

避免跨步访问：如果非跨步不可，先把数据重排成紧凑布局，或者在 kernel 外做一次转置。

优先使用向量类型：float4 一次读 16B，配合对齐基址，32 个线程能覆盖 512B 连续区域，减少指令数。

分配时保证对齐：cudaMalloc 返回的地址通常已按 256B 对齐，但自己在结构体内做偏移时要小心。

善用 L2 缓存：某些访问模式虽不能完全合并，但如果数据在 L2 中被复用，实际惩罚会小很多；可以用 cudaFuncAttributePreferredSharedMemoryCarveout 之类方式调优。

用 profiler 复核：每次改动后都看一眼有效带宽和事务数，而不是只看 kernel 总耗时。

这些规则就像仓库门口的“装车规范”：先看提货单是否连续，再看货架是否对齐，最后用计数器确认车次有没有减少。牢记这条清单，你就不会再让搬运车空跑了。

---

AI 应用场景

场景 1：图像预处理中的像素归一化

阶段：数据预处理与特征工程

解决的问题：图像数据量巨大，如果每个线程读像素时地址散乱，预处理阶段就会成为 pipeline 瓶颈。传统按像素逐点处理的 CPU 方案慢，GPU 方案若忽略合并访问也同样快不起来。

具体案例：在 ImageNet 分类训练前，需要用 CUDA 对 1920×1080 的 RGB 图像做减均值、除方差。每个线程处理一个像素，连续线程读取连续像素字节；如果直接用交错 RGB 布局，访问是连续的，归一化吞吐接近理论峰值。若改为按通道分离的平面布局，则每个通道单独连续，合并效果同样优秀。

场景 2：AI 推理引擎的张量布局优化

阶段：模型推理优化

解决的问题：深度学习张量在内存中有 NCHW 和 NHWC 等不同排布。NCHW 中相邻线程访问同一 channel 的相邻像素是合并的，但某些算子切换 channel 时会跨大步；NHWC 则让空间位置连续的线程访问连续内存，更利于合并。用 SoA/AoS 的思路理解，就是在选一种“货架摆放方式”。

具体案例：TensorRT 在优化卷积网络时，会根据算子选择最合适的张量布局，有时把 NCHW 重排成 NHWC，就是为了让 GPU 访存合并、减少内存事务。推理服务上线后，一次 batch 的 throughput 能提升 20% 以上。

Python 演示

python26 行

# -*- coding: utf-8 -*-# 模拟 CUDA 全局内存的合并与非合并访问 def count_transactions(addresses, segment_size=128):    segments = set()    for addr in addresses:        segments.add(addr // segment_size)    return len(segments) def simulate_access(stride_elems, thread_count=32, elem_size=4):    base = 1024    addresses = [base + i * stride_elems * elem_size for i in range(thread_count)]    transactions = count_transactions(addresses)    useful_bytes = thread_count * elem_size    moved_bytes = transactions * 128    efficiency = useful_bytes / moved_bytes    return transactions, moved_bytes, efficiency def main():    print("stride_elems  transactions  moved_bytes  efficiency")    for stride in [1, 2, 4, 8, 32]:        trans, moved, eff = simulate_access(stride)        print("%5d           %2d            %5d        %.2f%%" % (stride, trans, moved, eff * 100)) if __name__ == "__main__":    main()

目的：用 Python 模拟不同跨步下的全局内存访问，展示合并访问如何减少硬件事务数、提升有效带宽。

segments.add(addr // segment_size)把每个地址归入所属的 128B 段，统计被触碰的不同段数。
addresses = [base + i * stride_elems * elem_size ...]生成 32 个线程的字节地址，stride=1 代表合并访问。
transactions = count_transactions(addresses)一次内存事务对应一个段，段数越多，搬运车出车次数越多。
moved_bytes = transactions * 128硬件实际搬运的字节数，往往远超有效字节。
efficiency = useful_bytes / moved_bytes合并且对齐时接近 1，跨步越大，效率越低。

常见误区

误区我听说“全局内存未缓存”，于是认为所有全局访问都会立刻打到显存，缓存优化没有意义。

正解早期 CUDA 文档强调全局内存不像共享内存那样由软件管理，但现代 GPU 已经有 L2 缓存，部分访问也能被 L1 缓存命中。更重要的是，即使命中缓存，第一次 miss 仍要付出高延迟；合并访问能从源头减少事务数，让冷数据也来得更快。

为何常见：初学者把“软件不可管理”误读成“硬件完全没有缓存”，忽略了现代架构的演进。

误区每个线程读 4 字节，32 个线程一共 128 字节，那内存事务肯定只搬 128 字节，性能一定好。

正解只有当天线访问地址连续且对齐时，128B 才恰好是一个事务。如果地址跨步大、散落在多个 128B 段，硬件会搬运 2 个甚至更多个事务，实际搬运量可能是 256B、512B 甚至更多。

为何常见：人眼习惯用“总请求量”思考，忽略了硬件按段搬运的物理规则。

误区反正程序结果是对的，对齐不对齐无所谓，顶多慢 5%。

正解非对齐访问可能让一个请求横跨两个段，触发额外事务，慢的幅度可能远不止 5%。尤其在带宽受限的 kernel 里，非对齐能把有效带宽砍掉一半甚至更多。性能也是正确性的一部分。

为何常见：CPU 对非对齐访问容忍度较高，GPU 的段式搬运机制放大了这个代价，但从代码上看又很隐蔽。

误区SoA 是“银弹”，只要把结构体数组拆成字段数组，访存性能立刻起飞。

正解SoA 只在“某个 kernel 只访问其中一个字段”时最有利。如果 kernel 同时访问多个字段，AoS 反而可能让每个线程一次性拿齐数据，减少指令和缓存行。真正要做的是根据热点访问模式选择布局，而不是无脑转换。

为何常见：网上教程常以单个字段的场景举例，读者容易把局部经验推广成万能结论。

面试题

Q1. 全局内存最符合下面哪个描述？

单选题 · easy

A.容量最大但延迟最高← 正确答案
B.容量最小但速度最快
C.只能被单个线程访问
D.由程序员手工管理缓存

答案：A。全局内存在 GPU 存储层次中容量最大、延迟最高，是所有线程都能访问的公共存储。它不像共享内存那样由程序员手工管理，所以 B、C、D 都不对。把它想成城外大仓库，什么都大，但路远。

Q2. 合并访问是指：

单选题 · easy

A.相邻线程访问相邻地址← 正确答案
B.所有线程访问同一地址
C.线程随机访问任意地址
D.每个线程访问独立缓存

答案：A。合并访问的核心是让 warp 内相邻线程访问连续地址，使硬件一个事务就能覆盖所有请求。广播访问是一种特殊高效模式，但不是合并访问的定义，C 和 D 更不对。

Q3. 一个 warp 的 32 个线程连续读取 32 个 float（共 128B），若地址对齐，最少需要几个 128B 内存事务？

单选题 · easy

A.只需要 1 个事务← 正确答案
B.需要 32 个事务
C.需要 128 个事务
D.需要 4 个事务

答案：A。32 个 float 正好 128 字节，连续且对齐时全部落在一个 128B 段内，一次内存事务就能完成。就像 32 个拣货员站在同一排，一辆托盘车全部搞定。

Q4. 跨步访问（stride > 1）性能差的主要原因是：

单选题 · mid

A.产生了多个内存事务← 正确答案
B.线程产生了分支
C.寄存器发生溢出
D.共享内存冲突增加

答案：A。跨步访问让地址散落在不同内存段，硬件必须发起多个事务，实际搬了很多用不上的数据。它不是分支问题，也不是寄存器和共享内存的问题。

Q5. 把结构体数组 AoS 改成数组结构体 SoA，最主要好处是：

单选题 · mid

A.同字段连续存放利于合并← 正确答案
B.整体数据量显著减小
C.省去线程同步开销
D.自动扩大显存容量

答案：A。SoA 把所有对象的同一字段放在连续内存里，线程按编号读同一字段时自然合并。数据总量不变，也不会减少同步，更不会扩大显存。

Q6. 下列哪种地址访问最可能造成非合并访问？

单选题 · mid

A.每隔 8 个 float 取一个← 正确答案
B.每个线程取相邻 float
C.所有线程读同一 float
D.地址按 128B 对齐

答案：A。每隔 8 个 float 取一个会让 32 个线程的地址拉得很散，跨过多个 128B 段。相邻和广播都是高效模式，地址对齐本身也不会导致非合并。

Q7. 有效带宽的计算方式是：

单选题 · mid

A.有效数据字节数 ÷ 耗时← 正确答案
B.理论带宽 ÷ 耗时
C.总指令数 × 时钟频率
D.线程数 ÷ 存储段数

答案：A。有效带宽衡量“真正有用数据”的搬运速度，用有效字节数除以耗时。理论带宽是硬件上限，不能简单除以耗时；线程数、指令数与带宽不是这种关系。

Q8. 一个 warp 读取连续 64 个 int（256B），不考虑缓存，最少会触发几个 128B 内存事务？

单选题 · hard

A.2 个← 正确答案
B.4 个
C.8 个
D.16 个

答案：A。连续 256B 恰好覆盖两个对齐的 128B 段，所以最少 2 个事务。这道题考的是把“请求字节数 ÷ 段大小”和“是否对齐”结合起来，不能只凭直觉选 8 或 16。

Q9. 以下哪些思路有助于提升全局内存合并访问效率？

多选题 · hard

A.让线程 ID 与地址增量对应← 正确答案
B.把所有字段塞进一个大结构体
C.保证基址按 128B 对齐← 正确答案
D.避免大跨步的访问模式← 正确答案

答案：ACD。线程 ID 对应地址增量是典型线性化访问，对齐基址让事务不跨段，避免大跨步是根本优化方向。把所有字段塞进大结构体往往会让单字段访问变得非合并，反而有害。

Q10. 关于全局内存，下列说法正确的有：

多选题 · extreme

A.所有线程都可以访问← 正确答案
B.延迟通常高于共享内存← 正确答案
C.容量通常小于寄存器
D.合并访问能减少事务数← 正确答案

答案：ABD。全局内存是全局可见的公共存储，延迟比共享内存高很多，容量也远大于寄存器。合并访问能让多个线程的请求落在更少的事务段中，从而减少事务数。