DAY81：阅读Compute Capability 5.x

H.4. Compute Capability 5.x

H.4.1. Architecture

A multiprocessor consists of:

• 128 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 one instruction 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,
• a unified L1/texture cache of 24 KB used to cache reads from global memory,
• 64 KB of shared memory for devices of compute capability 5.0 or 96 KB of shared memory for devices of compute capability 5.2.

The unified L1/texture cache is also used by the texture unit that implements the various addressing modes and data filtering mentioned in Texture and Surface Memory.

There is also an L2 cache shared by all multiprocessors that is used to cache accesses to local or global memory, including temporary register spills. Applications may query the L2 cache size by checking the l2CacheSize device property (see Device Enumeration).

The cache behavior (e.g., whether reads are cached in both the unified L1/texture cache and L2 or in L2 only) can be partially configured on a per-access basis using modifiers to the load instruction.

H.4.2. Global Memory

Global memory accesses are always cached in L2 and caching in L2 behaves in the same way as for devices of compute capability 3.x (see Global Memory).

Data that is read-only for the entire lifetime of the kernel can also be cached in the unified L1/texture 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.

Data that is not read-only for the entire lifetime of the kernel cannot be cached in the unified L1/texture cache for devices of compute capability 5.0. For devices of compute capability 5.2, it is, by default, not cached in the unified L1/texture cache, but caching may be enabled using the following mechanisms:

• Perform the read using inline assembly with the appropriate modifier as described in the PTX reference manual;
• Compile with the -Xptxas -dlcm=ca compilation flag, in which case all reads are cached, except reads that are performed using inline assembly with a modifier that disables caching;
• Compile with the -Xptxas -fscm=ca compilation flag, in which case all reads are cached, including reads that are performed using inline assembly regardless of the modifier used.

When caching is enabled using some the three mechanisms listed above, devices of compute capability 5.2 will cache global memory reads in the unified L1/texture cache for all kernel launches except for the kernel launches for which thread blocks consume too much of the multiprocessor's resources. These exceptions are reported by the profiler.

H.4.3. Shared Memory

Shared memory has 32 banks that are organized such that successive 32-bit words map to successive banks. Each bank has a bandwidth of 32 bits per clock cycle.

A shared memory request for a warp does not generate a bank conflict between two threads that access any address within the same 32-bit word (even though the two addresses fall in the same bank): In that case, for read accesses, the word is broadcast to the requesting threads and for write accesses, each address is written by only one of the threads (which thread performs the write is undefined).

Figure 19 shows some examples of strided access.

Figure 20 shows some examples of memory read accesses that involve the broadcast mechanism.

Figure 19. Strided Shared Memory Accesses. Examples for devices of compute capability 3.x (in 32-bit mode) or compute capability 5.x and 6.x

• Left
• Linear addressing with a stride of one 32-bit word (no bank conflict).
• Middle
• Linear addressing with a stride of two 32-bit words (two-way bank conflict).
• Right
• Linear addressing with a stride of three 32-bit words (no bank conflict).

Figure 20. Irregular Shared Memory Accesses. Examples for devices of compute capability 3.x, 5.x, or 6.x.

• Left
• Conflict-free access via random permutation.
• Middle
• Conflict-free access since threads 3, 4, 6, 7, and 9 access the same word within bank 5.
• Right
• Conflict-free broadcast access (threads access the same word within a bank).

文备注/经验分享：

