DAY47:阅读read only cache和Time Function

B.10. Read-Only Data Cache Load Function

The read-only data cache load function is only supported by devices of compute capability 3.5 and higher.

T __ldg(const T* address);

returns the data of type T located at address address, where T is char, short, int, long longunsigned char, unsigned short, unsigned int, unsigned long long, int2, int4, uint2, uint4, float, float2, float4, double, or double2. The operation is cached in the read-only data cache (see Global Memory).

B.11. Time Function

clock_t clock();long long int clock64();

when executed in device code, returns the value of a per-multiprocessor counter that is incremented every clock cycle. Sampling this counter at the beginning and at the end of a kernel, taking the difference of the two samples, and recording the result per thread provides a measure for each thread of the number of clock cycles taken by the device to completely execute the thread, but not of the number of clock cycles the device actually spent executing thread instructions. The former number is greater than the latter since threads are time sliced.

本文备注/经验分享：

先说read only cache了. 这东西Kepler专有的, 而且因为有两种Kepler的存在(精确的说, 三种, 还有K80单独列出)，这东西也分成前后两个阶段.计算能力3.0的Kepler将L1和shard memory做成一体的, 然后取消了L1 cache的缓冲普通global memory访问的能力, 改为只能缓冲local memory传输。

因为当时的计划是, Kepler上来就能1个SM里面有192个SP, 将来还可能会更多, 例如一个SM里面假设能有几百个SP. 然后推出市场后, 发现性能表现的很不理想, 于是慢慢取消了这么做, 改称5.X/6.X/7.X的128个SP到64个SP/SM这种较少数量的. (A卡也是64个SP每个CU)， 基于当时的这种考虑(是不是有点像Intel当年失败的NetBurst架构——P4使用过, 当时P4说一直推出到高频率10Ghz版本的U), 一个SM里面的SP数量会变的越来越多, 里面SP们对global memory上的工作集大小可能会变的越来越大, L1对global memory的缓冲效果也越来越差. 因此将L1改成了只能缓冲local memory, 而同时引入了read only cache， Kepler 3.0上, L1 cache默认有16KB, 32KB, 48KB大小可调.

(对应的, 和它一体的shared memory则相应的缩水到48KB, 32KB, 甚至只有16KB)，而引入的新read only cache却很大, 足足有48KB，但是3.0上, 这个cache只能服务texture/surface访问,因此3.0的卡实际上如果想利用这个cache, 就必须回推到最初的计算能力1.x那样, 必须使用texture了.这实际上是一种倒退(还不如2.x的fermi)，这也是经典问题的由来, "我的XXX代码, 是使用texture读取好? 还是普通读取好?"答案是不一定的. 如果是1.X和3.0, 那么可能会texture好(也要看情况)，如果是2.x/5.x/6.x/7.x,那么可能texture反而不好.这就是为何很多人天天气势汹汹的在论坛上发文, "我使用了texture, 按照某书(2009年)的说法, 怎么性能下降了呢"。这必须要看具体的代码访问方式, 和具体的计算能力了.因为3.0的这个48KB每SM的read only cache, 必须通过texture才能用.很多人进行了批评和表示不能接受.加上kepler在表现上的失败，从2代Kepler开始(计算能力3.5), read-only cache进行了改良,引入了__ldg()指令, 也就是本章节的这个.允许对于普通的global memory方式, 能利用上这个read-only cache，请注意, 3.0的老卡依然不能这样利用. 只有3.5/3.7的Kepler才能.这样read-only cache就在表现上更像一个只读的普通L1 cache了.

用户可以直接将__ldg()看成是一个作用于隐形的texture, 其后备存储覆盖了整个global memory的线性存储, 这样来用。这里还需要说明的是:

(1)3.5/3.7的硬件能自动发现一些只读的数据, 自动尝试利用该read-only cache. 特别里手册中提到的const type * __restrict__ p形式的指针.

(2)如果怕编译器的自动处理不保险, 你总是可以要求使用手工__ldg()函数来手工利用该cache.

