（四）全局内存

本篇笔记参考如下：

https://blog.csdn.net/weixin_33298352/article/details/156397919?ops_request_misc=&request_id=&biz_id=102&utm_term=统一内存&utm_medium=distribute.pc_search_result.none-task-blog-2~all~sobaiduweb~default-3-156397919.142

https://face2ai.com/CUDA-F-4-5-使用统一内存的向量加法/

在本章，我们将剖析核函数与全局内存的联系及其对性能的影响。通过分析不同的全局内存访问模式来实现通过核函数高效地利用全局内存。

4.1 CUDA内存模型概述

内存的访问和管理是所有编程语言的重要部分。因为多数工作负载被加载和存储数据的速度所限制，所以有大量低延迟、高带宽的内存对性能是十分有利的。

4.1.1 内存层次结构的优点

应用程序不会在某一时间点访问任意数据或运行任意代码。应用程序往往遵循局部性原则，这表明它们可以在任意时间点访问相对较小的局部地址空间。有两种不同类型的局部性：

时间局部性
空间局部性

时间局部性认为如果一个数据位置被引用，那么该数据在较短的时间周期内很可能会再次被引用，随着时间流逝，该数据被引用的可能性逐渐降低。空间局部性认为如果一个内存位置被引用，则附近的位置也可能会被引用。

一个内存层次结构由具有不同延迟、带宽和容量的多级内存组成。通常，随着从处理器到内存延迟的增加，内存的容量也在增加。一个典型的层次结构如图

CPU和GPU的主存都采用的是DRAM（动态随机存取存储器），而低延迟内存（如CPU一级缓存）使用的则是SRAM（静态随机存取存储器）。

特性	DRAM (显存/内存)	SRAM (缓存/寄存器)
速度 (延迟)	较慢 (~50-100 纳秒)	极快 (~1-10 纳秒)
存储密度	极高（结构简单，能塞更多数据）	较低（晶体管多，占地方）
造价	便宜（性价比高）	极其昂贵
功耗	较低（但由于刷新会有底噪功耗）	较高（工作时快，待机时漏电流大）

在这种内存层次结构中，当数据被处理器频繁使用时，该数据保存在低延迟、低容量的存储器中；而当该数据被存储起来以备后用时，数据就存储在高延迟、大容量的存储器中。这种内存层次结构符合大内存低延迟的设想。

4.1.2 CUDA内存模型

对编程者来说，一般有两种类型的存储器：

·可编程的：需要显式地控制哪些数据存放在可编程内存中

·不可编程的：你不能决定数据的存放位置，程序将自动生成存放位置以获得良好的性能

在CPU内存层次结构中，一级缓存和二级缓存都是不可编程的存储器。另一方面，CUDA内存模型提出了多种可编程内存的类型：

寄存器
共享内存
本地内存
常量内存
纹理内存
全局内存

一个核函数中的线程都有自己私有的本地内存。一个线程块有自己的共享内存，对同一线程块中所有线程都可见，其内容持续线程块的整个生命周期。所有线程都可以访问全局内存。所有线程都能访问的只读内存空间有：常量内存空间和纹理内存空间。全局内存、常量内存和纹理内存空间有不同的用途。纹理内存为各种数据布局提供了不同的寻址模式和滤波模式。对于一个应用程序来说，全局内存、常量内存和纹理内存中的内容具有相同的生命周期。

4.1.2.1 寄存器

寄存器是GPU上运行速度最快的内存空间。核函数中声明的一个没有其他修饰符的自变量，通常存储在寄存器中。在核函数声明的数组中，如果用于引用该数组的索引是常量且能在编译时确定，那么该数组也存储在寄存器中。注意：如果一个核函数使用了超过硬件限制数量的寄存器，则会用本地内存替代多占用的寄存器，这种寄存器溢出会给性能带来不利影响。

4.1.2.2 本地内存

核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中。编译器可能存放到本地内存中的变

量有：

在编译时使用未知索引引用的本地数组
可能会占用大量寄存器空间的较大本地结构体或数组
任何不满足核函数寄存器限定条件的变量

“本地内存”这一名词是有歧义的：溢出到本地内存中的变量本质上与全局内存在同一块存储区域。

本地内存访问符合高效内存访问要求。对于计算能力2.0及以上的GPU来说，本地内存数据也是存储在每个SM的一级缓存和每个设备的二级缓存中。

4.1.2.3 共享内存

在核函数中使用如下修饰符修饰的变量存放在共享内存中：

1	__shared__

因为共享内存是片上内存，所以与本地内存或全局内存相比，它具有更高的带宽和更低的延迟。它的使用类似于CPU一级缓存，但它是可编程的。

共享内存在核函数的范围内声明，其生命周期伴随着整个线程块。当一个线程块执行结束后，其分配的共享内存将被释放并重新分配给其他线程块。

一个块内的线程通过使用共享内存中的数据可以相互合作。访问共享内存必须同步使用如下调用

1	void __syncthreads();

该函数设立了一个执行障碍点，即同一个线程块中的所有线程必须在其他线程被允许执行前达到该处。为线程块里所有线程设立障碍点，这样可以避免潜在的数据冲突。

4.1.2.4 常量内存

常量内存驻留在设备内存中，并在每个SM专用的常量缓存中缓存。常量变量用如下修饰符来修饰：

1	__constant__

常量变量必须在全局空间内和所有核函数之外进行声明。对于所有计算能力的设备，都只可以声明64KB的常量内存。常量内存是静态声明的，并对同一编译单元中的所有核函数可见。

4.1.2.5 纹理内存

纹理内存驻留在设备内存中，并在每个SM的只读缓存中缓存。纹理内存是一种通过指定的只读缓存访问的全局内存。只读缓存包括硬件滤波的支持，它可以将浮点插入作为读过程的一部分来执行。纹理内存是对二维空间局部性的优化，所以线程束里使用纹理内存访问二维数据的线程可以达到最优性能。

