(转)GPU的存储结构

最新推荐文章于 2024-07-26 10:42:33 发布

qq_22677425

最新推荐文章于 2024-07-26 10:42:33 发布

阅读量2.7k

点赞数 5

学习了一下GPU的存储结构，有一篇很好的文章，收藏一下
【CUDA】学习记录（7）- Global Memory

Memory
kernel性能高低不仅仅和线程的执行方式相关，还和存储器的访问和管理密切相关。众所周知，memory的操作在讲求效率的语言中占有极重的地位。low-latency和high-bandwidth是高性能的理想情况。但是购买拥有大容量，高性能的memory是不现实的，或者不经济的。因此，我们就要尽量依靠软件层面来获取最优latency和bandwidth。CUDA将memory model unit分为device和host两个系统，充分暴露了其内存结构以供我们操作，给予用户充足的使用灵活性。

GPU和CPU的主存都是用DRAM实现，cache则是用lower-latency的SRAM来实现。GPU和CPU的存储结构基本一样。但是CUDA将memory结构更好的呈现给用户，从而能更灵活的控制程序行为。

CUDA Memory Model
对于编程人员来讲，memory分为两类：
➤ Programmable: 我们可以灵活操作的部分。
➤ Non-programmable: 不能控制的部分。
对CPU而言，L1和L2缓存对我们而言是non-programmable memory.
CUDA将memory完全暴露给了用户：
➤ Registers
➤ Shared memory
➤ Local memory
➤ Constant memory
➤ Texture memory
➤ Global memory

在这里插入图片描述

每个thread有自己独立的registers和local memory，每个block中的所有threads共享share memory，所有的线程都可以访问global memory，其中constant和texture是只读内存。

Registers
寄存器是GPU最快的memory，kernel中没有什么特殊声明的自动变量都是放在寄存器中的。当数组的索引是constant类型且在编译期能被确定的话，就是内置类型，数组也是放在寄存器中。
寄存器是每个thread的私有变量，一旦thread执行结束，寄存器变量就会失效。寄存器也是稀缺资源，在Fermi每个thread最多63个registers, Kepler(255).在每个kernel中使用较少的寄存器，可以使更多的block 驻留在SM上，实现更多并发的blocks，进而提高occupy和性能。
如果kernel使用的register超过硬件限制，这部分会使用local memory来代替register，即所谓的register spilling，我们应该尽量避免这种情况。编译器有相应策略来最小化register的使用并且避免register spilling。
-Xptxas -v,-abi=no选项可以查看每个thread使用的寄存器数量，shared memory和constant memory的大小。

Local Memory
有时候，如果register不够用了，那么就会使用local memory来代替这部分寄存器空间。除此外，下面几种情况，编译器可能会把变量放置在local memory：
➤ 编译期间无法确定值的本地数组
➤消耗太多寄存器的较大的结构体或数组
➤ 任何超过寄存器限制的变量
local memory这个名字是有歧义的：在local memory中的变量本质上跟global memory在同一块存储区。所以，local memory有很高的latency和较低的bandwidth。在CC2.0以上，GPU针对local memory会有L1（per-SM）和L2（per-device）两级cache。

Shared Memory
用shared修饰符修饰的变量存放在shared memory。因为shared memory是on-chip的，他相比localMemory和global memory来说，拥有高的多bandwidth和低很多的latency。他的使用和CPU的L1cache非常类似，但是他是programmable的。
按惯例，像这类性能这么好的memory都是有限制的，shared memory是以block为单位分配的，如果每个block占用的share memory过多，那么每个SM上驻留的blocks就少，active warp的数目也会减少。
不同于register，shared memory尽管在kernel里声明的，但是他的生命周期是伴随整个block，而不是单个thread。当该block执行完毕，所拥有的资源就会被释放，重新分配给别的block。
shared memory是同一个block 中thread交流的基本方式。同一个block中的thread通过shared memory中的数据来相互合作。获取shared memory的数据前必须先用__syncthreads()同步。L1 cache和shared memory使用相同的64KB on-chip memory，我们也可以使用API来动态配置二者的大小。

Constant memory GPU编程自学7 —— 常量内存与事件
常量内存具有以下特点：

