CUDA C编程权威指南-第四章：全局内存

系列导航：导读 | 上一篇：第3章 CUDA执行模型 | 下一篇：第5章 共享内存和常量内存

系列第 4 篇。第 3 章里有个现象：线程块最内层维度取线程束一半时，内存加载效率会明显变差——根因是对全局内存的访问模式不好。这一章就讲核函数和全局内存怎么打交道：内存模型、对齐与合并、带宽优化，以及固定内存、零拷贝、统一内存等用法。

---

前言：从执行到访存

书中开篇即指出：在上章「用 nvprof 检测内存选项」中，将线程块的最内层维度设为线程束大小的一半时，内存加载效率会大幅下降——这一性能损失无法通过线程束调度或暴露的并行量来解释。真正的原因是对全局内存的访问模式不佳。

本章将帮助你（书中原意）：

• 剖析内核与全局内存的交互，理解这些交互如何影响性能
• 解释 CUDA 内存模型，通过分析不同的全局内存访问模式，学会在核函数中高效使用全局内存

大多数 GPU 应用受内存带宽限制；若不能正确组织对全局内存的访问，其他优化往往事倍功半。本章将系统覆盖：内存模型、内存管理（含固定内存、零拷贝、统一寻址）、对齐与合并、读写路径、矩阵转置案例、带宽测量、AoS/SoA，为第 5 章共享内存与常量内存打下基础。

---

一、本章在全书中的位置与学习目标

1.1 为什么要学「全局内存」

第 3 章从执行模型回答了「线程如何被调度、为何 32 很重要、占用率与延迟隐藏」；但性能瓶颈往往不仅在计算，更在访存。需要搞清楚的是：将线程块最内层维度设为线程束大小的一半时，内存加载效率会大幅下降——这一现象无法仅用线程束调度或并行量解释，根本原因是对全局内存的访问模式不佳。因此，理解 CUDA 内存模型、掌握对齐与合并、学会选择合适的内存分配方式（固定内存、零拷贝、统一内存），是写出高性能核函数的基础；本章也为第 5 章「用共享内存改善全局内存访问」做铺垫。

1.2 学完本章，你应该能回答

学习目标	检验方式
理解 CUDA 内存层次（寄存器、本地、共享、全局、常量、纹理）及 L1/L2/只读缓存	能说出各内存的作用域、生命周期及何时用谁
掌握 cudaMalloc / cudaMemcpy / 固定内存 / 零拷贝 / 统一内存 的用法与适用场景	能解释为何 H2D/D2H 要尽量减少、何时用固定内存、零拷贝的局限
理解 对齐与合并的定义及对事务数、带宽利用率的影响	能判断一段访问是否对齐、是否合并，并估算事务数
理解 128 字节与 32 字节事务、L1/L2/只读缓存三种读取路径	能说出写操作为何不经 L1、何时用 __ldg 或 const __restrict__
掌握 矩阵转置中「合并读+非合并写」与「非合并读+合并写」的取舍	能解释为何写合并往往更关键，并知道第 5 章用共享内存可双端合并
会用 有效带宽公式与 nvprof 判断带宽受限、计算受限	能计算有效带宽并与理论峰值对比
理解 AoS vs SoA 对合并与带宽的影响	能解释为何只访问部分字段时 AoS 约 50% 带宽浪费

1.3 博客阅读导图（本章架构）

第 4 章  全局内存
│
├── 二、CUDA 内存模型概述（书 4.1 节）
│   ├── 内存层次结构的必要性、可编程内存详解、GPU 缓存
│   └── 重点：全局内存的「对齐+合并」决定带宽利用率
│
├── 三、内存管理（书 4.2 节）
│   ├── 设备内存分配与释放、主机与设备间数据传输
│   ├── 固定内存（cudaMallocHost）、零拷贝（cudaHostAlloc Mapped）
│   └── 统一虚拟寻址与统一内存（cudaMallocManaged）
│
├── 四、内存访问模式（书 4.3 节）★ 重难点
│   ├── 内存事务、对齐与合并的概念与四种模式表
│   ├── 全局内存读取路径（L1/L2、只读缓存、__ldg/__restrict__）
│   ├── 全局内存写入（仅 L2、32 字节段）
│   └── 矩阵转置：合并读+非合并写 vs 非合并读+合并写
│
├── 五、核函数能达到的带宽与性能分析（书 4.4 节）
│   └── 有效带宽公式、与理论峰值对比、计算受限 vs 带宽受限
│
├── 六、使用统一内存的矩阵加法（书 4.5 节）
│   └── cudaMallocManaged 与传统方式的对比
│
├── 七、结构体数组（AoS）与数组结构体（SoA）（书 4.6 节）
│   └── 只访问部分字段时 AoS 的带宽浪费与 SoA 的合并优势
│
└── 八、本章小结与重难点回顾

---

二、CUDA 内存模型概述（书 4.1 节）

2.1 内存层次结构的必要性