4.1.2.6 全局内存

全局内存是GPU中最大、延迟最高并且最常使用的内存。global指的是其作用域和生命周期。它的声明可以在任何SM设备上被访问到，并且贯穿应用程序的整个生命周期。一个全局内存变量可以被静态声明或动态声明。可以使用如下修饰符在设备代码中静态地声明一个变量：

1	__device__

全局内存常驻于设备内存中，可通过32字节、64字节或128字节的内存事务进行访问。这些内存事务必须自然对齐，也就是说，首地址必须是32字节、64字节或128字节的倍数。优化内存事务对于获得最优性能来说是至关重要的。

4.1.2.7 GPU缓存

跟CPU缓存一样，GPU缓存是不可编程的内存。在GPU上有4种缓存：

一级缓存
二级缓存
只读常量缓存
只读纹理缓存

每个SM都有一个一级缓存，所有的SM共享一个二级缓存。一级和二级缓存都被用来在存储本地内存和全局内存中的数据，也包括寄存器溢出的部分。

4.1.2.8 CUDA 内存模型核心对比

内存类型	物理位置	访问权限（作用域）	访问速度	典型用途
1. 寄存器	片上	单线程 (Thread)	极快 (0 cycle)	局部变量、循环计数器
2. 共享内存	片上	线程块 (Block)	极快 (1-10 cycles)	块内线程通信、缓存重用数据
3. 本地内存	片外	单线程 (Thread)	慢 (DRAM 速度)	寄存器溢出、大数组
4. 常量内存	片外 (带缓存)	全局 (Grid)	快 (命中缓存时)	只读常数、内核参数
5. 纹理内存	片外 (带缓存)	全局 (Grid)	快 (具有空间局部性)	图像处理、空间查找表
6. 全局内存	片外	全局 (Grid)	慢 (200-800 cycles)	初始输入数据、最终结果

4.1.2.9 静态全局内存

我们通过代码globalVariable.cu说明如何静态声明一个全局变量。

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * An example of using a statically declared global variable (devData) to store
 * a floating-point value on the device.
 */

__device__ float devData;

__global__ void checkGlobalVariable()
{
    // display the original value
    printf("Device: the value of the global variable is %f\n", devData);

    // alter the value
    devData += 2.0f;
}

int main(void)
{
    // initialize the global variable
    float value = 3.14f;
    CHECK(cudaMemcpyToSymbol(devData, &value, sizeof(float)));
    printf("Host:   copied %f to the global variable\n", value);

    // invoke the kernel
    checkGlobalVariable<<<1, 1>>>();

    // copy the global variable back to the host
    CHECK(cudaMemcpyFromSymbol(&value, devData, sizeof(float)));
    printf("Host:   the value changed by the kernel to %f\n", value);

    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

编译并执行程序

1 2	nvcc -arch=sm_120 globalVariable.cu -o globalVariable ./globalVariable

得到结果如下：

1
2
3

Host:   copied 3.140000 to the global variable
Device: the value of the global variable is 3.140000
Host:   the value changed by the kernel to 5.140000

尽管主机和设备的代码存储在同一个文件中，它们的执行却是完全不同的。即使在同一文件内可见，主机代码也不能直接访问设备变量。类似地，设备代码也不能直接访问主机变量。

如果我们在核函数外部定义了一个 __device__ int my_count = 10; 或者 __constant__ float pi = 3.14;，在 Host 端（CPU）的代码里，我们并不知道 my_count 在 GPU 上的具体物理内存地址。

符号映射： my_count 只是一个符号（Symbol）。cudaMemcpyFromSymbol 会自动去查找这个符号对应的 GPU 物理地址，然后把数据搬过来。

文件作用域中的变量：可见性与可访问性

在CUDA编程中，你需要控制主机和设备这两个地方的操作。一般情况下，设备核函数不能访问主机变量，并且主机函数也不能访问设备变量，即使这些变量在同一文件作用域内被声明。

4.2 内存管理

CUDA编程的内存管理与C语言的类似，需要程序员显式地管理主机和设备之间的数据移动。现在，工作重点在于如何使用CUDA函数来显式地管理内存和数据移动。

分配和释放设备内存
在主机和设备之间传输数据

4.2.1 内存分配和释放

CUDA编程模型假设了一个包含一个主机和一个设备的异构系统，每一个异构系统都有自己独立的内存空间。核函数在设备内存空间中运行，CUDA运行时提供函数以分配和释放设备内存。

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

void **devPtr (指向指针的指针)：需要传入一个指针的地址。函数执行成功后，它会将分配好的 GPU 内存首地址写入你提供的指针中。

size_t count (字节数)：想要申请的内存大小，单位是字节（Byte）。通常计算方式为：元素个数 * sizeof(数据类型)。

返回值 cudaError_t：返回一个错误代码。如果成功，返回 cudaSuccess；如果显存不足或参数错误，会返回对应的错误类型。

一旦一个应用程序不再使用已分配的全局内存，那么可以以下代码释放该内存空间：

1	cudaError_t cudaFree(void *devPtr);

参数：直接传入之前 cudaMalloc 得到的那个 GPU 指针即可。

作用：回收该指针指向的全局内存空间。

4.2.2 内存传输

一旦分配好了全局内存，就可以使用下列函数从主机向设备传输数据：

1	cudaErrort cudaMemcpy(void dst, const void src, size_t count, enum cudaMemcpyKind kind);

void *dst (目的地)：指向接收数据内存地址的指针。

const void *src (源地址)：指向发送数据内存地址的指针。

size_t count (字节数)：要拷贝的数据大小，单位是字节。

enum cudaMemcpyKind kind (拷贝方向)：这是一个枚举类型，告诉 CUDA 数据往哪儿搬。

函数从内存位置src复制了count字节到内存位置dst。变量kind指定了复制的方向，可以有下列取值：

cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice

这里通过代码（memTransfer.cu来举例

int main(int argc, char **argv)
{
    // set up device
    int dev = 0;
    CHECK(cudaSetDevice(dev));

    // memory size
    unsigned int isize = 1 << 22;
    unsigned int nbytes = isize * sizeof(float);

    // get device information
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("%s starting at ", argv[0]);
    printf("device %d: %s memory size %d nbyte %5.2fMB\n", dev,
           deviceProp.name, isize, nbytes / (1024.0f * 1024.0f));

    // allocate the host memory
    float *h_a = (float *)malloc(nbytes);

    // allocate the device memory
    float *d_a;
    CHECK(cudaMalloc((float **)&d_a, nbytes));

    // initialize the host memory
    for(unsigned int i = 0; i < isize; i++) h_a[i] = 0.5f;

    // transfer data from the host to the device
    CHECK(cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice));

    // transfer data from the device to the host
    CHECK(cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost));

    // free memory
    CHECK(cudaFree(d_a));
    free(h_a);

    // reset device
    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

