【CUDA】内存模型

原创已于 2025-05-16 10:41:20 修改 · 990 阅读

30 ·

CC 4.0 BY-SA版权

文章标签：

#gpu算力 #c++

于 2025-02-06 17:36:53 首次发布

GPU 专栏收录该内容

11 篇文章

订阅专栏

一、Programmable

1.1 寄存器(Registers)

1.2 本地内存(Local Memory)

1.3 共享内存(shared Memory)

1.4 常量内存(Constant Memory)

1.5 全局内存(Global Memory)

1.6 纹理内存(Textrue Memory)

1.7 总结

二、Cache(Non-programmable)

对于程序员而言，memory可以分为下面两类：

Programmable：可以灵活操作的部分。
Non-programmable：不能操作，由一套自动机制来达到很好的性能

一、Programmable

在CUDA中可编程内存的类型有：

寄存器(Registers)
本地内存(Local Memory)
共享内存(Shared Memory)
常量内存(Constant Memory)
纹理内存(Texture Memory)
全局内存(Global Memory)

一个线程拥有自己私有的寄存器和局部内存(local memory)
一个block内的所有线程可以读写该block专有的共享内存(shared memory)
所有线程可以读写全局内存(Global memory)，只读常量内存(constant memory)和纹理内存(texture memory)

1.1 寄存器(Registers)

在内核函数中声明且没有其他修饰符修饰的变量通常存放在GPU的寄存器中。如：下面代码中的线程索引变量i。寄存器通常用于存放内核函数中需要频繁访问的线程私有变量，这些变量与内核函数的生命周期相同，内核函数执行完毕后，就不能再进行访问了

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

寄存器是稀有资源。在Fermi上，每个thread限制最多拥有63个register，Kepler则是255个。让kernel使用较少的register就能够允许更多block驻留在SM中，就增加了Occupancy，提升了性能

可以通过如下编译选项查看线程的寄存器，共享内存的使用情况

-Xptxas -v,-abi=no

若使用的寄存器超出了硬件限制，那么多出来的部分就会存放在Local memory里

CUDA编译器提供了关键字__launch_bound__启发式地限制寄存器的使用，从而提高SM内active block数量

__global__ void 
__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) 
kernel(...) {
 // your kernel body
}

heuristics(启发式)一词，说通俗一点，就是在代码中显示告知编译器，该Kernel函数每个Block分配的最大线程数以及SM最小支持的active block数，这样编译器就可以调整每个线程拥有的寄存器数，优化SM中active block的数量了

maxThreadPerblock：参数指定了一个Block内最大线程数
minBlockkPerMultiprocessor：一个SM最少支持的active block数。这个参数不是必需的，编译器会根据不同的GPU架构赋值

另一种方法限制线程使用的寄存器的数量，采用如下编译参数，

-maxrregcount=32

所有CUDA线程的寄存器都被限制在32个，除非某个Kernel函数显示的使用__launch_bounds__

1.2 本地内存(Local Memory)

线程私有

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

编译时使用未知索引引用的本地数组（无法在编译期间确定具体的访问模式）
可能会占用大量寄存器空间的较大本地结构体或者数组
任何不满足内核函数寄存器限定条件的变量

物理上并不存在单独的Local memory。编译器会将其放入片外的DRAM中(与Global Memory相同)，访存延迟大，带宽小

1.3 共享内存(shared Memory)

共享内存是SM的私有资源，属于on-chip内存，带宽相对较高，延迟也较低

在内核函数中被__shared__修饰符修饰的变量被存储到共享内存中。每个SM都有一定数量供线程块分配的共享内存，在内核函数内进行声明，生命周期伴随整个线程块，一个线程块执行结束后，为其分配的共享内存也被释放以便重新分配给其他线程块进行使用。线程块中的线程通过使用共享内存中的数据可以实现互相之间的协作，不过使用共享内存可通过如下函数进行同步：

void __sybcthreads()

该函数为线程块中的所有线程设置了一个执行障碍点，使得同一线程块中的所有线程必须都执行到该障碍点才能往下执行，这样就可以避免一些潜在的数据冲突

1.4 常量内存(Constant Memory)

offchip内存，只读，拥有SM私有的constant cache，因此在cache hit的情况下速度快。常量内存是全局的，对所有Kernel函数可见。因此声明要在Kernel函数外

__constant__ float variable;

常量变量存储在常量内存中，内核函数只能从常量内存中读取数据，常量内存必须在host端代码中进行初始化

cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src,size_t count);

下面的例子展示了如何声明常量内存并与之进行数据交换：

__constant__ float const_data[256];
float data[256];
cudaMemcpyToSymbol(const_data, data, sizeof(data));
cudaMemcpyFromSymbol(data, const_data, sizeof(data));