程序具有时间局部性（刚用过的数据很快会再用）和空间局部性（用过的数据附近的数据也会被用）。内存层次结构在「容量大、速度慢」与「容量小、速度快」之间折中，把热数据放在更快、更小的存储中。CUDA 与 CPU 的一大区别是：CUDA 将部分内存层次暴露给程序员（如共享内存、常量内存），便于我们显式控制数据放置与访问方式。

CUDA 编程模型从 GPU 架构中抽象出一个内存层次结构（书中图 2-3 为简化的 GPU 内存结构，主要包含全局内存和共享内存）。在 GPU 内存层次结构中，最主要的两种内存是全局内存和共享内存：全局内存类似于 CPU 的系统内存，共享内存类似于 CPU 的缓存，但 GPU 的共享内存可以由 CUDA C 的内核直接控制。

理解与体会：CPU 的缓存对程序员是透明的；而 CUDA 把共享内存、常量内存等暴露出来，意味着我们既要理解「谁快谁慢」，也要在代码里显式决定数据放哪、怎么访问。这是「以空间换时间」和「以控制换性能」的典型体现。

2.2 可编程内存详解

书中按作用域、生命周期和缓存行为对可编程内存进行了分类。下表与书中描述一致，便于查阅：

内存类型	修饰符/分配方式	作用域	生命周期	特点
寄存器	核函数内普通变量	线程私有	线程	最快；Fermi 每线程最多 63，Kepler 扩展至 255；溢出则进本地内存
本地内存	寄存器溢出、大数组、未知索引	线程私有	线程	实际在全局内存区，高延迟、低带宽
共享内存	__shared__	线程块	块	片上、可编程；与 L1 共享 64KB（Fermi），需注意 bank 冲突
全局内存	cudaMalloc / __device__	全部	应用	容量大、延迟高；访问需对齐与合并
常量内存	__constant__ + cudaMemcpyToSymbol	全部	应用	只读、有专用常量缓存；同址广播优
纹理内存	纹理 API	全部	应用	只读、滤波、二维空间局部性优

要点补充（书中对应内容）：

• 寄存器：核函数内未加特殊修饰的局部变量通常存放在寄存器中。寄存器是 SM 的稀缺资源；溢出部分会放到本地内存，对性能影响很大。
• 本地内存：编译器可能将以下变量放在本地内存——不满足寄存器限定条件的变量、占用大量寄存器空间的大数组或结构体、使用未知索引引用的本地数组。本地内存与全局内存在同一物理存储区域，访问特性为高延迟、低带宽。
• 共享内存：在核函数中用 __shared__ 声明；生命周期与线程块一致。Fermi 上可通过 cudaFuncSetCacheConfig 在共享内存与 L1 之间调整 64KB 的划分（如 prefer shared / prefer L1 / prefer equal）。
• 常量内存：用 __constant__ 在全局范围声明，通过 cudaMemcpyToSymbol 由主机初始化；对所有设备只读。当线程束中所有线程从相同地址取数据时，常量内存表现最好（一次读取广播给整个 warp）；不同线程取不同地址时不如合并的全局内存访问。
• 纹理内存：驻留在设备内存，通过专用只读缓存访问，支持滤波与二维局部性优化。
• 全局内存：对应设备 DRAM，可动态或静态分配（__device__）；访问时对齐与跨线程的地址分布共同决定事务数量与带宽利用率。这一节会说明：一般情况下，满足内存请求的事务越多，未使用的字节被传输的可能性越大，数据吞吐量就会降低。

2.3 GPU 缓存（不可编程）

与 CPU 缓存类似，GPU 缓存不可编程，其行为在出厂时已设定。书中提到的四种缓存为：

• 一级缓存：每个 SM 一块，与共享内存共享片上存储（如 Fermi 64KB 可配置）。用于缓存全局内存和本地内存的加载；存储操作不经过 L1。
• 二级缓存：所有 SM 共享，全局内存访问必经 L2（除非走只读缓存路径）。
• 只读常量缓存：每 SM 一块，用于常量内存。
• 只读纹理/只读缓存：用于纹理与（计算能力 3.5+）全局内存只读加载，粒度 32 字节。

Fermi 及之后架构允许通过编译选项控制全局内存加载是否经 L1（见第四节）。

---

三、内存管理（书 4.2 节）

CUDA 编程模型假设系统由一个主机和一个设备组成，且各自拥有独立的内存。核函数在设备上运行；为使你有充分的控制权并使系统达到最佳性能，CUDA 运行时负责分配与释放设备内存，并在主机内存和设备内存之间传输数据。书中表 2-1 列出了标准 C 函数与对应的 CUDA C 内存操作函数。下表与书中一致，便于本章查阅与对照：

标准 C 函数	CUDA C 函数	功能说明
malloc	cudaMalloc	分配内存（设备端线性内存）
memcpy	cudaMemcpy	拷贝数据（主机↔设备或设备↔设备）
memset	cudaMemset	将设备内存某区域设为指定值
free	cudaFree	释放设备内存

主机端分配仍可用 malloc/free；设备端必须使用上表中的 CUDA API。不能在主机代码中解引用设备指针（如 *d_A），否则会导致崩溃；CUDA 6.0 起的统一寻址/统一内存提供了单指针模型，详见 3.5 小节。

3.1 设备内存的分配与释放