编译并执行代码

1 2	nvcc -O3 memTransfer.cu -o memTransfer nsys profile --stats=true ./memTransfer

得到关键输出结果

Time (%)  Total Time (ns)  Num Calls   Avg (ns)    Med (ns)   Min (ns)  Max (ns)  StdDev (ns)               Name             
--------  ---------------  ---------  ----------  ----------  --------  --------  -----------  ------------------------------
    85.1         23839444          1  23839444.0  23839444.0  23839444  23839444          0.0  cudaDeviceReset               
    12.9          3613021          2   1806510.5   1806510.5   1735295   1877726     100713.9  cudaMemcpy                    
     1.0           286902          1    286902.0    286902.0    286902    286902          0.0  cudaFree                      
     1.0           270127          1    270127.0    270127.0    270127    270127          0.0  cudaMalloc                    
     0.0             8615          1      8615.0      8615.0      8615      8615          0.0  cuCtxSynchronize              
     0.0              439          1       439.0       439.0       439       439          0.0  cuModuleGetLoadingMode        
     0.0              151          1       151.0       151.0       151       151          0.0  cudaGetDeviceProperties_v12000

API 名称	时间占比 (除去 Reset)	含义	诊断
`cudaMemcpy`	~86% (剩余时间中)	主机与设备间数据传输	当前的瓶颈。数据传输耗时约 3.6ms，在小型任务中这非常常见。
`cudaMalloc` / `cudaFree`	~14% (剩余时间中)	显存申请与释放	申请和释放各占约 0.28ms。在循环中频繁调用会导致严重卡顿。
内核函数 (Kernel)	0%	实际的 GPU 计算	关键异常点：列表中没有看到自定义的 Kernel（如 `my_kernel<<<...>>>`）。

4.2.3 固定内存

分配的主机内存默认是pageable（可分页），它的意思也就是因页面错误导致的操作，该操作按照操作系统的要求将主机虚拟内存上的数据移动到不同的物理位置。

GPU不能在可分页主机内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数据时，它无法控制。当从可分页主机内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，然后从固定内存传输数据给设备内存。

1	cudaError_t cudaMallocHost(void **devPtr, size_t count);

函数分配了count字节的主机内存，这些内存是页面锁定的并且对设备来说是可访问的。由于固定内存能被设备直接访问，所以它能用比可分页内存高得多的带宽进行读写。然而，分配过多的固定内存可能会降低主机系统的性能

固定主机内存必须通过下述指令来释放：

1	cudaError_t cudaFreeHost(void *ptr);

编译并执行代码memTransfer.cu，其中用固定主机内存替换可分页内存

1 2	nvcc -O3 pinMemTransfer.cu -o pinMemTransfer nsys profile --stats=true ./memTransfer

得到结果如下

Time (%)  Total Time (ns)  Num Calls   Avg (ns)    Med (ns)   Min (ns)  Max (ns)  StdDev (ns)               Name             
--------  ---------------  ---------  ----------  ----------  --------  --------  -----------  ------------------------------
    63.6         25200135          1  25200135.0  25200135.0  25200135  25200135          0.0  cudaDeviceReset               
    21.4          8467790          1   8467790.0   8467790.0   8467790   8467790          0.0  cudaHostAlloc                 
     8.1          3197511          2   1598755.5   1598755.5   1229537   1967974     522153.8  cudaMemcpy                    
     4.8          1900381          1   1900381.0   1900381.0   1900381   1900381          0.0  cudaFreeHost                  
     1.6           618805          1    618805.0    618805.0    618805    618805          0.0  cudaMalloc                    
     0.6           252393          1    252393.0    252393.0    252393    252393          0.0  cudaFree                      
     0.0            10658          1     10658.0     10658.0     10658     10658          0.0  cuCtxSynchronize              
     0.0             1036          1      1036.0      1036.0      1036      1036          0.0  cuModuleGetLoadingMode        
     0.0              233          1       233.0       233.0       233       233          0.0  cudaGetDeviceProperties_v12000

与可分页内存相比，固定内存的分配和释放成本更高，但是它为大规模数据传输提供了更高的传输吞吐量。

4.2.4 零拷贝内存

通常来说，主机不能直接访问设备变量，同时设备也不能直接访问主机变量。但有一个例外：零拷贝内存。主机和设备都可以访问零拷贝内存。

在CUDA核函数中使用零拷贝内存有以下几个优势。

当设备内存不足时可利用主机内存
避免主机和设备间的显式数据传输
提高PCIe传输率