常量内存适合用于线程束中的所有线程都需要从相同的内存地址中读取数据的情况，如：所有线程都需要的常量参数，每个GPU只可以声明不超过64KB的常量内存

1.5 全局内存(Global Memory)

offchip内存，所有线程可见

一个全局内存变量可以在host代码中使用cudaMalloc函数进行动态声明，或者使用__device__修饰符在device代码中静态声明。全局内存变量可以在任何SM设备中被访问到，其生命周期贯穿应用程序的整个生命周期

静态声明并使用全局变量

#include <cuda_runtime.h>
#include <stdio.h>
 
__device__ float dev_data;
 
__global__ void AddGlobalVariable(void) {
  printf("device, global variable before add: %.2f\n", dev_data);
  dev_data += 2.0f;
  printf("device, global variable after add: %.2f\n", dev_data);
}
 
int main(void) {
  float host_data = 4.0f;
  cudaMemcpyToSymbol(dev_data, &host_data, sizeof(float));
  printf("host, copy %.2f to global variable\n", host_data);
  AddGlobalVariable<<<1, 1>>>();
  cudaMemcpyFromSymbol(&host_data, dev_data, sizeof(float));
  printf("host, get %.2f from global variable\n", host_data);
  cudaDeviceReset();
  return 0;
}

注意：变量 dev_data 只是作为一个标识符存在，并不是 device 端的全局内存变量地址，所以不能直接使用 cudaMemcpy 函数将 host 上的数据拷贝到 device 端。不能直接在 host 端的代码中使用运算符&对 device 端的变量进行取地址操作，因为其只是一个表示 device 端物理位置的符号

不过可以使用如下函数来获取其地址：

cudaError_t cudaGetSymbolAddress(void** devPtr, const void* symbol);

获取地址后，就可以使用 cudaMemcpy 函数进行操作

int main(void) 
{
  float host_data = 4.0f;
  float *dev_ptr = NULL;
  cudaGetSymbolAddress((void **)&dev_ptr, dev_data);
  cudaMemcpy(dev_ptr, &host_data, sizeof(float), cudaMemcpyHostToDevice);
  printf("host, copy %.2f to global variable\n", host_data);
  AddGlobalVariable<<<1, 1>>>();
  cudaMemcpy(&host_data, dev_ptr, sizeof(float), cudaMemcpyDeviceToHost);
  printf("host, get %.2f from global variable\n", host_data);
  cudaDeviceReset();
  return 0;
}

在CUDA编程中，一般情况下 device 端的内核函数不能访问 host 端声明的变量，host 端的函数也不能直接访问 device 端的变量，即使是在同一个文件内声明的

1.6 纹理内存(Textrue Memory)

offchip内存，拥有SM私有的cache，在cache hit的情况下访存速度快，对所有线程可见。纹理内存是一种通过指定的只读缓存访问的全局内存，是对二维空间局部性的优化，所以使用纹理内存访问二维数据的线程可以达到最优性能

texture<type, dim> tex_var;        //Initialize
cudaChannelFormatDesc();           //Options
cudaBindTexture2D(...);            //Bind
tex2D(tex_var, x_index, y_index);  //Fetch

1.7 总结

Local，Global，Contant，Texture为片外DRAM，其中Global，Constant，Texture内存在Host端代码声明，所有线程可见
SM拥有私有的Registers和Shared Memory（其实还有SM私有的L1 cache以及共有的L2 cache），Constant和Texture内存有专有的Caches(片上)

二、Cache(Non-programmable)

GPU上有4种缓存：

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

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

三、固定内存

页锁定内存（Pinned Memory）或固定内存（Fixed Memory），属于主机内存中的一种特殊内存

默认的 host 端的内存是可分页的，其按照操作系统的要求将主机虚拟内存上的数据移动到不同的物理位置。GPU不能在可分页的 host 端内存上安全地访问数据，因为当 host 端操作系统在物理位置上移动该数据时它无法控制。当从可分页的 host 端内存传输数据到 device 端内存时，CUDA 驱动程序会先临时分配页面锁定的或固定的 host 端内存，再将 host 端的数据复制到该内存中，最后从该内存中把数据拷贝到 device 端的内存中

CUDA提供下面的函数，可以直接分配固定的主机内存：

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

device 端可以用很高的带宽进行读写操作。不过，过多地分配固定内存会降低 host 系统的性能，因为能用于虚拟内存的可分页内存数量减少了。固定内存通过下面的函数进行释放：

cudaError_t cudaFreeHost(void *ptr);

cudaHostRegister

cudaHostRegister可将现有的主机内存区域注册为可被GPU访问的页锁定内存

允许将已经分配的内存（如malloc等）注册为锁页内存。适用于需要将现有内存区域用作GPU访问的场景。注册的内存同样受到锁页内存资源限制的影响。在使用完毕后，需要使用cudaUnregisterHostMemory() 来注销内存