需要由 constant 限定符来声明
只读
硬件上并没有特殊的常量内存块，常量内存只是只是全局内存的一种虚拟地址形式
目前的GPU常量内存大小都只有64K，主要是因为常量内存采用了更快速的16位地址寻址(2^16 = 65536 = 64K)
对于数据不太集中或者重用率不高的内存访问，尽量不要使用常量内存，否则甚至会慢于使用全局内存
常量内存无需cudaMalloc()来开辟，而是在声明时直接提交一个固定大小，比如 constant float mdata[1000]
常量内存的赋值不能再用cudaMemcpy()，而是使用cudaMemcpyToSymbol()
常量内存带来性能提升的原因主要有两个：

对常量内存的单次读操作可以广播到其他的“邻近(nearby)”线程，这将节约15次读取操作
常量内存的数据将缓存起来，因此对于相同地址的连续操作将不会产生额外的内存通信量。

对于原因1，涉及到线程束(Warp)的概念。

在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。即线程束中的每个线程都将在不同数据上执行相同的指令。

当处理常量内存时，NVIDIA硬件将把单次内存读取操作广播到每个半线程束(Half-Warp)。在半线程束中包含16个线程，即线程束中线程数量的一半。如果在半线程束中的每个线程从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求并在随后将数据广播到每个线程。如果从常量内存中读取大量数据，那么这种方式产生的内存流量只是使用全局内存时的1/16。

对于原因2，涉及到缓存的管理

由于常量内存的内容是不发生变化的，因此硬件将主动把这个常量数据缓存在GPU上。在第一次从常量内存的某个地址上读取后，当其他半线程束请求同一个地址时，那么将命中缓存，这同样减少了额外的内存流量。

另一方面, 常量内存的使用也可能会对性能产生负面的影响。半线程束广播功能实际上是一把双刃剑。虽然当所有16个线程都读取相同地址时，这个功能可以极大提升性能，但当所有16个线程分别读取不同的地址时，它实际上会降低性能。因为这16次不同的读取操作会被串行化，从而需要16倍的时间来发出请求。但如果从全局内存中读取，那么这些请求会同时发出。

Texture Memory
texture Memory实际上也是global Memory在一块，但是他有自己专有的只读cache。这个cache在浮点运算很有用，texture Memory是针对2D空间局部性的优化策略，所以thread要获取2D数据就可以使用texture Memory来达到很高的性能。

Global Memory
global Memory是空间最大，latency最高，GPU最基础的memory。“global”指明了其生命周期。任意SM都可以在整个程序的生命期中获取其状态。global中的变量既可以是静态也可以是动态声明。可以使用device修饰符来限定其属性。global memory的分配就是之前频繁使用的cudaMalloc，释放使用cudaFree。global memory驻留在devicememory，可以通过32-byte、64-byte或者128-byte三种格式传输。这些memory transaction必须是对齐的，也就是说首地址必须是32、64或者128的倍数。优化memory transaction对于性能提升至关重要。当warp执行memory load/store时，需要的transaction数量依赖于下面两个因素：
➤ Distribution of memory addresses across the threads of that warp.（没有读明白？？？）
➤ Alignment of memory addresses per transaction.地址对齐
一般来说，所需求的transaction越多，潜在的不必要数据传输就越多，从而导致throughput efficiency降低。
对于一个既定的warp memory请求，transaction的数量和throughput efficiency是由CC版本决定的。对于CC1.0和1.1来说，对于global memory的获取是非常严格的。而1.1以上，由于cache的存在，获取要轻松的多。

GPU Caches
跟CPU的cache一样，GPU cache也是non-programmable的。在GPU上包含以下几种cache：
➤ L1
➤ L2
➤ Read-only constant
➤ Read-only texture
每个SM都有一个L1 cache，所有SM共享一个L2 cache。二者都是用来缓存local和global memory的，当然也包括register spilling的那部分。在Fermi GPus 和 Kepler K40或者之后的GPU，CUDA允许我们配置读操作的数据是否使用L1和L2或者只使用L2。
在CPU方面，memory的load/store都可以被cache。但是在GPU上，只有load操作会被cache，store则不会。
每个SM都有一个只读constant cache和texture cache来提升性能

wiki/CUDA参数