cudaMalloc — 在设备上分配线性内存：

cudaError_t cudaMalloc(void **devPtr, size_t count);

devPtr 为指针的指针，用于返回所分配内存的设备指针；count 为字节数。与标准 C 的 malloc 类似，只是在 GPU 内存中分配。分配失败时返回 cudaErrorMemoryAllocation。使用前应检查返回值；书中推荐用统一的错误处理宏（如 CHECK(cudaMalloc(...))）封装。

cudaMemset — 设备内存初始化：

cudaError_t cudaMemset(void *devPtr, int value, size_t count);

cudaFree — 释放设备内存：

cudaError_t cudaFree(void *devPtr);

注意：devPtr 必须是由 cudaMalloc（或其它设备分配 API）得到的指针；重复释放或非法指针会报错。设备内存的分配和释放非常影响性能，应尽量重复利用已分配缓冲区。

3.2 主机与设备间的数据传输

cudaMemcpy — 在主机与设备之间拷贝数据：

cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind);

kind 常用取值：

• cudaMemcpyHostToDevice：主机 → 设备
• cudaMemcpyDeviceToHost：设备 → 主机
• cudaMemcpyDeviceToDevice：设备 → 设备

此函数以同步方式执行：在 cudaMemcpy 返回且传输完成之前，主机应用程序是阻塞的。除内核启动外，大多数 CUDA API 会返回 cudaError_t；可用 cudaGetErrorString 将错误码转为可读字符串（类似 C 的 strerror）。

带宽对比（书中以 Fermi 为例）：

• 设备内存（GDDR5）：理论峰值约 144 GB/s（随型号不同）。
• PCIe 总线：约 8 GB/s 量级。

因此减少主机与设备之间的拷贝次数与数据量是优化的重要原则；应尽量在设备上完成多步计算后再回传，并善用固定内存、流重叠（第 6 章）或统一内存。

理解与体会：PCIe 带宽远低于设备内存带宽，因此「能少拷就少拷、能重叠就重叠」是金科玉律。固定内存的意义不仅在于单次传输更快，更在于为第 6 章的异步传输与流重叠创造条件。

3.3 固定内存（Pinned Memory）

普通主机内存是可分页的：操作系统可能将物理页换出或移动。在向设备传输时，驱动可能先锁页，将源数据复制到一块临时的「锁页」缓冲区，再从该缓冲区 DMA 到设备——多一次拷贝且难以与计算重叠。

固定内存（页锁定内存）在分配时就保证不会被换页或移动，设备可通过 DMA 直接访问，从而：

• 提高 H2D/D2H 传输带宽
• 便于与异步拷贝和流结合，实现传输与计算重叠（第 6 章）

cudaError_t cudaMallocHost(void **ptr, size_t count);
cudaError_t cudaFreeHost(void *ptr);

这一节会说明：固定内存的分配与释放成本比可分页内存高，且会占用不可换页的物理页，不宜无节制使用。适合对传输性能敏感、且会多次传输的大块数据。下面给出书中风格的完整主函数示例：使用固定内存分配主机数组，拷贝到设备、执行核函数、再拷回主机，并与错误处理宏配合使用。核函数仍为简单的向量加法（与第 2 章一致）。

#define CHECK(call)                                                            \
    do {                                                                       \
        cudaError_t err = call;                                                 \
        if (err != cudaSuccess) {                                               \
            printf("Error %s:%d, code %d, reason: %s\n", __FILE__, __LINE__,   \
                   err, cudaGetErrorString(err));                               \
            exit(1);                                                            \
        }                                                                      \
    } while (0)

__global__ void sumArraysGPU(float *a, float *b, float *c) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

int main() {
    int nElem = 1 << 14;  // 16384
    size_t nBytes = nElem * sizeof(float);

    // 固定内存分配（主机端）
    float *h_a, *h_b, *h_c;
    CHECK(cudaMallocHost((void**)&h_a, nBytes));
    CHECK(cudaMallocHost((void**)&h_b, nBytes));
    CHECK(cudaMallocHost((void**)&h_c, nBytes));
    // 初始化 h_a, h_b ...

    float *d_a, *d_b, *d_c;
    CHECK(cudaMalloc((void**)&d_a, nBytes));
    CHECK(cudaMalloc((void**)&d_b, nBytes));
    CHECK(cudaMalloc((void**)&d_c, nBytes));

    CHECK(cudaMemcpy(d_a, h_a, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, nBytes, cudaMemcpyHostToDevice));

    dim3 block(1024);
    dim3 grid(nElem / block.x);
    sumArraysGPU<<<grid, block>>>(d_a, d_b, d_c);

    CHECK(cudaMemcpy(h_c, d_c, nBytes, cudaMemcpyDeviceToHost));

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    cudaFreeHost(h_a);
    cudaFreeHost(h_b);
    cudaFreeHost(h_c);
    return 0;
}

使用 nvprof 对比「可分页内存 + cudaMemcpy」与「固定内存 + cudaMemcpy」时，可看到固定内存路径的 H2D/D2H 耗时更短；当数据规模较大时，传输时间的节省往往能抵消固定内存分配的成本。