四、零拷贝内存

一般情况下 host 不能直接访问 device 端的变量，device 也不能直接访问 host 端的变量。有一种例外的情况，那就是零拷贝内存，host 和 device 都可以访问零拷贝内存。在内核函数中使用零拷贝内存有以下几个优势：

当 device 内存不足时使用 host 内存
避免 device 和 host 之间显示的数据传输
提高 PCIe 传输率

零拷贝内存是固定内存，CUDA 提供下面的函数创建一个固定内存到 device 地址空间的映射：

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

flags 参数可以选择以下几种

cudaHostAllocDefault：使 cudaHostAlloc 函数的行为与 cudaMallocHost 一致
cudaHostAllocPortable：返回能被所有 CUDA 上下文使用的固定内存
cudaHostAllocWriteCombined：返回写结合内存，该内存可以在某些系统配置上通过PCIe总线更快地传输
cudaHostAllocMapped：返回被映射到 device 地址空间的 host 端内存

使用下面的函数可以获取映射到固定内存的 device 端指针：

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

若需要在 host 和 device 之间共享少量的数据，那么零拷贝内存会是一个不错的选择。不过对于需频繁读写的操作，使用零拷贝内存会显著地降低程序的性能，因为每一次映射到内存的传输都需要通过 PCIe 总线进行。另外，使用零拷贝内存必须同步 host 和 device 的内存访问操作以避免潜在的数据冲突

五、统一虚拟寻址

在CC2.0以上的设备支持一种新特性：Unified Virtual Addressing（UVA）。这个特性在CUDA4.0中首次介绍，并被64位Linux系统支持。如下图所示，在使用UVA的情况下，CPU和GPU使用同一块连续的地址空间：

在UVA之前，需要分别管理指向 host memory 和 device memory 的指针。使用UVA之后，实际指向内存空间的指针对用户来说是透明的，看到的是同一块连续地址空间

开启UVA，使用 cudaHostAlloc 分配的 pinned memory 获得的地址对于 device 和 host 来说是通用的。可以直接在kernel里使用这个地址。回看前文，对于zero-copy的处理过程是：

分配已经映射到device的pinned memory
根据获得的host地址，获取device的映射地址
在kernel中使用该映射地址

使用UVA之后，就没必要获取device的映射地址了，直接使用一个地址就可以，如下代码所示：

// allocate zero-copy memory at the host side
cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped);
cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped);
// initialize data at the host side
initialData(h_A, nElem);
initialData(h_B, nElem);
// invoke the kernel with zero-copy memory
sumArraysZeroCopy<<<grid, block>>>(h_A, h_B, d_C, nElem);

六、统一内存

cudaMallocManaged 可分配统一内存（UM），这种内存可以在主机和设备之间自动迁移

基于按需页面迁移的机制。当GPU需访问统一内存时，若数据不在GPU内存中，会触发页面迁移。简化了内存管理，因为无需手动管理主机和设备之间的数据传输

统一内存的性能可能受到页面迁移开销的影响。若主机和设备频繁地对同一块内存进行访问，可能会导致"抖动"现象，降低性能

异步内存存取

分配统一内存 (UM) 时，内存尚未驻留在主机或设备上。主机或设备尝试访问内存时会发生页错误，此时主机或设备会批量迁移所需的数据。能够执行页错误并按需迁移内存对于加速应用程序简化开发流程大有助益。在处理展示稀疏访问模式的数据时（如：在应用程序实际运行之前无法得知需要处理的数据时），以及数据可能由多个 GPU 设备访问时，按需迁移内存将会带来显著优势

有些情况下（如：在运行时之前需要得知数据，以及需要大量连续的内存块时），可以有效规避页错误和按需数据迁移所产生的开销

通过异步内存存取，可以在应用程序代码使用统一内存 (UM) 前，在后台将其异步迁移至系统中的任何 CPU 或 GPU 设备。减少页错误和按需数据迁移所带来的成本，并进而提高 GPU 核函数和 CPU 函数的性能。预取往往会以更大的数据块来迁移数据，因此其迁移次数要低于按需迁移。此技术非常适用于以下情况：在运行时之前已知数据访问需求且数据访问并未采用稀疏模式

使用cudaMemPrefetchAsync函数将数据预取到当前处于活动状态的 GPU 设备，再预取到 CPU

int deviceId;
cudaGetDevice(&deviceId);                                         // The ID of the currently active GPU device
 
cudaMemPrefetchAsync(pointerToSomeUMData, size, deviceId);        // Prefetch to GPU device
cudaMemPrefetchAsync(pointerToSomeUMData, size, cudaCpuDeviceId); // Prefetch to host