找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 81|回复: 2

DAY83:阅读Compute Capability 3.x

[复制链接]
发表于 2018-10-11 12:03:47 | 显示全部楼层 |阅读模式
H.3. Compute Capability 3.xH.3.1. ArchitectureA multiprocessor consists of:
  • 192 CUDA cores for arithmetic operations (see Arithmetic Instructions for throughputs of arithmetic operations),
  • 32 special function units for single-precision floating-point transcendental functions,
  • 4 warp schedulers.

When a multiprocessor is given warps to execute, it first distributes them among the four schedulers. Then, at every instruction issue time, each scheduler issues two independent instructions for one of its assigned warps that is ready to execute, if any.
A multiprocessor has a read-only constant cache that is shared by all functional units and speeds up reads from the constant memory space, which resides in device memory.
There is an L1 cache for each multiprocessor and an L2 cache shared by all multiprocessors. The L1 cache is used to cache accesses to local memory, including temporary register spills. The L2 cache is used to cache accesses to local and global memory. The cache behavior (e.g., whether reads are cached in both L1 and L2 or in L2 only) can be partially configured on a per-access basis using modifiers to the load or store instruction. Some devices of compute capability 3.5 and devices of compute capability 3.7 allow opt-in to caching of global memory in both L1 and L2 via compiler options.
The same on-chip memory is used for both L1 and shared memory: It can be configured as 48 KB of shared memory and 16 KB of L1 cache or as 16 KB of shared memory and 48 KB of L1 cache or as 32 KB of shared memory and 32 KB of L1 cache, using cudaFuncSetCacheConfig()/cuFuncSetCacheConfig():
  1. // Device code
  2. __global__ void MyKernel()
  3. {
  4.     ...
  5. }

  6. // Host code

  7. // Runtime API
  8. // cudaFuncCachePreferShared: shared memory is 48 KB
  9. // cudaFuncCachePreferEqual: shared memory is 32 KB
  10. // cudaFuncCachePreferL1: shared memory is 16 KB
  11. // cudaFuncCachePreferNone: no preference
  12. cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferShared)
复制代码



The default cache configuration is "prefer none," meaning "no preference." If a kernel is configured to have no preference, then it will default to the preference of the current thread/context, which is set using cudaDeviceSetCacheConfig()/cuCtxSetCacheConfig() (see the reference manual for details). If the current thread/context also has no preference (which is again the default setting), then whichever cache configuration was most recently used for any kernel will be the one that is used, unless a different cache configuration is required to launch the kernel (e.g., due to shared memory requirements). The initial configuration is 48 KB of shared memory and 16 KB of L1 cache.
Note: Devices of compute capability 3.7 add an additional 64 KB of shared memory to each of the above configurations, yielding 112 KB, 96 KB, and 80 KB shared memory per multiprocessor, respectively. However, the maximum shared memory per thread block remains 48 KB.
Applications may query the L2 cache size by checking the l2CacheSize device property (see Device Enumeration). The maximum L2 cache size is 1.5 MB.
Each multiprocessor has a read-only data cache of 48 KB to speed up reads from device memory. It accesses this cache either directly (for devices of compute capability 3.5 or 3.7), or via a texture unit that implements the various addressing modes and data filtering mentioned in Texture and Surface Memory. When accessed via the texture unit, the read-only data cache is also referred to as texture cache.


H.3.2. Global Memory
Global memory accesses for devices of compute capability 3.x are cached in L2 and for devices of compute capability 3.5 or 3.7, may also be cached in the read-only data cache described in the previous section; they are normally not cached in L1. Some devices of compute capability 3.5 and devices of compute capability 3.7 allow opt-in to caching of global memory accesses in L1 via the -Xptxas -dlcm=ca option to nvcc.
A cache line is 128 bytes and maps to a 128 byte aligned segment in device memory. Memory accesses that are cached in both L1 and L2 are serviced with 128-byte memory transactions whereas memory accesses that are cached in L2 only are serviced with 32-byte memory transactions. Caching in L2 only can therefore reduce over-fetch, for example, in the case of scattered memory accesses.
If the size of the words accessed by each thread is more than 4 bytes, a memory request by a warp is first split into separate 128-byte memory requests that are issued independently:
  • Two memory requests, one for each half-warp, if the size is 8 bytes,
  • Four memory requests, one for each quarter-warp, if the size is 16 bytes.