3.4 零拷贝内存（Zero-Copy / Pinned Mapped Memory）

零拷贝内存是固定在主机上的内存，且设备端可通过另一套指针直接访问，无需显式 cudaMemcpy。书中适用场景包括：

• 提高 PCIe 传输率（某些访问模式下）
• 避免主机和设备之间的显式内存传输
• 设备内存不足时利用主机内存

创建零拷贝内存需使用 cudaHostAlloc，并传入 cudaHostAllocMapped 标志：

cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);

常用 flags：cudaHostAllocMapped（产生零拷贝）、cudaHostAllocPortable（所有 CUDA 上下文可用）、cudaHostAllocWriteCombined（某些设备上写传输更优）。设备端不能直接使用 pHost，必须先取得「设备可见」的指针：

cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);

核函数中使用 pDevice 访问的是主机上 pHost 指向的那块内存。注意：使用零拷贝时，主机与设备可能同时访问同一块内存，需注意主机与设备之间的竞争条件。零拷贝内存在物理上仍在主机，每次设备访问都经 PCIe；书中实验表明，对向量加法这类需要多次读写的内核，常规「先 cudaMemcpy 到设备再计算」往往比零拷贝更快，因为零拷贝下每次加载/存储都走 PCIe。零拷贝更适合「一次读、一次写」或设备显存不足时的权宜之计；在集成架构（如 Tegra，CPU 与 GPU 共享物理内存）上，零拷贝效果会好很多。

书中给出了零拷贝与常规设备内存的运行时间对比表（不同数据规模 (2^n) 元素）。下表与书中实验对应（单位：μs，为总耗时示例）：

数据规模 (n)（(2^n) 元素）	常规设备内存（μs）	零拷贝内存（μs）
10	2.5	3.0
12	3.0	4.1
14	7.8	8.6
16	23.1	25.8
18	86.5	98.2
20	290.9	310.5

零拷贝流程的典型主机端代码（书中风格）如下。注意：先用 cudaHostAlloc(..., cudaHostAllocMapped) 分配，再通过 cudaHostGetDevicePointer 取得设备端指针传给核函数；核函数执行时每次对 a_dev/b_dev 的访问都经 PCIe，因此数据量大时总时间更长。

float *a_host, *b_host, *res_d;
CHECK(cudaHostAlloc((void**)&a_host, nBytes, cudaHostAllocMapped));
CHECK(cudaHostAlloc((void**)&b_host, nBytes, cudaHostAllocMapped));
CHECK(cudaMalloc((void**)&res_d, nBytes));
// 初始化 a_host, b_host ...

float *a_dev, *b_dev;
CHECK(cudaHostGetDevicePointer((void**)&a_dev, a_host, 0));
CHECK(cudaHostGetDevicePointer((void**)&b_dev, b_host, 0));
sumArraysGPU<<<grid, block>>>(a_dev, b_dev, res_d);
CHECK(cudaMemcpy(res_from_gpu_h, res_d, nBytes, cudaMemcpyDeviceToHost));
// ...
cudaFreeHost(a_host);
cudaFreeHost(b_host);
cudaFree(res_d);

随规模增大，零拷贝的劣势更明显，因为内核执行时每次访存都经 PCIe。

重难点：零拷贝的物理位置始终在主机，设备每次访问都经 PCIe；因此对「多次读写」的内核（如向量加），常规「先 cudaMemcpy 再计算」通常更快。零拷贝更适合「一次读、一次写」或设备显存不足时的权宜之计；在集成 GPU（CPU 与 GPU 共享物理内存）上效果会好很多。

3.5 统一虚拟寻址（UVA）与统一内存（Unified Memory）

统一虚拟寻址（UVA）：自 CUDA 4.0（设备架构 2.0）起，设备内存和主机内存在同一虚拟地址空间中映射。这样，由 cudaHostAlloc 分配的固定主机内存在主机和设备上可使用同一地址，无需再通过 cudaHostGetDevicePointer 获取设备指针，可直接将主机返回的指针传给核函数（在支持 UVA 的平台上）。这简化了零拷贝的使用。

统一内存（Unified Memory，CUDA 6+）：书中提到，从 CUDA 6.0 开始，NVIDIA 提出了名为「统一寻址」的编程模型改进，连接了主机内存和设备内存空间，可使用单个指针访问 CPU 和 GPU 内存，无须彼此之间手动拷贝数据。更详细的实现方式包括 cudaMallocManaged 或 managed 变量：分配得到的内存由驱动和运行时统一管理，主机与设备均可通过同一指针访问，访问时若数据不在当前处理器则触发迁移。统一内存在语义上简化编程（无需显式 cudaMemcpy），适合快速原型与教学；与零拷贝的区别在于：零拷贝的物理位置始终在主机，而统一内存的物理位置可由系统迁移，以优化访问延迟和一致性。统一内存的详细行为与限制（如计算能力 6.x 的按需页面迁移）见 NVIDIA 文档；本章重点是建立「存在这样一种简化模型」的认识。书中在 4.5 节用矩阵加法对比了传统 cudaMalloc+cudaMemcpy 与使用统一内存的写法。