该函数实际上是利用了内联的ptx, 手册也多处提到了内联ptx的使用, 很多时候很有用. 特别是底层PTX有的能力, 但是没有导出到CUDA C中的, 我们总是可以通过这些方式来尝试使用它们.

然后还需要说明的是:

到了计算能力5.0+(Maxwell/Pascal...), 形式又发生了变化, 虽然取消了read-only cache, 但是__ldg依然有效.在5.0的Maxwell上, 取消了read-only cache, 拆分了原本一体的L1 + shared cache, 变成了:

(1)新的Unified cache: 是原本的texture cache + L1 data cache (2)新的独立的shared memory: 总是固定大小的64KB. 不可调(5.2和6.1的卡将这个扩大到量96KB)。

而此时的unified cache的利用方式很像3.5/3.7的read-only cache,同样是可以自动缓冲编译器发现的一些只读数据, 也同样可以用户手工的通过__ldg()来使用该cache.

然后到了Maxwell第二代和Pascal 6.1上, 形式继续发生了变化.除了刚才说的独立的shared memory增大到了96KB外.该cache还能服务于可写的数据.

也就是在这些卡上, 实际上unified cache变成了原本的texture cache + 完整的L1 cache，此时无论你是否使用__ldg, 还是只是要求编译器自动利用unified cache. 都可以了.在这些卡上(也包括只读的情况下的5.0的maxwell), 你在使用profiler的时候,如果发现你完全没有使用texture, 就是在普通指针访存.profiler却报告了Texture访问的大量指标, 不要惊讶. 正常现象.大家一体了.也就是这些卡上, 除了独立的shared memory, 其他SM内部数据缓存都融合了(除了constant cache---这是另外一个话题)。

然后到量7.X, 情况继续发生了变化.texture cache + L1 cache + shared memory, 完全融合为统一的128KB的缓存了，(其中的分配给shared memory的部分大小可调)，此时已经只有一种cache了.但是历年来的__ldg依然可以正常使用.只是不同的卡上, 性能不同而已了.

需要注意的是, 当年有个奇特的卡, 叫K80，K80的1个SM实际上等于普通的Kepler的2个SM的资源(除了SP数量. 或者说, 资源不变, 只有等效的一半数量的SP)，它具有特别巨大(在当时的那个年代来说)的L1 + shared memory，可以配置成16KB L1 + 112KB的shared memory, 或者选择96KB和80KB的shared memory。在当时很是不错了。

需要注意部分计算能力上使用L1 cache可能会导致的额外问题，例如无法启动kernel---CUDA Runtime/Driver此时会自动禁用相关卡上的L1 cache---原因暂时未知(原因知道, 但是NV不愿意公布, 所以这里也不说)),这个在之前的阅读的时候有过提到.

总之虽然read-only cache不在了, 但是__ldg的效果依然存在.

第二部分，关于clock和64-bit计算值版本的clock64()， 它可以用来计算一个线程里面的单个代码片段的执行时间(周期数)，但是请注意的是, 实际上本章也说过, N个线程, 每个线程执行的时间和相加, 并不等于总执行时间.(总执行时间一般总是要远远小于所有的线程的时间累加的. 例如相差1个甚至多个数量级)，所以该函数的主要测量用途实际上并不那么实用.但在很多场合, 例如进行微架构测试的时候, 还是有用的. 正常用户一般情况下, 不需要使用它.(偶尔可能需要使用, 例如你在自行实现了一个互斥锁的时候, 你可以通过clock()来避让一段时间再尝试竞争锁的所有权)， 但是需要说明的是(虽然手册上有这么一个例子), 在GPU上使用mutex之类的东西本身可能是一个设计上的问题) 以及, 需要补充的是. 7.X新引入了休眠机制, 可以更有效的在类似应用场合进行避让.所以你看. NV还是很厚到的. 只要硬件增加什么新功能(例如TensorCore), 就会立刻让你用。

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

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