当使用零拷贝内存来共享主机和设备间的数据时，你必须同步主机和设备间的内存访问，同时更改主机和设备的零拷贝内存中的数据将导致不可预知的后果。

通常情况下，GPU 只能访问自己的显存 (VRAM)。如果你想处理 CPU 里的数据，必须经历：

CPU 内存 $\rightarrow$ 驱动固定内存 $\rightarrow$ GPU 显存。

而零拷贝内存通过将主机内存映射到 GPU 的地址空间，使得 GPU 在执行内核函数（Kernel）时，如果需要某个数据，就直接发一个请求给 PCIe 总线，从 CPU 内存里“抓”过来。

实现零拷贝主要靠 cudaHostAlloc 函数，并配合 cudaHostAllocMapped 标志。

分配内存： 使用 cudaHostAlloc 分配页锁定内存（Pinned Memory）。
设置标志： 传入 cudaHostAllocMapped，告诉驱动这块内存要映射给 GPU。
获取设备指针： 使用 cudaHostGetDevicePointer 获取一个供 GPU Kernel 使用的地址。

4.2.5 统一虚拟寻址

计算能力为2.0及以上版本的设备支持一种特殊的寻址方式，称为统一虚拟寻址（UVA）。有了UVA，由指针指向的内存空间对应用程序代码来说是透明的。通过UVA，由cudaHostAlloc分配的固定主机内存具有相同的主机和设备指针。因此，可以将返回的指针直接传递给核函数。

4.2.6 统一内存寻址

统一内存中创建了一个托管内存池，内存池中已分配的空间可以用相同的内存地址（即指针）在CPU和GPU上进行访问。统一内存寻址依赖于UVA的支持，但它们是完全不同的技术。UVA为系统中的所有处理器提供了一个单一的虚拟内存地址空间。但是，UVA不会自动将数据从一个物理位置转移到另一个位置，这是统一内存寻址的一个特有功能。

4.3 内存访问模式

大多数设备端数据访问都是从全局内存开始的，并且多数GPU应用程序容易受内存带宽的限制。因此，最大限度地利用全局内存带宽是调控核函数性能的基本。

CUDA执行模型的显著特征之一就是指令必须以线程束为单位进行发布和执行。存储操作也是同样。在执行内存指令时，线程束中的每个线程都提供了一个正在加载或存储的内存地址。在线程束的32个线程中，每个线程都提出了一个包含请求地址的单一内存访问请求，它并由一个或多个设备内存传输提供服务。

4.3.1 对齐与合并访问

全局内存通过缓存来实现加载/存储。全局内存是一个逻辑内存空间，我们可以通过核函数访问它。所有的应用程序数据最初存在于DRAM上，即物理设备内存中。核函数的内存请求通常是在DRAM设备和片上内存间以128字节或32字节内存事务来实现的。

一行一级缓存是128个字节，它映射到设备内存中一个128字节的对齐段。如果线程束中的每个线程请求一个4字节的值，那么每次请求就会获取128字节的数据，这恰好与缓存行和设备内存段的大小相契合。

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

·对齐内存访问

·合并内存访问

当设备内存事务的第一个地址是用于事务服务的缓存粒度的偶数倍时（32字节的二级缓存或128字节的一级缓存），就会出现对齐内存访问。运行非对齐的加载会造成带宽浪费。

对齐合并内存访问的理想状态是线程束从对齐内存地址开始访问一个连续的内存块。为了最大化全局内存吞吐量，为了组织内存操作进行

对齐合并是很重要的。

对齐与合并内存的加载操作如图：

在这种情况下，只需要一个128字节的内存事务从设备内存中读取数据。

非对齐和未合并的内存访问如图：

在这种情况下，需要3个128字节的内存事务来从设备内存中读取数据：一个在偏移量为0的地方开始，读取连续地址之后的数据；一个在偏移量为256的地方开始，读取连续地址之前的数据；另一个在偏移量为128的地方开始读取大量的数据。

4.3.2 全局内存读取

在SM中，数据通过以下3种缓存/缓冲路径进行传输：

一级和二级缓存
常量缓存
只读缓存

内存加载访问模式：

内存加载可以分为两类：缓存加载（启用一级缓存）和没有缓存的加载（禁用一级缓存）

内存加载的访问模式可以分为：

有缓存与没有缓存：如果启用一级缓存，则内存加载被缓存
对齐与非对齐：如果内存访问的第一个地址是32字节的倍数，则对齐加载
合并与未合并：如果线程束访问一个连续的数据块，则加载合并

4.3.2.1 缓存加载

缓存加载操作经过一级缓存，在粒度为128字节的一级缓存行上由设备内存事务进行传输。缓存加载可以分为对齐/非对齐及合并/非合并。

**1）**理想情况下：线程束中所有线程请求的地址都在128字节的缓存行范围内。完成内存加载操作只需要一个128字节的事务。总线的使用率为100%。

**2）**另一种情况：访问是对齐的，引用的地址不是连续的线程ID，而是128字节范围内的随机值。**由于线程束中线程请求的地址仍然在一个缓存行范围内，所以只需要一个128字节的事务来完成这一内存加载操作。**总线利用率仍然是100%

**3）**还有一种情况就是：线程束请求32个连续4个字节的非对齐数据元素。在全局内存中线程束的线程请求的地址落在两个128字节段范围内。

原因：当启用一级缓存时，由SM执行的物理加载操作必须在128个字节的界线上对齐，所以要求有两个128字节的事务来执行这段内存加载操作。总线利用率为50%。

**4）**线程束中所有线程请求相同的地址。因为被引用的字节落在一个缓存行范围内，所以只需请求一个内存事务，但总线利用率非常低。如果加载的值是4字节的，则总线利用率是4字节请求/128字节加载＝3.125%。

