CUDA 内存概述

GPU的内存包括：

全局内存（global memory）
常量内存（constant memory）
纹理内存核表面内存（texture memory）
寄存器（register）
局部内存（local memory）
共享内存（shared memory）
L1、L2缓存（从费米架构开始有了SM层次的 L1 cache 和设备层次的 L2 cache）

在这里插入图片描述
速度快慢如下图所示：

CUDA 内存详解

Global Memory

Global Memory在某种意义上等同于GPU显存，kernel函数通过Global Memory来读写显存。Global Memory是kernel函数输入数据和写入结果的唯一来源。

Rigisters 寄存器

寄存器是GPU最快的memory，kernel中没有什么特殊声明的自动变量都是放在寄存器中的。当数组的索引是constant类型且在编译期能被确定的话，就是内置类型，数组也是放在寄存器中。

寄存器变量是每个线程私有的，一旦thread执行结束，寄存器变量就会失效
寄存器是稀有资源。(省着点用，能让更多的block驻留在SM中，增加Occupancy)
--maxrregcount 可以设置大小
不同设备架构，数量不同

Shared Memory

Shared Memory位于GPU芯片上，访问延迟仅次于寄存器。Shared Memory是可以被一个Block中的所有Thread来进行访问的，可以实现Block内的线程间的低开销通信。在SMX中，L1 Cache跟Shared Memory是共享一个64KB的告诉存储单元的，他们之间的大小划分不同的GPU结构不太一样；
用 __shared__ 修饰符修饰的变量存放在shared memory:

On-chip
拥有高的多bandwidth和低很多的latencyo
同一个Block中的线程共享一块Shared Memoryo
需要使用 __syncthreads() 同步。
比较小，要节省着使用，不然会限制活动warp的数量

Local Memory

Local Memory本身在硬件中没有特定的存储单元，而是从Global Memory 虚拟出来的地址空间。** Local Memory 是为寄存器无法满足存储需求的情况而设计的，主要是用于存放单线程的大型数组和变量。** Local Memory是线程私有的，线程之间是不可见的。由于GPU硬件单位没有Local Memory的存储单元，所以，针对它的访问是比较慢的。

但是更多在以下情况，会使用 Local Memory：

无法确定其索引是否为常量的数组。
会消耗太多寄存器空间的大型结构或数组。
如果内核使用了多于可用寄存器的任何变量(这也称为寄存器溢出)
--ptxas-options=-v

Constant Memory

Constant Memory （常量内存）类似于 Local Memory，也是没有特定的存储单元的，只是Global Memory 的虚拟地址。因为它是只读的，所以简化了缓存管理，硬件无需管理复杂的回写策略。Constant Memory 启动的条件是同一个warp所有的线程同时访问同样的常量数据。

其具有以下几个特点：

constant的范围是全局的，针对所有kernel。
在同一个编译单元，constant对所有kernel可见。
kernel只能从constant Memory读取数据，因此其初始化必须在host端使用下面的函数调用: cudaError_t cudaMemcpyToSymbollconst void* symbol, const void* src,size t count);
当一个warp中所有thread都从同一个Memory地址读取数据时，constant Memory表现会非常好会触发广播机制。

Texture Memory

Texture Memory是GPU的重要特性之一，也是GPU编程优化的关键。Texture Memory实际上也是Global Memory的一部分，但是它有自己专用的只读cache。这个cache在浮点运算很有用，Texture Memory是针对2D空间局部性的优化策略，所以thread要获取2D数据就可以使用texture Memory来达到很高的性能。从读取性能的角度跟Constant Memory类似。

在这里插入图片描述

Host Memory

主机端存储器主要是内存可以分为两类：可分页内存（Pageable）和页面（Page-Locked 或 Pinned）内存。

可分页内存通过操作系统 API(malloc/free) 分配存储器空间，该内存是可以换页的，即内存页可以被置换到磁盘中。可分页内存是不可用使用DMA（Direct Memory Acess)来进行访问的，普通的C程序使用的内存就是这个内存

例子

下面例子讲解如何使用统一内存：

__device__ __managed__ int x[2];
__device__ __managed__ int y;
__global__ void kernel(){
    x[1] = x[0] + y;
}

int main(){
    x[0] = 3;
    y = 5;
    kernel<<< 1, 1 >>>();
    cudaDeviceSynchronize();
    printf("result=%d\n", x[1]);
    return 0;
}