找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 146|回复: 3

DAY29:阅读最大化存储器的吞吐率

[复制链接]
发表于 2018-6-9 13:27:20 | 显示全部楼层 |阅读模式
5.3. Maximize Memory Throughput
The first step in maximizing overall memory throughput for the application is to minimize data transfers with low bandwidth.
That means minimizing data transfers between the host and the device, as detailed in Data Transfer between Host and Device, since these have much lower bandwidth than data transfers between global memory and the device.
That also means minimizing data transfers between global memory and the device by maximizing use of on-chip memory: shared memory and caches (i.e., L1 cache and L2 cache available on devices of compute capability 2.x and higher, texture cache and constant cache available on all devices).
Shared memory is equivalent to a user-managed cache: The application explicitly allocates and accesses it. As illustrated in CUDA C Runtime, a typical programming pattern is to stage data coming from device memory into shared memory; in other words, to have each thread of a block:
  • Load data from device memory to shared memory,
  • Synchronize with all the other threads of the block so that each thread can safely read shared memory locations that were populated by different threads,
  • Process the data in shared memory,
  • Synchronize again if necessary to make sure that shared memory has been updated with the results,
  • Write the results back to device memory.
For some applications (e.g., for which global memory access patterns are data-dependent), a traditional hardware-managed cache is more appropriate to exploit data locality. As mentioned in Compute Capability 3.x and Compute Capability 7.x, for devices of compute capability 3.x and 7.x, the same on-chip memory is used for both L1 and shared memory, and how much of it is dedicated to L1 versus shared memory is configurable for each kernel call.
The throughput of memory accesses by a kernel can vary by an order of magnitude depending on access pattern for each type of memory. The next step in maximizing memory throughput is therefore to organize memory accesses as optimally as possible based on the optimal memory access patterns described in Device Memory Accesses. This optimization is especially important for global memory accesses as global memory bandwidth is low, so non-optimal global memory accesses have a higher impact on performance.

回复

使用道具 举报

 楼主| 发表于 2018-6-9 13:52:27 | 显示全部楼层
这章节主要来说如何优化存储器的吞吐率,但说的很简略。首先说,对于一张卡来说, 它主要的大容量存储器有两种:一种是做在卡的PCB板上的显存颗粒,这种存储器是距离GPU较近, 带宽较高的。另外一种则是内存或者来自其他显卡的存储器(P2P Access),这种距离GPU较远,带宽较低。因此历史和多种实际原因,我们常说的global memory实际上不是显存。它在CUDA中的精确含义来说,可能包括:显存,映射的内存(或者unified memory下的自动管理的内存/显存一体自动迁移+缓冲),以及,对方卡的显存。

除了自己的显存速度较快外, 其他的都较慢(哪怕是DGX这种有NVLink的,虽然访问对方的卡比PCI-E版本的同样的卡快很多,但依然不能和本地的自己的显存比),所以基于这种速度上的考虑,在实际的使用中,应当尽量使用快速的,而不应当使用慢速的存储器。

在使用中,这种尽量的使用。包含两种含义:一种是kernel在访存,程序中的一行行代码在一点点的读写(例如,一个线程读取1个INT之类的)。另外一种往往是程序员要求一次性移动一个大范围的数据(例如要求进行一次cudaMemcpy,从内存往显存移动4GB数据)。

前者我们应当要注意,这种一点点的kernel读写应当尽量使用本地存储器(自己的显存分配出来的global memory), 而尽量要减少kernel读写内存映射成的global memory, 或者对方的其他卡上的显存映射成的global memory.(对方的卡上的显存是指P2P Access的,一张卡可以访问其他卡上的显存,就如同是自己的显存一样。只是速度慢)


而后者我们应当尽量避免显存和内存间的cudaMemcpy*(), 如果一个问题可以使用自己的显存上的数据整理移动操作完成(Device到Device的内部传输),那就不要从内存传输过来(PCI-E较慢,自己显存内部的传输较快);但如果一个问题必须需要从内存读取(例如刚刚CPU端从磁盘上读取了一个文件),那么应当考虑尽量减少传输次数,每次传输较大的内容(例如可以考虑10张图片传输一次),因为CUDA的传输,小字节的传输量没有优势,一次传输较多的内容,性能才能提速上来。还有就是在不能避免这种传输的时候,不仅仅要一次传输较大的内容,还应当考虑使用page-locked memory(用cudaMallocHost(), cudaHostAlloc(), 或者cudaHostRegister()得到的内存),这种内存传输起来较快(少了一步内存中转挪移的过程,这个后面会遇到的)

