一、CUDA内存模型概述
寄存器、共享内存、本地内存、常量内存、纹理内存和全局内存
一个核函数中的线程都有自己私有的本地内存。
一个线程块有自己的共享内存，对同一个线程块中所有的线程都可见，其内容持续线程块的整个生命周期。
所有线程都可以访问全局内存。
所有线程都可以访问的只读空间有：常量内存空间和纹理内存空间。其中纹理内存为各种数据布局提供了不同的寻址模式和滤波模式。
对于一个应用程序来说，全局内存、常量内存和纹理内存中的内容具有相同的生命周期。
1. 寄存器
核函数中声明的一个没有其他修饰符的自变量，通常存储在寄存器中。在核函数声明的数组中，如果用于引用该数组的索引是常量且能在编译时确定，那么该数组也存储在寄存器中。
寄存器是一个在SM中由活跃线程束划分出的较少资源。在核函数中使用较少的寄存器将使在SM上有更多的常驻线程块。每个SM上并发线程块越多，使用率和性能就越高。
如果一个核函数使用了超过硬件限制数量的寄存器，则会用本地内存代替多占用的寄存器。这种寄存器溢出会给性能带来不利影响。
2. 本地内存
核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中。
溢出到本地内存中的变量本质上与全局内存在同一块存储区域，因此本地内存访问的特点是高延迟核低带宽。
3. 共享内存
在核函数中使用修饰符__shared__修饰的变量存放在共享内存中。
因为共享内存是片上内存，所以与本地内存或全局内存相比，其具有更高的带宽和更低的延迟。其使用类似于CPU一级缓存，是可编程的。
每一个SM都有一定数量的由线程块分配的共享内存。
共享内存在核函数的范围内声明，其生命周期伴随着整个线程块。当一个线程块执行结束后，其分配的共享内存将被释放并重新分配给其他线程块。
共享内存是线程间互相通信的基本方式，一个块内的线程通过使用共享内存中的数据可以相互合作。访问共享内存必须同步使用如下调用：void __syncthreads();
4. 常量内存
__constant__
常量内存驻留在设备内存中，并在每个SM专用的常量中缓存。
常量变量必须在全局空间内和所有核函数之外进行声明。
在大多数情况下这个函数是同步的。
线程束中所有线程从相同的内存地址中读取数据时，常量内存表现最好。数学公式中的系数是可以使用常量内存的。
5. 纹理内存
6. 全局内存
__device__
在主机端使用cudaMalloc函数分配全局内存，使用cudaFree函数释放全局内存。
7. GPU缓存
8. CUDA变量声明总结
9. 静态全局内存
二、内存管理
1. 内存分配和释放
在主机上使用cudaMalloc()函数分配全局内存
从主机上传输数据填充所分配的全局内存：cudaMemset()
一旦一个应用程序不再使用已分配的全局内存，可以释放该内存空间：cudaFree()
设备内存的分配和释放操作成本较高，所以应用程序应重利用设备内存，以减少对整体性能的影响。
2. 内存传输
cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind);
CUDA编程的一个基本原则应是尽可能地减少主机和设备之间的传输。
3. 固定内存
GPU不能在可分页主机内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数据时，它无法控制。当从可分页主机内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的或固定的主机内存，将树脂基源数据复制到固定内存中，然后从固定内存传输数据给设备内存。
CUDA运行时允许直接分配主机内存：
cudaMallocHost(void **devPtr, size_t count);
固定主机内存必须通过下述指令来释放：
cudaFreeHost(void *ptr);
与可分页内存相比，固定内存的分配和释放成本更高，但是它为大规模数据传输提供了更高的传输吞吐量。
4. 零拷贝内存
零拷贝内存：主机和设备都可以访问零拷贝内存。
零拷贝内存是固定内存，该内存映射到设备地址空间中。可以通过下列函数创建一个到固定内存的映射：
cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);
释放：cudaFreeHost()
flags:
（1）cudaHostAllocDefault；（2）cudaHostAllocPortable；（3）cudaHostAllocWriteCombined；（4）cudaHostAllocMapped文章来源：https://www.toymoban.com/news/detail-648401.html
cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);文章来源地址https://www.toymoban.com/news/detail-648401.html
5. 统一虚拟寻址（UVA）
到了这里，关于笔记04：全局内存的文章就介绍完了。如果您还想了解更多内容，请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章，希望大家以后多多支持TOY模板网！
![[开发语言][c++]：Static关键字和全局变量](https://imgs.yssmx.com/Uploads/2024/01/809372-1.png)