---

四、内存访问模式（书 4.3 节）★ 本章重难点

大多数设备端数据访问都是从全局内存开始的，且多数 GPU 应用程序容易受内存带宽的限制。因此，最大限度地利用全局内存带宽是调控核函数性能的基本。若不能正确调控全局内存的使用，其他优化方案很可能也收效甚微。

为了在读写数据时达到最佳性能，内存访问操作必须满足一定条件。CUDA 执行模型的一个显著特征是：指令必须以线程束为单位发布和执行，存储操作也是同样。在执行内存指令时，线程束中的每个线程都提供一个正在加载或存储的内存地址；根据线程束中内存地址的分布，内存访问可被分成不同的模式。

4.1 内存事务与对齐、合并的概念

内存事务：从核函数发起一次加载或存储请求，到硬件完成该请求所涉及的总线/缓存操作。事务有固定的粒度。这一节会说明：核函数的内存请求通常是在 DRAM 设备和片上内存之间以 128 字节或 32 字节内存事务来实现的。所有对全局内存的访问都会通过二级缓存；许多访问还会通过一级缓存，取决于访问类型和 GPU 架构。若两级缓存都用到，则内存访问由** 128 字节**的设备内存事务实现；若只使用二级缓存，则由 32 字节的事务实现。

在优化应用程序时，需要注意设备内存访问的两个特性：

1. 对齐内存访问（Aligned）
2. 合并内存访问（Coalesced）

对齐：当一个内存事务的首个访问地址是缓存粒度（32 或 128 字节）的整数倍时，称为对齐内存访问；否则为非对齐。判定方式可归纳为：

• 若使用 L1 缓存（128 字节段）：首地址为 128 的整数倍（如 0、128、256）时为对齐。
• 若仅用 L2（32 字节段）：首地址为 32 的整数倍（如 0、32、64、128）时为对齐。

即：首地址 (\bmod 32 = 0)（或 (\bmod 128 = 0)） 为对齐条件。非对齐访问会带来带宽浪费（多取整段却只用一部分，或需多段才能覆盖请求）。

合并：当线程束内的线程访问的内存都在一个内存块里时，就会出现合并访问。对齐且合并是理想状态：线程束内所有线程访问的数据在一个内存块内，且从内存块的首地址开始被需要，此时用最少的事务次数满足最多的内存请求，效率最高。

带宽利用率可理解为：有用字节数 / 实际传输字节数。设线程束请求 32 个 float（128 字节），若 1 个 128 字节事务即可满足且首地址对齐，则利用率 100%；若需 2 个 128 字节事务（例如连续但非对齐），则利用率约 50%。书中用图示说明：若一个事务加载的数据分布在多个对齐段上（例如部分在 0～127，部分在 128～255），则需要多次事务，利用率下降；最坏情况下每个线程访问不同段，需要多达 32 次事务，利用率可低至 (128/(32\times 128)=1/32) 或更差（按 128 字节段、每线程 4 字节计）。

理解与体会：第 3 章「块最内层维度为 16 时效率下降」的本质在这里揭晓：16 个线程只占半个 warp，若按连续索引访问，半 warp 只需 64 字节，可能仍落在 1 个 128 字节段内，但若与相邻半 warp 合在一起看，整体访问模式可能跨段或未对齐，导致事务数增加、利用率下降。所以让块最内层维度为 32 的倍数，本质是让每个 warp 的访问尽量落在同一对齐段内，实现合并。

4.2 对齐与合并的直观例子与四种模式表

书中以一个线程束加载数据为例（使用一级缓存，即 128 字节段）说明不同访问模式对事务数和利用率的影响。下表将四种典型情况归纳（书中多处图示对应）：

访问模式	描述（32 线程 × 4 字节）	事务数（128B 段）	总线利用率（约）
对齐 + 合并	连续 128 字节，首地址对齐 128	1	100%
对齐但不合并	32 个线程访问同一 4 字节	1	4/128≈3.125%
合并但不对齐	连续 128 字节，首地址错开（如 1）	2	约 50%
不对齐且不合并	32 个线程落在 N 个不同 128B 段	N（1～32）	约 128/(N×128)=1/N

• 理想情况：32 个线程各请求 4 字节，共 128 字节，且这 128 字节落在一个 128 字节的对齐段内 → 1 个事务，总线利用率 100%。
• 连续但非对齐：32 个线程请求连续 128 字节，但首地址不对齐，例如横跨 0～127 和 128～255 两段 → 需要 2 个 128 字节事务，利用率约 50%。
• 分散：若 32 个线程访问的地址分布在 N 个不同的 128 字节段内（1≤N≤32），则需要 N 个事务，利用率约为 1/N。
• 同一地址：若线程束内所有线程请求同一 4 字节，仍只需 1 个 128 字节段，但有用数据仅 4 字节，利用率 4/128≈3.125%。

结论：用最少的事务次数满足最多的内存请求；事务数量和吞吐量需求随设备的计算能力变化。

4.3 全局内存读取路径：L1/L2 与只读缓存

