GPU 和CPU通过 VLink  或者PCLe 相连。

每个SM 又含有多个cuda core，多个SM共享全局内存，通过L2 高速缓存和全局内存进行相连，不同代的GPU之间的体系结构有所不同。

右边是一个SM 上的 内存，有共享内存，局部内存，寄存器，可以访问全局内存、常数内存、纹理内存。

每个线程 有自己的私有本地内存（local memory）、每个线程块包含共享内存，可以被线程块中的所有线程共享，其生命周期与线程块一致。共享内存的访问速度非常块，比全局内存的访问要快100倍。

所有的线程都可以访问全局内存，量大，但是访问速度是最慢的

常量和纹理内存用的比较少，所有的线程都能对它们进行读操作，不能进行写操作。

多个MP共享L2

 

存储量越大  访问速度越慢。CPU/GPU 的内存结构中，一级和二级缓存都是不可编程的，对用户是不开放的（完全不可控制的存储设备）。比如 L1 L2 缓存，寄存器，是不可编程的，是有cpu 、GPU进行管理控制的。

全局内存生命周期是从开始申请到释放，如果程序结束了，但是没有对释放内存的操作，全局内存的生命周期还是没有结束、

全局内存生命周期（显存）：

当一个线程开始执行，它会拥有自己的本地内存、寄存器，当线程结束时，对应的内存也就结束了。

线程块有自己的共享内存，对线程块内的所有线程可见，当线程块执行结束时，这些共享内存的生命周期就结束了。

寄存器的生命周期：

寄存器位于SM  中，访问的速度非常快，比全局内存的访问速度要快两个数量级，寄存器是稀缺资源，程序在执行过程中，会把一些最常用的变量放在寄存器中，每个线程中寄存器的数量是一定的，如果线程的变量太多，溢出的话，必须把一部分的数据放在全局内存中，效率就会变得很低。尽量少使用寄存器，这样可以使处于活跃状态的线程数更多。

核函数中，符合存储在寄存器，但不能进入被核函数分配的寄存器空间中的变量将存储在本地内存中，编译器中可能放在本地内存中的变量有以下几种：

使用未知索引的引用的本地数组

可能会占用大量寄存器空间的较大本地数组或者结构体

不满足核函数寄存器限定条件的变量

本地内存实质上和全局内存一样在同一块存储区域中，高延迟低带宽。

本地内存存储在每个SM 的一级缓存或者是设备的二级缓存。

 

共享内存

延迟低，带宽高。由线程块分配的共享内存，共享内存是片上内存，可以被编程。过度使用共享内存会导致SM 上活跃的线程束减少。也就是说，一个线程块使用的共享内存过多，导致更多的线程块没有办法被SM启动，这样活跃的线程束数量减少。

共享内存在核函数内声明，生命周期和线程块一致，线程块被运行，共享内存被分配，当此块结束，则共享内存被释放。共享内存是块内线程可见的，存在竞争问题，也可以通过共享内存进行通信。为了避免竞争，使用同步语句：_syncthreads()

频繁使用这样的语句会影响程序执行效率。

为了实现内存高带宽的同时访问，shared memory  被划分成了同时访问的等大小内存块。内存读写n 个地址的行为可以以b个独立的bank同时操作的方式进行，这样可以有效提高到b倍。如果多个线程请求的内存地址被映射到同一个bank上，那么这些请求就变成了窜行的。硬件将把这些请求分成x个没有冲突的请求序列，带宽就降为原来的x分之一。但是如果一个wrap内所有线程都访问一个内存地址的话，就会产生一次广播（boardcast），这些请求会一次完成。

全局内存访问对齐，一次要读取指定大小（32,64,128）整数倍字节的内存。内存请求的实物越多，未使用字节被传输的可能性越大。对齐的读写模式使得不需要的数据也被传输。

一级缓存

二级缓存

只读常量缓存

只读纹理缓存

CPU 读写过程都可能被缓存，GPU 只有读的过程被缓存，写的过程没有缓存。

常量缓存和纹理缓存，它们被用于设备内存来提高各自内存空间的读取性能。

GPU内存管理    申请 拷贝 释放

cudaMalloc((void **) & ptr, size)

cudaMemcpy(dst,src,size,directory)

cudathreadsynchronize()

cudaFree()

cudaMallocHost()//申请锁页内存