Each memory request is then broken down into cache line requests that are issued independently. A cache line request is serviced at the throughput of L1 or L2 cache in case of a cache hit, or at the throughput of device memory, otherwise.
Note that threads can access any words in any order, including the same words.
If a non-atomic instruction executed by a warp writes to the same location in global memory for more than one of the threads of the warp, only one thread performs a write and which thread does it is undefined.
Data that is read-only for the entire lifetime of the kernel can also be cached in the read-only data cache described in the previous section by reading it using the __ldg() function (see Read-Only Data Cache Load Function). When the compiler detects that the read-only condition is satisfied for some data, it will use __ldg() to read it. The compiler might not always be able to detect that the read-only condition is satisfied for some data. Marking pointers used for loading such data with both the const and __restrict__ qualifiers increases the likelihood that the compiler will detect the read-only condition.
Figure 18 shows some examples of global memory accesses and corresponding memory transactions.
Figure 18. Examples of Global Memory Accesses. Examples of Global Memory Accesses by a Warp, 4-Byte Word per Thread, and Associated Memory Transactions for Compute Capabilities 3.x and Beyond


                               
登录/注册后可看大图






H.3.3. Shared Memory
Shared memory has 32 banks with two addressing modes that are described below.
The addressing mode can be queried using cudaDeviceGetSharedMemConfig() and set using cudaDeviceSetSharedMemConfig() (see reference manual for more details). Each bank has a bandwidth of 64 bits per clock cycle.
Figure 19 shows some examples of strided access.
Figure 20 shows some examples of memory read accesses that involve the broadcast mechanism.
64-Bit Mode
Successive 64-bit words map to successive banks.
A shared memory request for a warp does not generate a bank conflict between two threads that access any sub-word within the same 64-bit word (even though the addresses of the two sub-words fall in the same bank): In that case, for read accesses, the 64-bit word is broadcast to the requesting threads and for write accesses, each sub-word is written by only one of the threads (which thread performs the write is undefined).

32-Bit Mode
Successive 32-bit words map to successive banks.
A shared memory request for a warp does not generate a bank conflict between two threads that access any sub-word within the same 32-bit word or within two 32-bit words whose indices i and jare in the same 64-word aligned segment (i.e., a segment whose first index is a multiple of 64) and such that j=i+32 (even though the addresses of the two sub-words fall in the same bank): In that case, for read accesses, the 32-bit words are broadcast to the requesting threads and for write accesses, each sub-word is written by only one of the threads (which thread performs the write is undefined).



回复

使用道具 举报

 楼主| 发表于 2018-10-11 14:40:42 | 显示全部楼层
嗯嗯。今天我们将从Kepler开始,详细的说一下目前还流行的各代显卡的主要差别,和他们的进化过程。

首先说,从这一代起,单精度浮点性能,并无本质变化:老的Kepler的泰坦(计算能力3.X),依然可以达到单精度大约8T的理论性能。

而今天的RTX2080,计算能力7.5, 也依然单精度在8T-10TFlops之间(后者是Boost)。也就是说, 如果只追求单精度性能的用户,在不考虑历代制程的变化和功耗的变化的情况下,实际上从Kepler的3.X开始到现在,基本没有太大的变化。这点用户首先需要注意了。

那么3.X/5.X/6.X/7.X这4代过来,NV更加注重的是新特性的增加(例如6.0开始的半精度),效率的提升(例如从5.0开始的同样制程下的功耗性能比),以及,适应了潮流的新特性(例如6.X开始的__dp4a, 7.X开始的两代Tensor Core等等。)

但是不能否认的是,Kepler是一代开创性的实验品,它所带来的很多特性,直到今天,依然发挥的重要作用,并逐渐成为的CUDA的基本构成部分。

我们简单的从GPU整体,以及SM内部,来看下Kepler都给我们带来了什么。
首先最重要的是,计算能力3.X引入了制程的提升(28nm),以及,引入了软件调度。相比之前的Fermi那一代,降低了大量的功耗。
  

你可能很难想象,在Fermi的当时,只有400多个SP的一张卡(例如GTX480),功耗就可以到200多瓦。当时人们都戏称,GTX480可以煎鸡蛋。犹如当年的Intel的失败的奔腾4的一代一样。

再看看到了计算能力3.X(Fermi是2.X),一张有2300多个SP的780,功耗也才200多W出头。提升的非常明显。这是因为本代的卡,试图从多个方面进行功耗的优化,效果还是显著的,但之所以给人留下了不好的印象,那是因为一些地方没有做的平衡,导致了性能的损失。首先是一个SM,在寄存器资源为256KB(64K个4B寄存器)的情况下,塞入了192个SP!
这是什么概念?下一代的5.X同样的资源,之塞入了128个SP,6.0计算能力更是只有64个SP。也包括今天在京东上热卖的RTX2080(7.5)同样只有64个SP。这导致了执行单元(SP)数量看起来非常惊人的多,但是可用的资源较少。