这一节会说明：全局内存加载操作是否会通过一级缓存取决于**(1) 编译器选项和(2) 设备的计算能力**（较老设备可能没有 L1）。在 Fermi 和 Kepler K40 及以后的 GPU（计算能力 3.5+）中，可通过编译器标志启用或禁用全局内存加载的一级缓存。

• 禁用 L1（仅用 L2）：
  -Xptxas -dlcm=cg
  所有对全局内存的加载请求直接进入 L2；若 L2 缺失则由 DRAM 完成。此时事务粒度为 32 字节，对非对齐或非合并的分散访问有时能获得更好的总线利用率（因为每次只取 32 字节，无效字节更少）。

• 启用 L1：
  -Xptxas -dlcm=ca
  全局内存加载请求先经 L1；L1 缺失则走 L2；再缺失则由 DRAM 完成。此种模式下，一个内存加载请求由一个 128 字节的设备内存事务实现。

书中特别说明：在 Kepler K10、K20、K20X 上，一级缓存不用于缓存全局内存加载，而是专门用于缓存寄存器溢出到本地内存的数据。

只读缓存（计算能力 3.5+）：最初预留给纹理内存；对计算能力 3.5 及以上的设备，只读缓存也支持使用全局内存加载代替 L1。只读缓存的加载粒度是 32 字节，对分散读取通常优于 L1。书中给出了两种使用方式，下面用完整核函数示例（书中风格）：

方式一：使用 __ldg 内建函数，强制通过只读缓存加载：

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

方式二：使用 const __restrict__ 修饰只读指针，让 nvcc 自动选择只读缓存（无别名、只读）：

__global__ void copyKernelRestrict(int * __restrict__ out, const int* __restrict__ in) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    out[idx] = in[idx];
}

这样编译器会识别 in 为只读且无别名，可自动通过只读缓存加载，无需改写成 __ldg。

4.4 全局内存写入

这一节会说明：存储操作在 Fermi 或 Kepler 上不能使用一级缓存；在发送到设备内存之前，存储操作只通过二级缓存。存储操作以 32 字节段的粒度执行；内存事务可被分为 1 段、2 段或 4 段。例如，若两个地址同属于一个 128 字节区域但不属于同一对齐的 64 字节区域，可能执行一个 4 段事务（比两次 1 段事务更高效）。对齐的、连续 128 字节写入可用一个 4 段事务完成；分散在 192 字节范围内则可能需 3 个 1 段事务。因此写端也应力求对齐与合并；在不少内核中，合并写对性能更关键，因为写不经过 L1。

重难点：读可以走 L1（128B 事务）或 L2/只读缓存（32B 事务），但写只走 L2、以 32 字节段为单位。因此「写合并」在带宽敏感型内核里往往比「读合并」更关键，矩阵转置中优先保证写合并正是基于这一点。

4.5 矩阵转置中的读/写模式（书中案例与完整核函数）

矩阵在全局内存中通常按行优先线性存储：元素 (row, col) 的线性下标为 row * nx + col。矩阵转置中，每个线程读一个元素、写一个元素，读和写的下标互换，因此必然出现「读合并则写不合并」或「读不合并则写合并」的取舍。

方案一：读按行（合并读）、写按列（非合并写）
线程 (ix, iy) 读 A[iy*nx+ix]（同一行连续），写 B[ix*ny+iy]（同一列，在行优先下不连续）。书中风格的核函数如下（nx×ny 矩阵，A 为输入，B 为输出）：

__global__ void copyRowReadColWrite(float *A, float *B, int nx, int ny) {
    unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
    if (ix < nx && iy < ny) {
        B[ix * ny + iy] = A[iy * nx + ix];  // 读 A 按行→合并；写 B 按列→非合并
    }
}

方案二：读按列（非合并读）、写按行（合并写）
线程 (ix, iy) 读 A[iy*nx+ix] 可理解为「按 A 的行读」即合并；若改为按 B 的列对应到 A，即读 A[ix*ny+iy]（按 A 的列），则为非合并读，写 B[iy*nx+ix] 为按行写即合并。下面写法为「读 A 的 (iy,ix)、写 B 的 (ix,iy)」，读的是 A 的列，写的是 B 的行：

__global__ void copyColReadRowWrite(float *A, float *B, int nx, int ny) {
    unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
    if (ix < nx && iy < ny) {
        B[iy * nx + ix] = A[ix * ny + iy];  // 读 A 按列→非合并；写 B 按行→合并
    }
}

书中结论：合并写往往更影响性能（写路径无 L1）；因此常优先保证写合并（即采用方案二），接受读端非合并；若使用只读缓存或 __ldg 读 A，可缓解非合并读的损失。第 5 章将用共享内存做中转：先合并读入块对应的 tile，在共享内存中转置，再合并写出，从而读写两端都接近合并。

重难点：矩阵转置是典型的「读合并则写不合并、写合并则读不合并」场景。写端更应保证合并，因为存储不经过 L1，写路径对带宽更敏感。实际优化时优先选「合并写」方案，读端用只读缓存或共享内存（第 5 章）弥补。

4.6 带宽利用率小结