Maxwell是比较神奇的一代，它类似Intel体系里面的Tick-Tock方式（分别是制程和架构的改良）。这一代也是NV的挽回颜面之做，分别在两个方面取得了长足的长进： 一个是功耗的极度优化，非常省电（折算到单位性能） 另外一个则是架构的调整，效率的极度提升，SP效率达到将近100%在Maxwell这一代变得非常现实。 然后本章节的开头部分你可以看到，每个SM（或者叫SMM，最后一个M是Maxwell的缩写）里面，从192个SP缩减到了128个SP。因为我们常规说卡的理论峰值总是从SP数量乘以频率来说起的，这样做同样的理论峰值下，因为SP数量的缩小，实际上可用的资源（如果非要折算到SP上的话）是提升的。具体的说，这一代的卡，同样的SP数量的情况下，可用资源提升了1/3. 这种可用资源的提升，伴随其他方面的改进，带来了Maxwell的高效率。在Kepler那个时期，曾经有过一款软件叫CUDA-Z，用来测试可能的峰值性能的，但是即便是这种综合性的测试里面，CUDA-Z依然在Kepler上只能表现出来67%左右的峰值性能。同样的软件不动，放到Maxwell上，却可以展现出来几乎100%的性能。这是一个很好的例子。 NV曾经在Maxwell的白皮书上说过，GM107（第一代的Maxwell）的设计目标有两个：一个是功耗的优化，一个是可以达到理论性能。 从今天的事后角度来看，这两点都达到了。它在没有改变28nm的Kepler的制程的情况下，功耗非常惊人的低。而软件方面的理论性能的容易达到，则是因为它减少了很多运算的延迟（基本延迟6周期），同时还取消了两组共享的CUDA Cores（64个SP，这是为何它变成了128SP）， 使得不再需要像Kepler那样的必须靠ILP才可能发挥性能， （还记得我们之前章节的ILP和TLP么？） 而ILP说实话，并不是一种对程序员友好的方式，它需要用户手工的进行向量化的数据运算（例如float2, int4这种），带来了很多用户角度的负担，也带来了编译器的负担。编译器很难在大量的向量化的情况下，安排好寄存器不发生bank conflict。实际上，向量化是当年AMD的VLIW架构年代，用OpenCL时候，被很多人骂的一点。好在NV即时的在Maxwell这一代改回来了。这样编译器不需要生成需要费心安排的代码，就可以取得不错的效率（因为Kepler本来那两组共享的SP就很难用上）。 目前的大家能从公开渠道看到的的Maxwell的架构图，实际上可以看到4个Scheduler每人只需要负责1组SP即可，单发射就可能满足SP的峰值。（实际上的Maxwell架构和手册里能看到的，以及和公开渠道能看到的，略微有区别，但是如果想得到真实的情况，可以混IRC上，有个疑似NV的员工，叫spec_, 在上面专门告诉你真实的架构是什么） 这一代还引入了强化版本的调度指令，一直沿用到计算能力7.X（精确的说，沿用到计算能力7.0（Volta）。但是NV额外的说明了7.5可以不加修改的执行7.0的二进制代码。所以等效于延伸到了7.5） 这代表了什么？ 用户从5.0开始的经验，特别是细微到较为深入的优化，几乎可以延伸到最新的卡上。有效的保护了单位和个人的智力投资。 此外，本章节还说明了一些其他方面的改动，例如Texture Cache和普通数据的读取，合并成了Unified Cache，大小为24KB. 这点需要额外的说明一下，避免很多人反复询问的困惑：Shared Memory和Unified Cache的确分家了，但是后者在profiler里经常依然被叫做texture cache. 经常用户会发现， 我并没有使用纹理， 然后profiler报告，纹理读取吞吐率为XXX GB/s，这是因为它们两个现在合并成一个了，用户按照本章节说明的const __restrict__修饰过的指针，或者像是__ldg读取（在Kepler 3.5+起有效），就可以直接利用tex/unified cache, 大部分情况下不需要手工写纹理读取的代码，就可以有效的利用上这个缓存，很是给力（但是对于2D, 3D有空间局部性访问的，依然应该使用纹理。这个自动化的只能适合普通的线性读取，手册这里没有说，但应该很容易发现）。 其次，从这一代起，Unified Cache（或者以后的L1）改成了32B最小访问粒度，一次在32B对齐， 32B大小的读取，只会导致Unified Cache(L1)从L2载入连续的32B，有效的规避了上一章节，手册说的Kepler具有的L1的过度读取的放大效果。 此外，手册这里给出的24KB的cache大小，实际上是有问题的。问题一在于，这24KB实际上是拆分成2个12KB的，用户用起来需要小心，这点很容易的从profiler中能看到，用户能从profiler中发现tex0和tex1两个指标，但是一般的，只要用户注意负载均衡，不会导致部分warp干的活多，部分warp干的活少 （例如warp编号除以4， 余下编号0和1的干的活多，编号2和3的warp干的活少），则profiler报告的maxwell上的这两个tex0和tex1开头的指标应该是基本一样多的。 如果用户发现了这两个一旦不平衡，则说明用户需要调整代码里的工作量，避免一个用的多，一个用的少，性能不好。这点是很容易发现的。 但是手册本章节没有说，这里需要说一下。 其次，Maxwell的卡实际上分成两种，另外一种的信息和本章节给出的信息不符合。是具有48KB的unified cache的（2个24KB的cache), 用户需要注意这点。 这种具有较大的unified cache的卡具有BUG（或者说Feature？），导致部分情况下，unified cache会突然失速，变得非常缓慢（实际上是被自动禁用掉了。后面说）。 然后本章节还类似的对shared memory画了图， 这个图实际上昨天就没有说（对Kepler），因为这个很常规的Shared Memory访问并无区别。用户应当直接参考我们之前的通用的Shared Memory章节。这里照例也不说明。但是Maxwell的Shared Memory有两处惊喜的改动。 1处是具有远程原子操作，或者说，Shared Memory本身具有运算能力，非常像GCN的A卡。Scheduler可以直接提交指令给Shared Memory(例如对某个或者某些单元+1，原子加法操作），而Shared Memory可以在无SP辅助的情况下，独立完成读取旧址，计算+1后的新值，回写新值等操作。Shared Memory上的性能相比Kepler，取得了长足的进步。实际上，NV在GTC上有过系列演示。各种图表展现在常规的Shared Memory上的原子操作的情况下，性能的剧烈提升。 请注意是剧烈提升。做图像处理的朋友可以实验一下手头的， 类似直方图统计之类的场合的性能变化。 这点是非常值得赞扬的。 从Maxwell这一代开始，N卡就将这个特性一直保留了下来。 另外一点则是Shared Memory上的Bank Conflict取得了很好的进步。在部分密集的Bank Conflict的情况下，访问延迟和吞吐率都得到了提升，（就如同部分shared Memory上的bank conflict消失了，或者被硬件自动解决了一样） 具体的原因尚未知道，但是用户能够直接的从代码的提速上感到这种惊喜（代码不变，Bank冲突率不变，但是速度就是快了），Arxiv上曾经有过文章，详细的比较过多代架构上的Shared Memory上的Bank Conflict时候的变化。读者感兴趣可以自行搜索。 除了SM内部的Shared Memory和Unified Cache这些缓存的变化外，L2这代比Kepler往往增加了不少。特别是最早出现的GM107（一代Kepler），桑心病狂的上了2MB的大L2 Cache。还是很体贴用户的。1080好像也不过是2MB的L2 Cache（记忆中的数据，可能不准，详情看NV的产品页面）。 而GM107不过是最低端的Maxwell卡。主席曾经说过，矫枉必须过正。从Kepler的改进必须这样！ 好在现在更新点的卡，理智了很多。 然后关于存储器方面的说完了，我们继续回到maxwell的SM里面看看。这里面手册本章节，隐晦的错过了一半的信息。实际上Maxwell分成两种，一代和二代，本章节的信息基本上都是关于1代的。 我说一下这两点的主要差异。 一代的Maxwell就是本章给出的信息，几乎毫无区别。 二代的Maxwell则就是Pascal，只是老28nm制程的Pascal。换句话说，用户应当知道，自己手头的Maxwell卡，计算能力5.0的应该看本章节内容。而计算能力5.2的应该直接看Pascal章节。 计算能力5.2的Maxwell等于Pascal，只是少了dp4a/dp2a这两条指令而已（相比6.1家用卡），或者少了通用FP16指令（相比6.0专业卡） 5.2的Maxwell我说一下为何和Pascal接近。 （1）Unified Cache（L1 Cache）几乎完全和Pascal（6.1）的一样。包括两者本身存在的BUG都没有改进。只是从28nm缩减到了16nm而已。 在你的CUDA安装目录下有个Maxwell和Pascal Tuning Guide，这本书里面各自都有一小段，写着：对于计算能力5.2（二代Maxwell）和6.1的某些卡， 部分情况下L1 Cache无法进行资源分配，强行启用会导致0%的Occupancy（即：kernel无法启动），这种情况下driver会自动为kernel禁用掉L1（Unified Cache），然后直接启动kernel。 这两代卡实际上可以观察到，例如某些较小的查找表，试图用__ldg在L1中查表，然后用户经常会惊喜（惊讶）的发现，改动改动代码，突然L1就完全失速了（表面），查找表的访问性能突然下降了一个数量级。这个时候就是命中了这个L1的BUG（或者说Feature）。 而二代Maxwell和Pascal在这方面毫无区别。 （2）此外，二代Maxwell的L1 Cache（Unified Cache）大小不是本章节说的24KB。而是和6.1一样的48KB（2x24KB）。 （3）二代Maxwell的Shared Memory大小和6.1的毫无区别（96KB）。而不是本章节介绍的Maxwell的64KB。

最后我们再将这一章做一个总结 回到SM内部的计算单元上。这一代具备4个双发射的调度器，SP计算密集型的代码，理论每个调度器单发射，就可以压满SP（用Profiler看到的IPC是4） 注意这里的4是warp单位的，你也可以乘以32换成得到128（线程单位），但是实际上Maxwell是如此的高效，设计的时候，SP的计算，可以和Shared Memory的访问，L1 Cache(Unified Cache)的访问，以及SFU的运算同时发射。因为SFU之类的本身数量的原因，最多可以取得将近6的IPC。 换成线程单位是192条指令/SM/周期。而此时你已经知道一个SMM里面只有128个SP。效率非常惊人（150%的效率哦）。所以Maxwell这一代真心好卡。然而Maxwell也有短板，就是双精度只有1/32.3%的双精度，几乎等于没有。 所以用户应当直接将二代的Maxwell看成6.1的Pascal，只是少了两条深度学习指令，功耗大点（28nm老制程么）。其他毫无区别。看上去NV是直接最后优化优化Maxwell，改个制程改名叫Pascal就放出来了。 而本章节Maxwell的介绍，故意不说二代Maxwell（二代Maxwell还是很多的，包括980， Titan之类的），还包括含有Maxwell字样的Quadro卡， （一代的Maxwell的Quadro卡，很多以Kepler的名义卖掉了。名字是Quadro KXXXX风格。用户买回去发现计算能力是5.X。而不是3.X） 所以本章节的手册，提供信息不准确也可以理解了

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

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