很多情况下,3.X这一代里面的192个SP,因为资源不饱它们,经常出现性能只有类似128个SP的效果(66.7%性能)。这是这一代的主要被人骂的地方,其他的地方都几乎很好。如果人们不那么贪心,用来和上一代的Fermi比,哪怕每次都是最坏的性能,依然SP数量在同样的功耗下,取得了长足的提升,还是值得考虑的。这是本代卡最显著的特点。其次,本代的卡对CUDA做出了很多基础性的贡献,例如:动态并行从这一代的卡开始引入(计算能力3.5),使得人们不在局限于从CPU上进行任务调度,能够让GPU自己调度自己,写代码的方便性提升了很多。

也使得原本很难进行并行化,或者对并行化感到棘手的应用场合,变得适用了。(详情请参考之前的动态并行章节。这是一个很重要的特性,感谢Kepler为我们引入)


其次,则是这一代的卡引入了Hyper-Q,Hyper-Q这个概念已经成为了现在使用CUDA的基础,就如同刚才的动态并行一样。在这代卡(精确的说,3.5, Kepler2代)之前,硬件队列只有1个,使得很多应用多流的场合,并不能起到效果(请参考之前的多流章节)

于是出现了很多CUDA参考书,引入了各种办法进行修补,例如多流时候的深度发布任务有限,和广度任务发布优先。(也就是先就着一个流发布完其中的任务,然后再对下一个流发布;还是每个流发布一个任务后,立刻切换到下一个流)

当年不同的发布方式,虽然都是在使用多流,导致了很多时候多流的效果消失。需要用户反复试验发布命令的方式。而Kepler引入了多个硬件任务队列,叫Hyper-Q的东西,使得用户的任何一种发布方式,往往都可以取得较好的多流并行效果。从而很大程度的减轻了CUDA用户编程时候的负担。这是一个很大的改进(详情参考我们之前的Hyper-Q章节)

如今这个特性也默默的沉淀在你所用的每一张N卡里面。成为了CUDA的标配。除了这个特性外,这一代还引入了软件调度,软件指令内部含有原本硬件应该负责的调度信息,软件指令内部含有原本硬件应该负责的调度信息,从而节省了很多的硬件晶体管,降低了功耗。

(这一代不明显,但是作为基础在Maxwell(5.X)中发扬了广大,同时一路沿用到6.X和7.0/7.5)

这个基础是从Kepler开始奠定的,
(所以这是为何今天N卡这么省电的原因)

从差不多Kepler的同时时期,竞争对手AMD抛弃了软件调度(它的VLIW)试图提升性能,最终今天变成火炉。

而NV这也从这个时期开始,放弃了自己原本的硬件调度,改用软件调用,降低了功耗变成了冰箱(开玩笑哈)

这两家大致从这个时候开始分道扬镳,AMD变得NV化,NV变得AMD话。不能不说是一个讽刺而又需要面对的事实。但是需要说明的是,为了配置这种改进,尽量的将晶体管分配给执行单元,而不是调度单元(执行单元才能发挥性能,调度单元只是为了调度,如果员工和领导的关系,NV的Kepler如同一个公司,引入了海量的员工,和少量的领导)







回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-10-11 15:14:21 | 显示全部楼层
sisiy 发表于 2018-10-11 14:40
嗯嗯。今天我们将从Kepler开始,详细的说一下目前还流行的各代显卡的主要差别,和他们的进化过程。

首先 ...

实际上本章节中说的,192个SP vs 4调度器 vs 256KB的寄存器,不足以能让这些SP(员工们)充分忙碌,

很多综合性的测试表明,这么多的SP(192)往往只表现出来了128个SP的性能。也就是说,Kepler有可能峰值,但是很困难。用户应当只认为有128/192的性能(很大原因在于寄存器的bank conflict)。

阿里曾经在GTC大会上,展现过他们自己编写的Kepler汇编器,能尽量让用户能发挥性能,实际上,在它们的阿里云的GPU实例中,也曾经提供过该汇编器给他们的云GPU实例的用户使用,并有过详细说明。我们现在的角度来看,Kepler和后续的代的最大的特点是,kepler的调度器需要寄存器配合,尽量能双发(至少需要50%的几率在双发射指令)才有可能峰值。而后续代的显卡则是,单发射指令就足以压满SP的峰值了,双发只是可选。这是一个很大的差异。