**5）**最坏的情况：线程束中线程请求分散于全局内存中的32个4字节地址。尽管线程束请求的字节总数仅为128个字节，但地址要占用N个缓存行（0＜N≤32）。

注意：

为了实现高带宽并行访问，Shared Memory 被物理地划分为 32 个等大小的存储模块，称为 Banks。

宽度：每个 Bank 的宽度通常是 4 字节（32-bit，对应一个 float 或 int）。
映射规则：地址是交错映射的。连续的 4 字节落在连续的 Bank 上。
- Bank 0 存储地址：0, 128, 256…
- Bank 1 存储地址：4, 132, 260…
- 计算公式： $\text{Bank Index} = (\text{Byte Address} \div 4) \pmod{32}$

例如这图中就会出现bank冲突

线程束（Warp）中的线程同时访问了第一个 128 字节段（深灰色）和第二个 128 字节段（中灰色）。

相同偏移量： 注意看箭头指向的位置。如果有两个线程分别指向了地址 $A$ 和地址 $A + 128$ （或者 $A + 256$ ），它们虽然在内存中相隔很远，但物理上映射到了同一个 Bank。这两个线程在同一个周期内请求同一个 Bank 的不同数据，硬件必须串行处理，这就是 2-way Bank Conflict。

CPU一级缓存和GPU一级缓存之间的差异：

CPU一级缓存优化了时间和空间局部性。GPU一级缓存是专为空间局部性而不是为时间局部性设计的。频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。

4.3.2.2 没有缓存的加载

没有缓存的加载不经过一级缓存，它在内存段的粒度上（32个字节）而非缓存池的粒度（128个字节）执行。这是更细粒度的加载，可以为非对齐或非合并的内存访问带来更好的总线利用率。

**1）**理想情况：对齐与合并内存访问。128个字节请求的地址占用了4个内存段，总线利用率为100%。

**2）**情况：内存访问是对齐的且线程访问是不连续的，而是在128个字节的范围内随机进行。

**3）**线程束请求32个连续的4字节元素但加载没有对齐到128个字节的边界。请求的地址最多落在5个内存段内，总线利用率至少为80%。

与这些类型的请求缓存加载相比，使用非缓存加载会提升性能，这是因为加载了更少的未请求字节。

**4）**线程束中所有线程请求相同的数据。地址落在一个内存段内，总线的利用率是请求的4字节/加载的32字节＝12.5%（4/32）。

在这种情况下，非缓存加载性能也是优于缓存加载的性能。

**5）**最坏的情况：线程束请求32个分散在全局内存中的4字节字。由于请求的128个字节最多落在N个32字节的内存分段内而不是N个128个字节的缓存行内

4.3.2.3 非对齐读取

因为访问模式往往是由应用程序实现的一个算法来决定的，所以对于某些应用程序来说合并内存加载是一个挑战。

我们以第三章的向量加法代码进行修改，同时为了说明核函数中非对齐访问对性能的影响，去掉所有的内存加载操作，来指定一个偏移量。

__global__ void readOffset(float *A, float *B, float *C, const int n,
                           int offset)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;

    if (k < n) C[i] = A[k] + B[k];
}

同时为与核函数同步，主机端代码也进行了修改

void sumArraysOnHost(float *A, float *B, float *C, const int n, int offset)
{
    for (int idx = offset, k = 0; idx < n; idx++, k++)
    {
        C[k] = A[idx] + B[idx];
    }
}

编译代码readSegment.cu

1	nvcc -O3 -arch=sm_120 readSegment.cu -o readSegment

由于规模较小，波动较大，单运行的结果很难看出程序的差异，因此我们通过ncu进行分析偏移0，11，128的结果：ncu 抓取到的底层扇区数据直接揭示了硬件效率的损耗。

1	ncu --metrics lts__t_sectors_srcunit_tex_op_read.sum,lts__t_sectors_srcunit_tex_op_write.sum -k readOffset ./readSegment 0/11/128

根据结果得到下面的表格

偏移量 (Offset)	字节偏移 (Bytes)	读取扇区 (Read)	写入扇区 (Write)	访存状态描述
0	0	262,144	131,072	完美合并与对齐
11	44	266,180	131,071	非对齐访问 (Misaligned)
128	512	262,112	131,056	重新对齐 (Realignment)

可以发现128 和 0 的结果几乎一样

这是因为： 512 是 128 的整数倍，虽然跳过了前 128 个元素，但线程 0 访问的地址依然是一个 Cache Line 的起始边界。

而当 offset = 11 时，字节偏移为 44 字节。读取扇区数增加到了 266,180。这多出来的约 4,000 个扇区，正是硬件为了补齐那些“跨界”访问而额外进行的读取操作。

通过以下命令得到程序的加载效率

1	ncu --metrics smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct -k readOffset ./readSegment 0/11/128

总结得到下面的表格

偏移量 (Offset)	字节对齐 (Offset×4)	全局加载效率 (Metric Value)	结论分析
0	0 Bytes	100%	完美合并：线程请求与硬件缓存行完全重合。
11	44 Bytes	80%	非对齐访问：跨越了缓存行边界，导致带宽浪费。
128	512 Bytes	100%	重新对齐：512 是 128（缓存行大小）的倍数，效率回归。

禁用一级缓存

为了强制执行没有缓存的加载，重新编译代码并增加了以下nvcc选项：

1	nvcc -O3 -Xptxas -dlcm=cg -arch=sm_120 readSegment.cu -o readSegment

我们同样通过ncu进行分析偏移0，11，128的结果

1	ncu --metrics lts__t_sectors_srcunit_tex_op_read.sum,lts__t_sectors_srcunit_tex_op_write.sum -k readOffset ./readSegment 0/11/128