• 对齐 + 合并：理想，单事务（或最少事务）、高利用率。
• 对齐但不合并：例如 32 个线程访问同一 4 字节，利用率极低。
• 合并但不对齐：连续 128 字节跨两段，需 2 笔 128 字节事务，利用率约 50%。
• 既不对齐也不合并：可能需多达 32 笔事务，利用率最低。

优化目标：最少事务数 + 每笔事务中请求的字节尽量多。

---

五、核函数能达到的带宽与性能分析（书 4.4 节）

本节对应书中「核函数能达到多少带宽」及与理论峰值的对比。书中通过 nvprof 等工具测量内核的实际带宽（如 GB/s），并与理论峰值对比，以判断内核是计算受限还是带宽受限，以及离理论上限还有多少空间。这是判断优化方向（改访存还是改计算、改配置）的重要依据。

实测带宽可按下式估算（书中思路）：

[
\text{有效带宽（GB/s）} = \frac{(\text{读字节数} + \text{写字节数}) \times 10^{-9}}{\text{核函数执行时间（秒）}}
]

例如，若某内核每个线程读 4 字节、写 4 字节，共 (N) 个线程，则总字节数 (= N \times 8)；用 nvprof 得到核函数时间 (T)（秒），则有效带宽 (\approx N \times 8 / (T \times 10^9)) GB/s。将该值与设备理论峰值带宽（如 Fermi 144 GB/s）比较，可得到带宽利用率，从而判断是否还有优化空间。

第 2 章已介绍「指令:字节」比（如 Tesla K10 约 13.6:1）：若每字节访问对应的指令数高于该比值，偏计算受限；否则偏带宽受限。大多数 HPC 工作负载是带宽受限的。

提高带宽利用的两类思路（与第 3 章一致）：

1. 足够的并发内存操作，以隐藏内存延迟（通过执行配置和每线程独立访存）。
2. 对齐及合并内存访问，减少带宽浪费。

这一节会说明：可通过 (1) 对核函数启动的执行配置进行试验 以充分体现每个 SM 的并行性，(2) 增加每个线程中执行独立内存操作的数量（如展开技术）来优化。将实测带宽与理论峰值比较，可判断瓶颈在算法（计算）还是带宽（访存）；再结合 gld_throughput、gst_throughput、gld_efficiency、gst_efficiency 等 nvprof 指标做针对性改进。

---

六、使用统一内存的矩阵加法（书 4.5 节）

书中用矩阵加法对比了传统方式与统一内存的写法。传统方式：主机分配 h_A、h_B、h_C，设备分配 d_A、d_B、d_C，执行 cudaMemcpy H2D → 核函数 → cudaMemcpy D2H。使用统一内存时，可用 cudaMallocManaged 分配一块既可被主机访问又可被设备访问的内存，主机端初始化后直接将该指针传给核函数，核函数执行完毕后主机可直接读取结果，无需显式 cudaMemcpy。这样代码更简洁，适合快速原型与教学；首次访问时由系统迁移数据，可能有一次性的迁移开销。追求极限性能时，仍可能需显式分配与流重叠（第 6 章）。

---

七、结构体数组（AoS）与数组结构体（SoA）（书 4.6 节）

书中以存储成对的浮点数据 (x, y) 为例，对比两种布局对合并与带宽的影响。

AoS（Array of Structures）：每个逻辑点是一个结构体，x、y 在内存中相邻存放。

struct InnerStruct {
    float x;
    float y;
};
struct InnerStruct myAoS[N];   // 内存布局: x0,y0, x1,y1, x2,y2, ...

若核函数只使用 x（例如对 x 做平方），典型写法如下。此时 32 个线程访问 myAoS[i].x，在内存上间隔 8 字节（每个结构体 8 字节），不是连续 4 字节，无法合并成少量事务；且每次加载会带出不需要的 y。

__global__ void useX_AoS(struct InnerStruct *data, float *out) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    out[i] = data[i].x * data[i].x;  // 只用到 x，但会加载 x 和 y
}

这一节会说明：在 GPU 上以 AoS 存储并执行只使用 x 的应用程序，将导致约 50% 的带宽损失（因为每 32 字节或 128 字节段中有一半是未使用的 y），且会在不需要的 y 上浪费二级缓存空间。

SoA（Structure of Arrays）：同一字段的所有值放在一个数组里，逻辑上同一「点」的 x、y 分布在不同数组。

struct InnerArray {
    float x[N];
    float y[N];
};
struct InnerArray mySoA;   // x[0..N-1] 连续，y[0..N-1] 连续

若 32 个线程访问 mySoA.x[i]，地址连续，易合并，且不加载 y：

__global__ void useX_SoA(float *x, float *out) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    out[i] = x[i] * x[i];  // 连续访问 x[]，合并
}

书中结论：并行编程范式（尤其是 SIMD/SIMT）对 SoA 更友好；CUDA 中普遍倾向于 SoA，因为这种内存访问可以有效地合并。

理解与体会：AoS 更符合「面向对象」的直觉（一个点就是一个 struct），但在 GPU 上若只访问部分字段，会白拉一整条 cache line、浪费带宽。SoA 是「数据导向」的布局，更贴合 SIMT 的合并访问需求，是性能敏感代码中的常见选择。