然后除了指令调度上的区别(本章节的前部分)和SP数量的暴增。Kepler还做了一些常用的用户可能方便的改进。
例如说,增宽的shared memory(本章的后面),从4B一个bank,扩大到8B。这是的一些特殊的应用场合,例如对double操作的读取能提升一倍。
也能让类似图像处理中的索引颜色,AES加密解密之类的场合进行提速。但是可能是用到的地方不多,在后续的代中取消了这点。

不过本章节还是对此进行了说明。因为如果你手头有这种卡的话(Kepler,3.X),适当的使用8B的宽度很多时候还是能提速你的代码的。
此外,本章节还试图对L1 Cache, Shared Memory, Texture Cache等进行改进,从本代开始,开始了这3者的分分合合的改动过程。(没错。NV改过去又改回来。也是醉了)

首先从这一代开始,L1 Cache不再像之前代的GPU一样,负责通用的读写缓冲了。当时Kepler 3.0一出来后,大家纷纷感叹,又回到了当年的计算能力1.X的时代了。为何这样说?因为3.0如果想利用SM内部的数据缓存(这里的read-only cache),必须使用纹理访问。这也是导致了论坛上持久不停的纷争,究竟是使用纹理好?还是直接读取好?是直接读取快?还是使用纹理快?也是目前市面上的很多CUDA书中这里, 对纹理的使用,乱作一团的原因。用户应当知道,如果是计算能力3.0的卡,应当考虑使用纹理。否则你的普通读写将不能使用SM内部的缓存,只能使用全局的次级的L2缓存。这也是本章节强调的原因。

但是NV很快发现这样好像用户们不适应,又在本章节的后面, 说了,如何在2代Kepler上,重新用回来L1 Cache的方式。此外,对普通用户最重要的是,对于只读的数据,提供了一个__ldg()函数。该函数能够在计算能力3.5+/5.X/6.X上,对很多应用起到显著的提速效果。(7.X我还不知道情况)

因为3.5之后,到7.0之前,几乎沿用了Kepler的这个特性,很多时候,如果你不特别的做一些处理(本章节有介绍),则L1(或者后续代的Unified Cache)是默认禁用的。性能比较惨。特别是对于一些很小的查找表。如果你不手工放入shared memory, 直接查表等于直接用全局的L2,性能可能有剧烈的损失。

这点用户需要注意。

用户可以直接将__ldg()看成是一个隐形的纹理,自动覆盖了整个显存。用它就可以在计算能力3.5+上自动的利用纹理(或者等效的)缓存,而不需要手写纹理访问的代码。很是方便。(当然,用这个只有纹理的缓存效果。其他的加速效果是没有的。详见我们之前的纹理章节)

然后还需要说明的是,这一代的L1 Cache(和Shared Memory合并在一起),是最后一代和L2 Cache line大小不同的N卡了。

不恰当的使用L1(例如你在3.5+上要求它进行普通读写缓冲),可能会起到反面效果。因为这代卡的L1是128B,L2是32B的最小传输大小。
不恰当的使用会导致过量传输,损失性能,详情可以简单的看一下本章节的描述,或者搜索更多信息(网上有很多的。特别是GTC的历年幻灯片)

然后这一代还引入了循环移位的支持,使得特定的场合下,以前的循环移位所需要的移位 + 反向移位 + 逻辑或拼合的三步操作,变成了单条指令即可。
(但是需要说明的是,该指令不是全速率的。循环移位直到现在家用的2080, 才变成全速率的)

而循环移位这种操作,在很多应用场合,特别是散列或者密码学计算中,有着广泛的应用。单具体的说,N卡的实现是通过Funnel Shift来完成的(漏斗型移位, 因为操作数的输入和输出是大口和小口的关系,形似漏斗而得名),而Funnel Shift来进行循环移位只是它的应用的一种,还有其他用途(例如用来拼接2个4B字节)
然后这一代的Shared Memory, 还是最后一代传统的N卡的原子操作风格的一代。

具体的说,这代上的原子操作都是通过读取-锁定-SP计算-回写解锁的过程来完成的。

而后续的从Maxwell起,都采用了A卡风格的Shared memory自行计算的方式(没错,后续代的显卡的shared memory存储器将带有计算功能)

NV一度在GTC的演讲中,叫它是:远程原子操作。所以大家如果看到==这种字样,不要惊讶,它就是目前最普通的shared memory上的原子操作。

这种方式很很多的好处,我们在明天到了现代的N卡的章节(Maxwell)的时候,再进行说明。


但是无可否认的是,我们目前用的很多特性的基础,都来自Kepler。







回复 支持 反对

使用道具 举报

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

本版积分规则

关闭

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

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