这就是为何手册说:
That means minimizing data transfers between the host and the device, as detailed in Data Transfer between Host and Device, since these have much lower bandwidth than data transfers between global memory and the device.



然后除了这两种常见的(嗯嗯。收到)大容量存储外,也就是除了内存和显存外(包括其他伙伴卡的显存),另外一种存储器叫片内存储器(或者片上),这种存储器没有独立的存储芯片,而是集成在GPU核心芯片里(显存会在显卡的基板PCB上看到的)









回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-6-9 13:57:33 | 显示全部楼层
这是一张拆开的1080:
QQ图片20180609150127.jpg

这些是显存的。
QQ图片20180609150137.jpg
中间的那个大芯片是GPU(核心)。每个显存芯片上面标记有容量。这些显存累加起来是8GB(对于1080来说是8GB。其他卡可能不是8GB)

你会看到GPU(中间的大芯片)能直接通过自己在PCB版上的走线,访问它们,而下面的金手指是PCI-E,访问内存或者其他卡上的显存,需要走PCI-E。
所以这是为何刚才说,尽量使用自己的显存(速度较快)的原因。然后除了这两种大容量存储器外---也就是你看到的PCB上的显存,和走PCI-E访问的内存和其他卡的显存。还有一种存储器是直接看不到的。它在这里面:
QQ图片20180609150145.jpg
如图。

GPU核心芯片自己还带有一些On Chip的存储器。这种存储器往往叫片内存储器,或者片上存储器。
看图可以看到,这种存储器和GPU核心的连接更加紧密(就在GPU核心里么),而这种On-Chip的存储器带宽非常高,延迟非常低(相比其他片外的,甚至需要跨PCI-E的),但可惜容量也非常小。一般包括我们常说的L1/L2 Cache, shared memroy等等它们,这些小而高速的片上的存储器,是提高性能的关键。例如很多人还在翻阅老书,非要死命的使用texture的原因,

就是因为texture cache这种,就是这些片上的存储器中的一种。优化的使用它的确是性能提升的关键。

回到GPU核心芯片内部,里面的这些缓存和shared memory之类的还能继续分,GPU上的存储器依然可以分成两种:SM内部的,包括L1 cache, L1 texture cache, Shared memory等等。和SM外部的,这主要是L2 cache。因为shared memory这种是SM内部的,blocks在上到SM后,距离更近,延迟更低,带宽更高(每个SM上的shared memory带宽 × 有多少个SM),所以应当尽量考虑利用SM内部的。

这就像俄罗斯套娃一样。这也很好理解吧?整个GPU服务器--->重要的是使用显卡的自身存储器--->重要的是使用显卡上的GPU核心里面的存储器--->重要的是使用SM里面的存储器。(好了。没有再下一层了。不用担心,套娃也有层数限制的)

所以本章节还简单说了一下如何使用shared memory, 例如从显存载入内容到shared memory,然后局部同步,再在shared memory内部多次使用这内容

(狂用即可。比显存快的多),最后如果需要,可以将最终结果再回写显存,而显存的结果如果需要,可以最终回写内存,而内存的结果如何需要,最后可以保存到磁盘上或者通过网络传输走。


本章节还表明(再往下),有的应用更使用使用普通缓存,而不是shared memory,例如一些数据访问的模型/模式是和具体的数据有关的,无法提前安排到shared memory,此时应当更多的考虑使用普通cache,例如SM内部的L1 cache。需要注意这里,不同的计算能力上的L1 cache安排不同,有独立的,和shared memory合并在一起的,也有和texture cache合并在一起的(例如昨天的GM204和一些Pascal卡是和texture cache合并在一起的,并且有导致occupancy会变成0的问题,此时昨日说了将自动禁用它)。



回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-6-9 14:12:06 | 显示全部楼层
sisiy 发表于 2018-6-9 13:57
这是一张拆开的1080。

这些是显存的。中间的那个大芯片是GPU(核心)。每个显存芯片上面标记有容量。这 ...

本章节还说了,显存的访存模式也很重要。以及,还说了不恰当的访存模式会严重降低性能,甚至降低一个数量级。例如shared memory的bank conflict,或者例如global memory上的严重不合并 等


回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表