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

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第29天,我们正在讲解性能,希望在接下来的71天里,您可以学习到原汁原味的CUDA,同时能养成英文阅读的习惯。

本文共计363字,阅读时间15分钟

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.

本文备注/经验分享:

这章节主要来说如何优化存储器的吞吐率,但说的很简略。首先说,对于一张卡来说, 它主要的大容量存储器有两种:一种是做在卡的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上看到的)

这是一张拆开的1080:

哪里是显存呢?

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

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

如图。 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的问题,此时昨日说了将自动禁用它)。

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

有不明白的地方,请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-06-11

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏加米谷大数据

HDFS如何读取文件以及写入文件

成都加米谷大数据科技有限公司是一家专注于大数据人才培养的机构。由来自阿里、华为、京东、星环等国内知名企业的多位技术大牛联合创办,技术底蕴丰厚,勤奋创新,精通主流...

833
来自专栏大大的微笑

Master-work模式

①.master负责手机客户端的请求,将任务分发给work,回收work处理结果, ②.work做具体的业务逻辑,并将结果通知到master 这里做了一个给每个...

2156
来自专栏程序人生

谈谈调度 - Linux O(1)

约莫十五年前,当我刚刚开始参加工作时,赶上 Linux 发布划时代的 2.6 内核。在这个大家都翘首期盼的内核版本中,最令人兴奋的便是 O(1) schedul...

2628
来自专栏吉浦迅科技

DAY52:阅读scheduling

今天的内容比较特殊,因为这个部分并没有出现在NVIDIA 在线版的《CUDA C Programming Guide》,但是如果你下载了CUDA,里面会带一份P...

401
来自专栏编程一生

架构师之路--搜索业务和技术介绍及容错机制

862
来自专栏吉浦迅科技

NVIDIA正式宣布CUDA 6:支持统一寻址

NVIDIA今天正式宣布了最新版并行计算开发工具CUDA 6,相比此前的CUDA 5.5有着革命性的巨大进步。 NVIDIA表示,CUDA 6可以让并行编程前所...

2808
来自专栏牛客网

2019届网易JAVA实习2面经历

1865
来自专栏腾讯云数据库(TencentDB)

TDSQL“相似查询工具MSQL+”入选VLDB论文

作者介绍:王晓宇,腾讯数据库TDSQL团队成员,目前参与TDSQL数据库内核研发工作。

1681
来自专栏欧阳大哥的轮子

深入iOS系统底层之指令集介绍

说到指令集以及CPU架构体系,大家就会想到计算机专业课程里面的计算机体系结构的方面的内容。既然课程中已经有了的内容我就不想那么枯燥的去复述一遍,而是先看一个类的...

491
来自专栏数据和云

极速体验:Oracle 18c 下载和Scalable Sequence新特性

Oracle 18c 已至,目前已经可以从Oracle Edelivery 网站下载。 该网站的网址是:https://edelivery.oracle.com...

2513

扫码关注云+社区