---

八、本章小结与重难点回顾

8.1 知识小结（与书中对应）

• CUDA 内存模型：寄存器、本地内存、共享内存、全局内存、常量内存、纹理内存各有作用域与生命周期；L1/L2 与只读缓存不可编程但影响访问粒度和路径；全局内存与共享内存是层次中的两个核心，本章重点在全局内存。
• 内存管理：cudaMalloc/cudaFree、cudaMemcpy 是基础；固定内存（cudaMallocHost）提升 H2D/D2H 带宽并利于与计算重叠；零拷贝（cudaHostAlloc Mapped + cudaHostGetDevicePointer）省显式拷贝但经 PCIe 访问，适合特定场景；统一虚拟寻址与统一内存（cudaMallocManaged）简化编程。
• 访问模式：对齐（首地址为 32/128 字节整数倍）和合并（线程束内访问落在尽可能少的内存段内）是提升全局内存带宽利用的关键；事务越少、每笔有用字节越多，有效带宽越高。
• 读取路径：L1+L2（128B 事务）、仅 L2（32B 事务）、只读缓存（__ldg/const __restrict__，32B）；写仅经 L2，以 32 字节段为单位，写合并常更关键。
• 矩阵转置：典型「合并读+非合并写」或「非合并读+合并写」；第 5 章用共享内存可实现读写双端合并。
• AoS vs SoA：并行内核中 SoA 更易实现合并访问，只访问部分字段时 AoS 约 50% 带宽浪费。

8.2 重难点速查

重难点	要点
对齐与合并	对齐看首地址是否为 32/128 的整数倍；合并看线程束内地址是否落在尽量少的内存段内；二者共同决定事务数与带宽利用率。
128B vs 32B 事务	经 L1 的加载以 128B 为单位；仅 L2 或只读缓存以 32B 为单位；写只经 L2，32B 段。
为何写合并更关键	存储不经过 L1，写路径对带宽更敏感；矩阵转置等场景优先保证写合并，读端可用只读缓存或共享内存弥补。
固定 vs 零拷贝 vs 统一内存	固定：主机锁页，DMA 快、利于重叠；零拷贝：设备直接访问主机内存，每次经 PCIe；统一内存：单指针、自动迁移，编程简单，极限性能仍可能用显式拷贝+流。
AoS 的 50% 浪费	只访问一个字段时，相邻元素在 AoS 中间隔一个字段，无法连续合并，且会加载无用字段。SoA 下同字段连续，易合并。

8.3 学习思考

• 第 3 章与本章的衔接：第 3 章「块最内层维度为 16 时效率下降」的根因，是全局内存访问未按 warp 对齐、合并；本章给出了具体判定方法（对齐、合并、事务数）和优化方向（对齐+合并、只读缓存、写优先）。
• 与第 5 章的关系：当算法天然导致「读合并写不合并」或反过来时，单靠全局内存只能二选一；第 5 章用共享内存做中转，先合并读入 tile、在块内重组、再合并写出，可实现读写双端高效，是矩阵转置、归约等模式的标准做法。
• 实践建议：写新内核时先保证索引连续、首地址对齐（如 cudaMalloc 自然对齐）；用 nvprof 看 gld_efficiency/gst_efficiency 判断是否合并；只读数据可尝试 __ldg 或 const __restrict__；结构体数据优先考虑 SoA 或混合布局。

---

下一章预告

在下一篇博客中，我们将进入第 5 章：共享内存和常量内存：

• 共享内存的布局、bank 与 bank 冲突，以及用 padding 消除冲突
• 用共享内存做矩阵转置的中转，实现读写双端合并
• 并行归约中共享内存的使用与 warp 内优化
• 常量内存与只读缓存的使用场景
• 线程束洗牌指令（warp shuffle）

从「全局内存怎么访问」到「把数据搬到片上、按需复用」，是突破带宽瓶颈的下一步。

---

本章自测

1. 什么是「合并访问」？为何它对全局内存带宽至关重要？
2. 矩阵转置时，若只使用全局内存，为何无法同时做到读合并与写合并？
3. 有效带宽公式是什么？如何用它判断内核是带宽受限还是计算受限？

答案与解析

1. 合并访问指同一 warp 内线程访问的全局内存地址落在尽量少的连续段（如 128 字节）内，可合并成少量事务，提高带宽利用率；否则会触发多次事务，带宽浪费。
2. 按行读则地址连续（读合并），但按行写转置结果时写的是不连续列（写不合并）；反之按列读则读不合并、按列写则写合并。单用全局内存只能二选一。
3. 有效带宽 = (读字节数 + 写字节数) / 核函数执行时间。与理论峰值对比：若接近理论带宽则带宽受限，远低于则可能计算受限或存在非合并访问等问题。

---

系列导航：导读 | 上一篇：第3章 CUDA执行模型 | 下一篇：第5章 共享内存和常量内存

---

本文为「CUDA C编程权威指南」系列博客第 4 篇，共 10 章。基于《Professional CUDA C Programming》by John Cheng, Max Grossman, Ty McKercher。