可以发现：禁用 L1 后，每一个线程束（Warp）的请求都会直接、无修饰地打在 L2 缓存的扇区（Sectors）上。

偏移量 (Offset)	读取扇区 (Read Sectors)	相比 Offset 0 的增量	访存效率状态
0	262,144	基准 (Baseline)	完美对齐
11	327,676	+ 65,532	非对齐惩罚（最大化）
128	262,112	- 32 (数据量减少)	重新对齐

在之前的实验（开启缓存）中，offset 11 只多出了约 4,000 个扇区；而现在多出了 65,000 个。这说明在现代架构上，L1 缓存在处理非对齐访问时起到了极强的缓冲作用。

可以得到这样的结论：缓存缺失对非对齐访问的性能影响更大。如果启用缓存，一个非对齐访问可能将数据存到一级缓存，这个一级缓存用于后续的非对齐内存访问。但是，如果没有一级缓存，那么每一次非对齐请求需要多个内存事务，并且对将来的请求没有作用。

4.3.2.4 只读缓存

只读缓存最初是预留给纹理内存加载使用的。对计算能力为3.5及以上的GPU来说，只读缓存也支持使用全局内存加载代替一级缓存。只读缓存（Read-Only Cache），也常被称为 LDG 缓存（因其对应的汇编指令 LDG 而得名），是一个专门为不经常变动的数据设计的读取路径。

只读缓存通常与 Texture Cache（纹理缓存） 合并，并与 L1 缓存共享硬件资源。

数据特性：它专门用于存储在内核（Kernel）执行期间不会被修改的数据。
硬件路径：当编译器确定数据是只读时，它会跳过标准的全局内存加载路径，转而使用特定的只读数据路径。这通常能提供更高的标签带宽（Tag Bandwidth）和更灵活的访存处理。

如何使用只读缓存：

A.使用 __restrict__ 和 const 指针

这是最推荐的方式。通过修饰符告诉编译器，该指针指向的数据在当前作用域内不会被修改且没有别名。

C++

1
2
3

__global__ void readOffset(const float* __restrict__ A, 
                           const float* __restrict__ B, 
                           float* C, ...)

B. 使用 __ldg() 内置函数

如果你不想修改函数签名，可以直接在读取时显式强制使用只读路径：

C++

1	C[i] = __ldg(&A[k]) + __ldg(&B[k]);

只读缓存 vs. 普通全局加载

特性	普通全局加载 (L1/L2)	只读缓存 (LDG/TEX)
数据一致性	需维护缓存一致性（处理读写冲突）	不维护一致性（假设数据不改）
访问粒度	通常为 32 字节扇区	同样为 32 字节，但对非对齐访问更友好
应用场景	频繁读写的通用数组	矩阵乘法中的权重、查找表、常量输入

4.3.3 全局内存写入

内存的存储操作相对简单。大多数情况下，全局内存写入会直接绕过 L1 缓存，直接写入 L2 缓存 或显存（DRAM）。

存储操作在32个字节段的粒度上被执行。内存事务可以同时被分为一段、两段或四段。例如，如果两个地址同属于一个128个字节区域，但是不属于一个对齐的64个字节区域，则会执行一个四段事务（也就是说，执行一个四段事务比执行两个一段事务效果更好）。

**1）**理想情况：内存访问是对齐的，并且线程束里所有的线程访问一个连续的128字节范围。存储请求由一个四段事务实现。

**2）**内存访问是对齐的，但地址分散在一个192个字节范围内的情况。存储请求由3个一段事务来实现。

**3）**内存访问是对齐的，并且地址访问在一个连续的64个字节范围内的情况。这种存储请求由一个两段事务来完成。

为了验证非对齐对内存存储效率的影响，按照下面的方式修改向量加法核函数。

__global__ void writeOffset(float *A, float *B, float *C, const int n,
                            int offset)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;

    if (k < n) C[k] = A[i] + B[i];
}

同样修改主机端代码

void sumArraysOnHost(float *A, float *B, float *C, const int n, int offset)
{
    for (int idx = offset, k = 0; idx < n; idx++, k++)
    {
        C[idx] = A[k] + B[k];
    }
}

编译代码writeSegment.cu`

1	nvcc -O3 -arch=sm_120 writeSegment.cu -o writeSegment

通过以下命令获取指标

ncu --metrics smsp__sass_average_data_bytes_per_sector_mem_global_op_st.pct,smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct,gpu__compute_memory_throughput.avg.pct -k writeOffset ./writeSegment 0/11/128

偏移量 (Offset)	加载效率 (LD %)	存储效率 (ST %)	状态诊断
0	100%	100%	完美合并写入。线程束写入的数据与硬件缓存行完全对齐。
11	100%	80.00%	非对齐写入。写入起始地址偏移了 44 字节，导致跨越了缓存行边界。
128	100%	100%	重新对齐写入。 $128 \times 4$ 字节正好是 512 字节（128 字节的倍数），效率回归 100%。

4.3.4 结构体数组与数组结构体

数组结构体（AoS）和结构体数组（SoA）。这是一个有趣的话题，因为当存储结构化数据集时，它们代表了可以采用的两种强大的数据组织方式（结构体和数组）。

AoS (结构体数组)：面向对象风格

struct Particle {
    float x, y, z;
};
Particle particles[N];

需要同时处理一个点的 x, y, z，它们在内存中是挨着的。

SoA (数组结构体)：面向硬件风格

struct Particles {
    float x[N];
    float y[N];
    float z[N];
};
Particles p;

访存局部性好：如果你只需要处理所有点的 x 坐标，它们在内存中是连续的。

AoS模式在GPU上存储示例数据并执行一个只有x字段的应用程序，将导致50%的带宽损失，因为y值在每32个字节段或128个字节缓存行上隐式地被加载。AoS格式也在不需要的y值上浪费了二级缓存空间。
SoA模式存储数据充分利用了GPU的内存带宽。由于没有相同字段元素的交叉存取，GPU上的SoA布局提供了合并内存访问，并且可以对全局内存实现更高效的利用。

4.3.4.1 AoS数据布局的简单数学运算

定义结构体如下

struct innerStruct
{
    float x;
    float y;
};

struct innerArray
{
    float x[LEN];
    float y[LEN];
};

核函数如下

__global__ void testInnerStruct(innerStruct *data, innerStruct * result,
                                const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        innerStruct tmp = data[i];
        tmp.x += 10.f;
        tmp.y += 20.f;
        result[i] = tmp;
    }
}

编译并执行代码simpleMathAoS.cu

1 2	nvcc -O3 -arch=sm_120 simpleMathAoS.cu -o simpleMathAoS ./simpleMathAoS

可以得到如下结果

1	innerstruct <<< 32768, 128 >>> elapsed 0.000430 sec

输入命令如下命令获取全局加载效率和全局存储效率

1	ncu --metrics smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct,smsp__sass_average_data_bytes_per_sector_mem_global_op_st.pct ./simpleMathAoS

可以得到以下结果

指标名称	测量值	硬件层面的真相
Global Load Efficiency (ld.pct)	50%	线程束（Warp）为了拿到需要的字段，搬运了 2 倍于逻辑需求的数据量。
Global Store Efficiency (st.pct)	50%	写回结果时，由于数据跨度（Stride）的存在，导致一半的写入带宽被“空跑”浪费了。

分析代码可以发现，结构体包含两个 float 成员（例如 float x, y;）。

当 Warp 中的 32 个线程尝试同时读取所有的 a 时，地址变得不连续。

线程 0 读地址 $N$ ，线程 1 读地址 $N+8$ 。

为了覆盖这 32 个 a 成员，原本 128 字节的数据散落在 256 字节的物理空间里。

计算公式： $\text{Efficiency} = \frac{128 \text{ (有效字节)}}{256 \text{ (搬运字节)}} = 50\%$ 。

4.3.4.2 SoA数据布局的简单数学运算

结构体组织：

struct InnerArray {
    float x[LEN];
    float y[LEN];
};

编译并执行代码simpleMathSoA.cu

1 2	nvcc -O3 -arch=sm_120 simpleMathSoA.cu -o simpleMathSoA ./simpleMathSoA

可以得到以下结果

1	innerarray <<< 32768, 128 >>> elapsed 0.000345 sec

输入命令如下命令获取全局加载效率和全局存储效率

1	ncu --metrics smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct,smsp__sass_average_data_bytes_per_sector_mem_global_op_st.pct ./simpleMathSoA

得到结果并与AoS对比

布局模式	加载效率 (ld.pct)	存储效率 (st.pct)	硬件行为描述
AoS (之前)	50%	50%	成员交错导致步长访问 (Stride)，浪费一半带宽。
SoA (当前)	100%	100%	合并访问，请求与物理搬运完全匹配。

在 SoA 模式下，所有的 x 元素在内存中是完全连续排布的，所有的 y 元素也是如此。

线程映射：当线程束 (Warp) 执行 float tmpx = data->x[i]; 时，线程 0 访问 x[0]，线程 1 访问 x[1]。

合并成功：这 32 个连续的 float 刚好构成了物理内存上的一个连续 128 字节块。硬件只需发起一次内存事务就能取回所有数据，没有任何带宽浪费。

4.3.5 性能调整

优化设备内存带宽利用率有两个目标：

对齐及合并内存访问，以减少带宽的浪费
足够的并发内存操作，以隐藏内存延迟

前面已经考虑了对齐合并的内存访问。后续则考虑并发内存访问最大化。

4.3.5.1 展开技术

每个线程都执行4个独立的内存操作。因为每个加载过程都是独立的，所以可以调用更多的并发内存访问。

__global__ void readOffsetUnroll4(float *A, float *B, float *C, const int n,
                                  int offset)
{
    unsigned int i = blockIdx.x * blockDim.x * 4 + threadIdx.x;
    unsigned int k = i + offset;

    if (k < n) C[i] = A[k] + B[k];
    if (k + blockDim.x < n) {
        C[i + blockDim.x] = A[k + blockDim.x] + B[k + blockDim.x];
    }
    if (k + 2 * blockDim.x < n) {
        C[i + 2 * blockDim.x] = A[k + 2 * blockDim.x] + B[k + 2 * blockDim.x];
    }
    if (k + 3 * blockDim.x < n) {
        C[i + 3 * blockDim.x] = A[k + 3 * blockDim.x] + B[k + 3 * blockDim.x];
    }
}

执行并编译代码readSegmentUnroll.cu

1 2	nvcc -arch=sm_120 readSegmentUnroll.cu -o readSegmentUnroll ./readSegmentUnroll

输入ncu命令

ncu --metrics \
smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct,\
smsp__sass_average_data_bytes_per_sector_mem_global_op_st.pct,\
lts__t_sectors_srcunit_tex_op_read.sum,\
lts__t_sectors_srcunit_tex_op_write.sum \
./writeSegment 11

核函数	读取效率 (ld.pct)	写入效率 (st.pct)	读取总扇区数 (Read Sum)	写入总扇区数 (Write Sum)
readOffset	80.00%	100.00%	266,184	131,071
Unroll2	80.00%	100.00%	264,188	131,071
Unroll4	80.00%	100.00%	263,436	131,071

随着展开倍数的增加，lts__t_sectors_srcunit_tex_op_read.sum 在轻微下降：

原始 (266,184) $\rightarrow$ Unroll4 (263,436)：减少了约 2,748 个扇区。
逻辑分析：循环展开允许编译器生成更强大的加载指令（如 LDG.E.128），并让 GPU 调度器有更大的机会在多个连续请求之间进行 L2 缓存合并。

4.3.5.2 增大并行性

为了充分体现并行性，可以用一个核函数启动的网格和线程块大小进行试验，以找到该核函数最佳的执行配置。此处不再赘述。

4.4 统一内存

传统 CUDA 编程中，数据迁移是显式的、粗粒度的。必须明确告诉系统：“把这块数据从 CPU 搬到 GPU”，然后启动核函数，最后再把结果搬回来。NVIDIA 的 CUDA Unified Memory（统一内存）改变了这一切。它的核心思想不是消除内存差异，而是为 CPU 和 GPU 构建一个共享的虚拟地址空间。

统一虚拟寻址（UVA）：共用一套地址命名体系

统一内存中创建一个托管内存池（CPU上有，GPU上也有），内存池中已分配的空间可以通过相同的指针直接被CPU和GPU访问，底层系统在统一的内存空间中自动的进行设备和主机间的传输。数据传输对应用是透明的，大大简化了代码。

4.4.1 懒加载 + 自动搬运

统一内存以 4KB 页面为单位进行管理。当你分配一块 cudaMallocManaged 内存时，系统并不会立即为其分配物理页，也不会预先把所有数据复制到 GPU。只有当某个线程首次访问某一页时，才会真正触发分配和迁移。

举个例子：你在 CPU 上初始化数组 a 和 b，此时它们驻留在主机内存；当 GPU 核函数第一次读取 a[0] 时，MMU 发现该页不在显存中，于是抛出缺页异常。CUDA 运行时捕获该异常，将对应页面从主机复制到设备显存，并更新 GPU 的页表映射。整个过程对核函数完全透明。

这种“按需加载”的策略显著降低了初始化开销，尤其适合那些仅部分数据会被实际使用的场景。

4.4.2 内存一致性保障

所有对统一内存区域的读写操作都保证全局顺序一致性。也就是说，一旦某个核函数修改了数据并完成执行，后续任何处理器（CPU 或其他 GPU）对该数据的读取都能看到最新值。

这得益于 CUDA 流水线中的隐式同步点。例如，cudaDeviceSynchronize() 不仅等待核函数完成，也确保所有相关的页面迁移和缓存刷新已完成。因此，开发者通常无需额外插入内存屏障指令。

4.4.3 统一内存用法

统一内存的基本思路就是减少指向同一个地址的指针，比如我们经常见到的，在本地分配内存，然后传输到设备，然后在从设备传输回来，使用统一内存，就没有这些显式的需求了，而是驱动程序帮我们完成。

1
2
3

cudaMallocManaged(&a, bytes);
cudaMallocManaged(&b, bytes);
cudaMallocManaged(&c, bytes);

CUDA 的内核启动（Kernel Launch）和部分内存拷贝是异步的。为保证CPU与GPU协同处理，采用下列函数阻塞主机端

1	cudaDeviceSynchronize()

该函数会阻塞主机端（CPU）线程，直到设备端（GPU）之前发出的所有任务（包括所有流中的内核、内存拷贝等）全部执行完毕。

然而虽然简洁，但也隐藏了一些潜在开销：每次跨端访问未驻留页面都会带来延迟。对于性能敏感的应用，可通过 cudaMemPrefetchAsync 主动预取数据到目标设备，避免运行时卡顿。例如：

// 预先将数据推送到 GPU
cudaMemPrefetchAsync(a, bytes, 0);  // 0 表示 GPU 0
cudaMemPrefetchAsync(b, bytes, 0);
cudaMemPrefetchAsync(c, bytes, 0);

下列完整代码：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vector_add(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];
    }
}

int main() {
    const int N = 1 << 20;
    size_t bytes = N * sizeof(float);

    float *a, *b, *c;
    cudaMallocManaged(&a, bytes);
    cudaMallocManaged(&b, bytes);
    cudaMallocManaged(&c, bytes);

    // CPU 初始化
    for (int i = 0; i < N; ++i) {
        a[i] = 1.0f;
        b[i] = 2.0f;
    }

    // 启动核函数
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    vector_add<<<gridSize, blockSize>>>(a, b, c, N);

    cudaDeviceSynchronize();

    printf("Result: %f\n", c[0]);  // 自动触发回迁（如有必要）

    cudaFree(a); cudaFree(b); cudaFree(c);
    return 0;
}

编译并运行代码Unified_Memory.cu

1 2	nvcc -arch=sm_120 Unified_Memory.cu -o Unified_Memory ./Unified_Memory

与全局内存对比来说

特性	全局内存 (Global Memory)	统一内存 (Unified Memory)	性能影响因素
分配方式	`cudaMalloc`	`cudaMallocManaged`	UM 涉及驱动层的页表映射。
数据传输	手动 `cudaMemcpy`	硬件/驱动按需迁移 (Page Fault)	UM 存在“首次访问”延迟。
访存效率	开发者完全控制，易实现 100% 合并	依赖 Page Migration 粒度	UM 在数据分布不均时易导致总线拥堵。
编程复杂度	高（需手动维护副本）	低（单一指针，CPU/GPU 共用）	UM 显著提升生产力，但可能隐藏性能瓶颈。

4.5 总结

本文深入剖析了 CUDA 全局内存的架构与优化机制，通过对比内存层次结构（如寄存器、共享内存、常量内存等）与访问模式（对齐、合并、非对齐），阐明了硬件底层事务如何影响指令带宽，并重点介绍了通过 SoA（数组结构体） 布局优化访存效率、利用 L1/L2 缓存与只读缓存 提升吞吐量，以及使用统一内存（Unified Memory） 简化异构编程复杂度的核心技术与性能调试方法。