打工人的时间是如何计算的
作者：GPUS开发者
链接：https://zhuanlan.zhihu.com/p/339478619
来源：知乎
著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。
今天主要说两点, 一点是如何正确的计算一段操作所用的时间。这里的一段操作是指的, GPU设备上的kernel计算, 以及, 数据传输操作。
正确的计时也是从今天开始的, CUDA优化章节的重要基础，因为你的代码干了什么, 例如对一张图片进行边缘查找, 或者颜色分布进行直方图统计, 这些工作量你本身, 作为代码的编写者, 是知道的. 此时再加上了正确的计时方法, 则你可以立刻衡量出来, “我的具体XXX操作过程, 在XXX ms内完成, 性能是XXX”(例如10张图片/秒)。
但是我们历年来, 很遗憾的看到, 大部分人的做法都是错误的. 甚至使用了错误的测时结果, 来气势汹汹询问一些问题. 此时, 因为你的基础部分(计时)是错误的, 从而导致了你的问题整体无效.
这点无论是从, 我们的论坛上的帖子中, 还是我们的直接的客户支持用, 用户给出的他们的代码中, 都可以看到这样的错误.
今天我们就说一下, 这些错误的根源, 和正确的计时方式该如何进行. 错误的计时根源往往有两种, 一种是对GPU上的代码片段的执行的特性, 具有误解。
例如在我们之前的文章中, 我们知道一个kernel的启动是异步的, 也就是一旦该kernel成功启动后, 它就开始在GPU上执行了. CPU这边的诸如<<<>>>()的菱形启动符, 是会在kernel完成了启动后, 就立刻返回CPU上的下一行代码执行的.
CPU并不自动等待GPU上的工作完成！
这点是CUDA在设计的时候, 为了充分能让GPU作为一个劳工的身份, 去完成一些重活而设计的； 而CPU作为CEO, 并不需要在”GPU劳工”辛苦忙碌的时候, 必须啥都不干的同步等待的。
就如同一家公司里的老板, 布置出来了活给员工, 那么员工在干活的期间, 老板并不是必须等待员工慢慢干完, 才能返回老板自己的下一个工作事项的。老板完全是布置完活后就没事了,然后可以继续给另外一个员工布置活, 或者自己悠闲地去喝着茶了。
这点说起来很简单, 但是很多人都在理解上犯了错. 我来举个例子.
https://bbs.gpuworld.cn/index.php?topic=73413.0
例如本帖, 本帖楼主犯了一个常见的错误, 没有等待kernel完成, 就立刻对它进行计时, 然后得出了错误的问题前提: “一个kernel如果被反复调用的话, 是会越来越慢的”。
我们看下该楼主的具体做法:
1 | start = clock(); |
楼主这里直接测了起始时刻start, 然后立刻用<<<>>>调用了自己的kernel, 然后不等该kernel”实际上的完成工作”, 就立刻测量了结束时间end, 然后就认为从start到end, 这两段时刻的差值, 是kernel的实际执行时间, 这是严重错误的。
这就像公司老板, 先看了一下手表, 现在的时刻是1点29分, 记录成Start; 然后叫了员工如花说,”如花，去把上次和我们合作活动的NV公司的联系人, 沟通一下XXX事宜”; 然后对如花说完这话后, 立刻又看了一下手表, 现在是1点30分, 记录成End.
然后老板认为,如花完成和某公司的沟通工作, 一共用时: 从1点29分到1点30分, 共总1分钟.
这显然是严重错误的. 这样的计时方式, 并不是员工实质完成一个工作的时间, 而只是老板(CPU)对员工(GPU)的派活, 所耗费的时间. 并不能实质衡量某工作的时间的.
类似的, 该帖子的楼主也犯了这个错误, 他也是立刻用<<<>>>给GPU派活后, 立刻看了一下表, 从而导致他理解得到了错误的信息, 从而让整个问题化为无意义. (错误的前提下, 给出的提问是无意义的)。
那么正确的做法是什么呢?
正确的做法(之一)是, CPU在给GPU派活前, 的确可以记录时刻Start; 但是一旦给GPU派活后, 必须等待GPU完成该活, 才能记录时刻End. 此时的End减去Start, 才是真正的干活耗时。
这就像公司老板给员工如花布置活前, 记录了1点29分为start时刻; 然后给如花布置了沟通联系的活了后, 老板等待, 例如2点00分, 如花届时完成了该活后, 才记录为end时刻.
此时的end - start = 2:00 - 1:29 = 31分钟, 才是如花真正干完该活所用的时间. 这样才是正确的.
不仅仅如此, 我们还会在今天的内容中看到, 除了老板自己去计时的方式, 我们还可以要求员工(GPU)去计时, 即员工如花自行在自己干活前记录一下开始时刻, 然后去干活, 然后员工如花在干完后, 自行也再记录一下结束时刻, 然后并将结束和开始的差值, 作为干活时间, 汇报给老板(CPU)即可.
回到该楼主的帖子, 我们很遗憾的看到, 该楼主在我们给出了两次回答和解决方式建议后, 即分别要求楼主用第三方工具验证他的计时错误的前提(这样他可以自行发现他的错误, 从而增长经验), 和直接给出了建议(即明确的告诉了他哪里是理解错了后), 他均无视了我们. 并继续在后续的跟帖中, 给出他自行认为的理解. 这点我们是感觉非常可惜的.
实际上人是互相尊重的, 特别是在作为提问者, 你更加应该尊重回答者给出的信息的. 无视这一点, 并取得”面子上”的好处, 是无益于事情的. 我们在这里今天严肃的提出这一点。
是希望其他的客户或者非客户, 在论坛提出了问题后, 在看到论坛给出的解答后, 不要为了”面子”, 带着有色眼镜, 从而实质上的无益于楼主们在论坛的经验的获取, 和以后遭遇相似问题时候的快速解决.
(反过来, 如果你尊重了论坛, 则你本次能反思得到了经验, 得到技术上的成长; 下次遇到后还能快速回忆场景, 快速解决, 节省干活时间, 增加在老板心中好的评价).
然后我们继续说一下该例子, 楼主的正当做法应该是:
- CPU记录开始时间
- CPU给GPU派活
- CPU等待GPU完成该活
CPU记录结束时间
我们在这里插入了步骤3, 也是手册上今天的CPU计时内容章节, 所推荐的做法(cudaDeviceSynchronize()同步等待, 或者其他任何等效的同步方式). 只有加上了该等待, 你的开始到结束的时间差, 才是真正的干活时间.
似是而非的计时方法
作者：GPUS开发者
链接：https://zhuanlan.zhihu.com/p/339698093
来源：知乎
著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。
我们继续回到今天的第一个大话题, 正确的计时. 因为这话题的确很重要了, 没有了正确的计时, 一切对工作效果(代码运行快慢)的衡量工作, 都会变成虚有.
首先, 我们已经说了, 正确的逻辑顺序. 即CPU开始时刻记录->CPU发布任务给GPU->CPU等待GPU完成->CPU记录结束时刻。 这4个步骤, 任何一个步骤错误了, 都会导致错误的结果。
我们还在论坛上经常看到有人会: 启动kernel->记录开始时刻->记录结束时刻, 这样的做法也是错误的(记录开始时刻必须在启动kernel前)。特别的, 在使用非专业卡, 或者在WDDM驱动下, 该问题可能会被掩盖，即哪怕你先启动了kernel, 然后才记录开始时刻, 有的时候看起来”结果是对的”, 这是因为WDDM驱动之类的有缓冲, 它可能在某种情况下会导致kernel的启动, 是被延后执行(正好插入到你记录完开始时刻后), 从而导致可能的错误的”咦, 我本次这样写也对了”的假象，从而隐藏了问题, 从而让书写者本次可能会认为, 这样写没问题的假象。我们这里严重强调一下, 必须保证正确的逻辑顺序, 才能得到正确答案。
回到今天的计时章节. 除了正确的逻辑顺序, 和工作逻辑流程上的保证(即在开始工作前计第1次时间, 必须等待工作实际完成后, 才能计第2次时间). 我们还需要对基本的计时工具本身, 进行讨论. 本实践指南手册上, 在今天这里简单的说了一下, 你应当在正确的操作系统平台上, 使用正确的计时工具/函数调用. 而我们准备详细的说一下这点， 分别对应常见的台式机上的Windows开发平台, 和我们目前在售的嵌入式上的Jetson平台(Linux)上的正确做法. 这里我们只推荐两种正确做法(也有其他的, 但这两种是推荐的).
(1) 在Windows上请使用QueryPerformanceCounter()/Frequency()这两个函数来进行计时. (技术上的理由: 它使用了主板上ACPI提供的HPET计时器, 该计时器在常见的主板上, 保证了至少几个Mhz+以上的时间分辨率, 足够)。
(2) 在Jetson嵌入式平台上, 使用gettimeofday()系统调用来进行时刻的记录. 它也同样具有很好的时间分辨率/精确性.
然后在任何一个平台上, 均不建议使用__rdtsc()或者clock()函数进行计时. 这里要重要说一下。这两个函数均是多年来, 在我们的用户中, 流行的两种方式, 可惜它们均存在一种问题(真的是好可惜. 对的选择的不多, 错误的大家都很喜欢).
我们还用上文提到的那个帖子的错误做法（https://bbs.gpuworld.cn/index.php?topic=73413.0）举例好了(没错, 除了之前的那个计时错误, 该帖子还有更多的计时错误). 该帖子是今天的实践手册上的正确做法的, 多个连续反面例子. 还是很起到很好的警示作用的.
该帖子中, 记录时刻使用了clock(). 很可惜, 该函数在不同的环境下, 有不同的返回值解释. 我们分别看一下MSDN上的clock解释, 和Linux上的man(3)手册中的解释.
首先是Linux上man手册的解释, 这是我们的jetson平台. 在该平台上, 它返回的是调用者进程所耗费的”CPU时间”, 什么叫CPU时间, 和我们常用的时间有什么不同?
举个例子说, 我在CPU上执行了一个可执行文件, 上去执行了3秒, 然后突然打开了一个巨大的磁盘文件(假设我们用的是普通的机械硬盘), 或者需要从网络读取一个资源, 而突然卡顿上了10秒无响应, 然后又立刻执行了4秒. 那么此时, 它实际上使用的CPU时间是7秒(4+3), 也就是中途它卡在磁盘读取或者等待网络数据返回的时候, 那10秒虽然流逝了, 基本上是不耗费CPU的。这和我们常用概念中的实际时间(一共经过了17秒, 4+10+3=17), 是不同的.
因此你看, 虽然实际上计算了很久, 但是从”CPU时间”的角度说, 可能很短, 因此该函数和我们常用的实际生活中的时间概念是不同的, 在我们的jetson的linux平台上, 是不能使用的.
我们平常日使用的时间, 在计算机中, 叫real time, 即中文翻译是”实时钟”;也叫wall time, 就是你挂载墙上的钟上经过的时间.
幸运的是, 微软的Windows平台上的clock()实现(请参考MSDN), 是返回的wall time, 也就是可用的我们的实际时间, 那么看上去, Windows上至少能用它?
实际上很可惜, 和jetson上一样, Windows也不能用. 这是为何? 因为很多客户在使用的时候, 只考虑了该计时方式的逻辑上的意义, 而没有考虑该计时方式的精度/时间分辨率. 在Windows平台上, 该函数的分辨率只有计时Hz到小于1Khz, 用人话说就是, 假设是50Hz, 它最小的分辨率只有20ms(1秒=1000ms, 分成50个周期). 而我们的代码运行的速度往往很快, 某个片段往往很短, 例如13ms, 和37ms(随意的举例), 用该函数得到的结果可能就会分别取整到了0ms, 和20或者40ms了.
此时时间都错误的离谱. 这就像我们用墙上的钟表的秒针(最长的那个指针)来计时一样, 它的分辨率只有1s级别, 如果我们的代码运行了300ms, 你会发现秒针没动, 运行时间为0;
或者我们的代码运行了1.7s, 你会发现秒针动了1下或者2下, 时间也错的离谱.
因为秒针的时间分辨率/精度不够, 所以不能用来计时。而clock()也存在类似的这个问题. 所以也不能用。
我们需要的是逻辑正确, 精度足够的计时器.
也是今天实践手册上, 在说”计时器”的选择章节, 强调的重要因素. 幸运的是, 我们的QueryPerformanceCounter()和gettimeofday()在2个平台上均可以满足这两点要求. 因此它们才变成了今天我们推荐的计时工具(CPU端, 或者你理解成老板专用工具). 此外, 今天的两个被拒绝的工具中(clock & __rdtsc), 后者也存在一处或者多处问题, 因此也不能用.
rdtsc主要是存在时基漂移(在后期的CPU和主板中逐渐的解决了), 不能跨核心同步, 以及, 还有rdtscp版本来解决其他”CPU乱序执行上”的其他问题. 这些问题或多或少的在后来的CPU/主板/操作系统中都逐渐解决了, 但是我们不敢打包票. 因此也不推荐使用.
GPU端的CUDA Event计时
好了. 你已经会了CPU端计时了, 记住, 正确的计时逻辑顺序, 和使用正确的计时工具, 这两点满足了, 你就会有正确的测时结果. 我们继续说一下GPU端的计时. 和CPU端的计时类似, 它同样需要2个方面: 正确的逻辑, 和正确的工具使用.
在开始这两点之前, 我们先说一下GPU端计时的优势和特色.
优势和特色主要有两点, 1个就是可以将计时本身当作命令发布下去, 而不需要一定在特定的时刻, CPU亲自动手去记录. 2个就是可以方便记录比较复杂的计时场景(特别是多流和传输/计算异步重叠的时候). 我们先说一下1点.
还记得我们之前的例子么? 老板让员工如花去完成一个活, 然后老板在如花开始动手之前, 和如花完整的完成了工作后, 分别进行了时间记录. 这个例子还可以这样做——老板: “如花，你去干XXX活. 干活前后你记下时间, 最后将这个活和用时都汇报给我”. 这种方式相当于是老板将计时本身的任务, 当成活布置给了员工, 这样老板可以在半夜12点突发奇想, 通过微信给员工如花布置任务: “明天9点上班后, 干YYY. 我晚点来, 你统计一下时间”. 而不需要老板必须在明天9点那一瞬间, 亲自不布置记录.
也不需要老板时刻的焦急的等待如花去完成, 最后在如花于11点完成的瞬间, 立刻找笔纸记录下来结束时间.
大大减轻了老板的调度成本, 和指挥公司运营的压力. 类似的, 我们的GPU作为一个劳力或者说协处理器的角色, CPU也需要调度它.
通过GPU端计时, 我们可以将计时本身的任务, 布置给GPU即可. 这样CPU上的调度(代码)可以有更自由的安排, 也减轻了用户们写代码上的逻辑安排的压力. 我们具体看看怎么做:
GPU上的计时, 是通过CUDA Event来完成的, 它可以理解成一种非常轻量的空白kernel, 只用来记录一下时间而已 (因此很多用户忧虑的, GPU上执行event的记录工作, 会不会拖慢GPU —- 完全不会的).
具体说, 是通过在特定的CUDA流中, 发布一种叫cudaEventRecord()的任务进去而已.
这样, 该流中的命令们, 一旦当GPU执行到”记录Event”的时刻, GPU就立刻记录一下当前的时间(注意, 是从GPU的角度, 有它的时间分辨率. 本实践手册保证了至少2Mhz+的分辨率/精度). 然后继续往下执行该流中的其他常规任务(例如kernel计算). 这种记录几乎完全不占用GPU的处理能力.
所以在GPU上, 我们可以知道, 该工具(CUDA Event)是精确可靠的计时工具, 那么只剩下来逻辑的正确性了. 保证了后者, 你就可以得到了GPU上的正确计时, 不能保证, 则一切无从谈起. 但是很遗憾的, 我们从这10年来的客户反馈上来看, 很多客户并不能合理的安排一个GPU上的计时逻辑. 从而导致了错误的解决.
我先说一下GPU上正确的逻辑安排应当是一个什么顺序的:
假设用户已经有了1个CUDA流stream, 2个CUDA Event分别是start和end, 现在需要对该流中的1个kernel K, 进行计时, 正确的逻辑是:
- cudaEventRecord(start, stream); //在流中发布计时命令, 要求记录start时间
- K<<<….stream>>>(); //在流中发布kernel K
- cudaEventRecord(end, stream); //在流中发布计时end时间
- 同步
其中第4点非常重要, 常见的有3中做法. 即cudaDeviceSynchronize()进行设备同步, cudaStreamSynchronize()进行流同步, cudaEventSynchronize()进行Event同步.
其中设备同步是大家喜闻乐见的, 相当于老板等待公司人员全部空闲下来的时候, 再检查两个start和end时间(的差). 例如老板可能会等待晚上9点, 发现都下班了, 然后再优先的拿出今天如花完成工作K的记录本, 查看一下K的前后时间, 得到一个用时.
这种方式虽然最简单方便, 但是老板可能会在一个很晚的时间后, 才能得到今天的工作汇总(因为你进行了设备同步, 等待设备(公司)上的所有工作完成后才能得到这个汇总), 很多时候不恰当, 或者导致GPU设备/公司运营效率低下.
第二种方式, 则是进行流同步, 大致相当于员工同步. 老板可以等待如花突然闲置下来了, 然后拿出如花的工作记录本, 查看一下她完成工作K的信息, 和前后工作的记录时刻. 从而知道了如花对工作K的计时. 这种方式好很多, 因为此时, 另外一个员工翠花可能依然有活在干, 时间也不过是下午3点, 老板及早的知道了, 还说不定有余力能调度其他事项. 提高公司运营效率.
第三种方式, 则是进行事件(Event)同步, 这相当于员工同步里的细项. 特别是在该员工有连续的多个活的时候非常好用(例如老板给如花布置了活K和K2, 并要求在K完成后立刻计时). 老板可以等待员工如花完成了工作K, 并记录了结束时刻的那一个瞬间, 立刻从沉睡的沙发上惊醒, 然后立刻检查如花该工作的信息和前后时刻. 而如花此时本身, 已经继续去干下一个活K2了.
这样老板不仅及时的在惊醒的瞬间, 慢慢开始泡茶喝(相当于CPU上的后续调度处理)检查如花的活K的相关信息的时候, 如花自身还在干下一个活. 提高了老板和该员工的同时的调度和工作效率.
所以你看, 最应当做的应该是方式3(对事件进行同步).
但是虽然事件同步很好用. 但是我们很遗憾的看到, 很多用户并不能正确的使用它.
毕竟这就如同很多家公司存在, 并不是所有的公司的老板, 都有能完善强力的调度协调能力的. 我们分析了一下历年来用户们不能正确的通过事件同步, 来计时的一些问题, 主要暴露出来的问题有这些点:
用户不能理解cudaEventRecord()只是发布了一个让GPU计时的”任务”. 这种发布并非是当前的CPU发布命令时候的时刻, 而是GPU上实际执行到了该计时任务处的时刻.
还用我们刚才的例子吧. 老板半夜在12点发布了微信命令, 如花在第二天的9点才开始干活, 那么实际上执行开始时间记录(cudaEventRecord(start, straem))的时刻, 是第二天的9点! 而不是半夜的12点!
这点相当多的用户都理解错了. 一定要注意.
其次则是, 必须要等待实际上的stream中的K任务完成了, 并记录了后续的stop时间后, 才能用两个时间做减法, 得到夹在中间的K任务的真正耗时.
也可以看我们之前的举例, 如花在9点开始干活, 然后干了2个小时的K任务, 完成于11点, 并记录完成事件stop; 然后她继续从11点又干了3个小时的任务K2, 以及其他各种任务到下午5点下班. 然后工作里的其他员工都干到了晚上11点才下班.
那么作为老板, 你在10点立刻去尝试减掉开始时刻9点是不对的, 因为该活并没有实际上的完成. 从晚上11点(设备同步)去检查, 发现是上午11点完成的, 得到11-9=2, 是对的; 从下午5点(如花下班, 流同步)去检查, 发现也是上午11点完成的, 也得到2个小时, 也是对的; 从上午11点整去检查(如花完成记录K完成后的stop事件时间), 也能得到2个小时, 也是对的.
这分别对应了我们的cudaEvent/Stream和DeviceSynchronize()三个同步调用.
读者们可以大致评估一下效果, 但不管怎样, 你要记住, 发布记录命令本身也是一个任务, 必须等到该任务实际上完成了记录才可以(用3大同步去等!). 以及, 切记任务实际上的完成记录的时间, 和你发布这一系列命令的时间毫无关系(你在半夜12点的微信上发布的好么!)
记录这两点, 大致你对GPU端的cuda event计时就没有大问题了.
GPU端Event计时的重要特色
作者：GPUS开发者
链接：https://zhuanlan.zhihu.com/p/340203355
来源：知乎
著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。
我们在上面的内容中说过, cuda event计时还有它的丰富的特色, 你已经看到了它能正确的计时, 还不耽误老板(CPU)上的提前半夜调度的便利. 我们下一个要讲的, 就是说它可以方便的跨流, 跨一堆任务进行计时. 但在说这个特色前, 我们需要将手册的一点说法进行修正。
本实践手册在今天的GPU计时章节, 说了, 一定要在默认流中进行调用啊, 等等的话, 这个说法实际上是不对的。
(坑：手册的这段有它的历史来源, 已经10年不管了，不用看手册里的解释，它的确是不对的。CUDA这10年来, 历经了v2 API变更, 从每个host线程独享一个context变化到共享; 历经了对非默认流同步的流的变更, 等等。）
手册这里的一定要在默认流中同步的做法, 减少了调度上的灵活性, 并实际上导致了性能的下降。我们不建议一定需要在默认流中记录并同步, 这是完全没有必要的. 什么手册上说的, 在其他流中记录会不准的, 在我们的多年使用中, 并不存在这个现象。
我们在修正了手册的这点说法后, 继续到GPU端的event计时上的其他特色。
这种其他特色是指, 可以对一个实际的操作序列整体, 进行计时。
我们还用一家公司来举例好了. 该公司今天有3个主题: 联系另外一家B, 得到信息. 然后根据信息开始工作K. 同时今天还有消防检查, 要一大早的就进行安全隐患排查。
于是老板想知道, 在有消防检查主题存在的情况下, 今天的工作效率(联系合作伙伴公司, 并后续的进行工作K)还能有多少, 这就涉及到了计时, 一个比较复杂的计时。
这实际上是比较实际的应用, 因为你单纯的今天公司啥都没有, 只专心干1-2个主题的情况比较少见，很多时候都是公司里有这种复杂场景的, 此时的工作效率评估很有必要。
这里可以用修复了手册上”必须在0流/默认流中”进行记录的说法, 此时完全可以有3个员工, Tina, Linda, Rabbit, 来同时开始做事, 它们大致对应了3个流。
老板首先对Tina下令, 序列如下: 记录你的开始工作时间Start, 你联系公司S光, 然后将结果告诉同事Linda, 你继续干你的日常其他活。
接着老板对Linda下令, 你先干你手工的日常活, 等Tina告诉你了S光的信息后, 你根据这个修改报价, 并完成当前积攒的出货. 然后你这两个干完后, 记录结束时间End。
最后老板对Rabbit下令, 你今天刷一天的墙, 并将灭火器的位置和容量都检查好, 等待迎接消防检查.
然后老板感兴趣的是这里的在Rabbit忙碌于今天的卫生环境处理的时候, Tina从联系S光开始(Start时刻), 到最后的Linda完成改价格和出货的时刻(End时刻), 这两个人都用了多久, 效率如何.(注意, 为了消防检查, 已经占用了一部分公司Rabbit资源).
这个时刻即可使用上刚才说的GPU端的Event计时, 配合多流同步操作. 即分别是:
流Tina中: Record Event Start -> 联系S光 -> Recrod Event 联系Done-> other jobs
流Linda中: 设备端同步Event: 联系Done -> 改价格并发货 -> 记录完成事件Stop -> Other jobs
流Rabbit: 刷墙 -> 消防任务
这是一个比较实际的GPU应用的流程, 你将这些都替换成流中的异步传输, 替换成kernel间的依赖, 替换Rabbit成一个不能压满设备的需要持续调用的小kernel. 即可符合实际中的常见情况。
这种常见情况必须这样操作, 才有可能充分利用GPU. 手册里的必须在特定流中做特定事情, 无法正确评估实际应用场景中的时间和性能. 所以我们举了这个例子.
好了, 回到这个例子. 注意我们这里有好多Event, 那么老板关系的主工作流程是从哪里到哪里? 是到Tina完成联系S光(Event 联系_Done)吗? 不是的.
实际上是到流Linda中的”发货完成事件”。
需要老板在至少等待到该事件完成后, 才可以评估联系同行—修改价格—发货(或者发布新报价)这一系列流程中的复杂任务, 的时间和工作性能表现. (当然, 老板也可能等晚上全部员工散场后再看, 但那样可能会影响老板本来调度能力, 因为她可能半下午就可以根据情况决定是否要发布新任务或者修改计划了)
你会看到, 我们这里复杂案例, 跨越了多个流来计时, 并在轻量背景GPU任务持续存在的情况下(Rabbit), 合理的评估了员工Tina和Linda的整体工作序列的表现. 这是修复了手册上的错误说法, 并给出的非常实际的例子。
GPU显存的特色
今天的内容将首先介绍GPU的显存的特色, 包括巨大的带宽, 巨大的延迟, 和最小传输粒度等. 这些都是很多讲CUDA的书容易忽视的地方(它们往往注重在计算本身)，然后再具体分析GPU的片外的显存, 片内的shared memory, 和片内的各级缓存的情况. 这些情况的分析对于写出一个效率良好的CUDA代码, 还是很重要的. 你只有先了解一个东西, 才能应对好这个东西. 而且今天的实践手册章节还介绍了一些坑和注意事项, 以及, 给出了一些经典的例子(但是今天我们可能不说这个例子).
我们先看看GPU的显存的特色.
GPU的显存的特色主要在于, 相比CPU依靠一级, 二级, 三级(4级)等各级缓存, 以及, 充分的硬件prefetch等措施, 来降低硬件所预估的某个小范围内的某个随机数据的交付时间(即, 降低一定范围内的读取延迟), GPU并不依靠这点。
GPU上的显存大部分情况, 是具有非常长的延迟的. 这种非常长的延迟并不会影响显存的巨大带宽的发挥, 手册今天你看到了一张7.0计算能力的V100的卡, 具有将近900GB/s的峰值带宽, 非常惊人的。
而之所以能不影响显存本身的巨大带宽的发挥, 是依靠我们上次说过的, GPU上能够海量并发执行的线程数量, 当一个线程在等待超长延迟的读取的结果的返回的时候, 能够切换到下一个线程(或者更高的执行调度单位)上去执行, 从而将这个延迟给掩盖掉(详见我们之前的CUDA C编程指南手册). 这样, GPU的显存和CPU的内存, 或者说GPU和CPU上的存储器的体系结构, 具有截然不同的鲜明特别. 也就是我们常说的, CPU(的各级大容量和快速访问)的缓存, 是为了降低延迟而设计的; 而GPU(的缺乏这种各级大容量的缓存), 则是为了尽量将有用的晶体管交给计算或者资源(例如寄存器)进行延迟掩盖, 从而能发挥接近峰值的吞吐率, 而设计的.
所以, 这种”高延迟”和”大带宽”并存的显存访问方式, 将会伴随我们几乎以后所有的CUDA代码的编写时候的考虑。
此外, 今日的实践手册还指出了. 大部分的CUDA代码在追求极致性能的时候, 所面临的主要矛盾, 往往是在显存访问上. (而不是在其他的几个矛盾, 例如计算性能上)， 不过这是不一定的, 虽然大部分的GPU程序可能如此, 但是很多还是强烈的需求计算性能的, 而不一定短板是在显存的带宽之类的上. 而大部分GPU程序只所以如此, 往往来源于以下3个方面:
(1) NV在生产显卡的时候, 往往搭配一定的计算性能和显存性能. 例如说, 低端的只有几个TFlops或者TIOPS的卡, 可能往往配备128-bit的GDDR5/6的显存, 带宽本身就较低. 一般这种搭配, 对于很多算法实现, 总是计算性能够, 而访存带宽不够的. 所以出现了这种现象。
这种现象往往需要通过深入的实践GPU上的CUDA代码书写, 才能逐渐在具体硬件上得到缓解(例如今天下面要讲的, 如果在固定的一定峰值性能的显存的卡上, 充分发挥出来global memory的性能. 或者充分使用shared/texture之类的, 来尽量节省显存的访问, 来发挥性能, 这样的优化)，总之你是直接买回来一张卡, 就试图认为我随意写写, 显存就不是问题, 问题在GPU核心芯片的计算上, 那是不可能的。
老黄的搭配总是有它的一定道理的, 总是你不好好优化存储器访问, 随意的写写, 是会卡住的. 当然, 这也是为何本手册的下面的各个不同的存储器/缓存的优化指南存在的意义。
(2) 很多代码是从CPU迁移过来的, 或者说老代码逐渐优化到GPU上了. 以前在CPU上的时候, 计算性能和访存性能拉得并不开, 但是随着GPU的到来和飞速发展, 它的计算比访存的差距, 拉得越来越大，例如原本在CPU上已经做好的东西, 在CPU上计算和显存是7:1(随意的假设)时候, 大致内存和CPU性能都没有太大的短板. 一切都很好; 而如果你迁移到GPU上的时候, 因为GPU的计算峰值提升的幅度, 远远比显存的带宽提升的幅度大, 这个比值可能变成了77:1，此时原本的各方面均衡的代码, 往往会优先的主要体现在访存上, 出现短板. 因此就需要优化。而且特别的, 如果我们之前说的, 因为CPU上有各级较大容量的缓存在为你服务着, 而GPU上(相比来说)往往没有甚至非常缺乏, 此时就算是提升的幅度比一样, 你也需要考虑优化访存了。
(3) 随着老黄历代的GPU的性能发展, GPU本身的前后代对比(而不是GPU比CPU), 依然会呈现一个计算峰值的提升比, 会远远超过显存带宽的提升比的问题. 也就是我们常说的, 随着一代一代的新卡的问世, 计算性能将比访存性能提升的更快。
例如我们知道大约10年前就有100GB/s-200GB/s的显存的卡了. 例如当年我们使用的经典老卡GTX480. 而今天手册的例子上的V100才不过到890多GB/s, 这是将近过去了一个10年了，只有几倍的提升，而考虑计算峰值, 当年480的1TFLOPS的卡, 而现在的卡不考虑TensorCore也有几十TFlops了(最新的30系列的卡),这就几十倍了. 算上TensorCore那就计算性能提升的更惊人了.
你会看到, 因为这种一代代的GPU卡, 计算性能总是更快速的提升的现象. 就自动会让GPU上的代码, 包括哪些从头就是为GPU设计和实现的代码, 本身会随着时间的发展, 自动趋向于卡访存的情况.
因此大部分的代码, 应当首先考虑本手册说的, 对显存的特性的了解和优化. (GPU上原生设计的代码本身如此, 更不用说那些从CPU迁移过来的代码了)。
这三点可能是主要的为何本手册上去先说这个的原因。
GPU显存的粒度
我们先看一下具体的, 一些使用显存所需要注意的事项, 以及, 一些简单的 技巧, 来提高显存的性能的使用和发挥.
我们回到之前的话题, 之前主要说了GPU显存的特点, 即高带宽和高延迟, 以及, 在GPU上面应用的主要容易遇到的矛盾或者说短板是显存带宽, 而往往不是计算。
这里还需要说一下, 显存的访问, 是具有最低的粒度要求的. 这个粒度很多时候是32B (但HBM的系列可能除外, 之前看到的第三方资料—-我没有使用过HBM的N卡, 说是64B),
粒度(granularity)实际上是指, 你访问了某段地址的哪怕1B, GPU也要至少传输32B的问题. 从已有的资料和历年的GTC的PDF和PPT来看, 这个32B, 指的是从显存(DRAM)传输到L2, 和L2传输到L1的粒度. 注意我们现在都是Pascal+的卡了, 以前的老卡还有L2->L1的overfetch的问题, 现在也没有了, 就是32B. 虽然32B很小, 但是不当的访存, 例如本手册上的例子, 大跨步(strided)的访问4B或者1B, 会导致周围的连续32B都会被读取的. 这样会严重的浪费有效带宽的. 然后手册上还说了, 除了读取的这种情况下, 对带宽的浪费, 在有ECC存在的情况下(依然先不考虑HBM显存的卡, 那个很特殊),
也会导致浪费. 为何? 因为ECC的不是对1B为单位的显存中存储的内容的保护和纠错, 而是对一个计算单位的. 对一个计算单位整体进行ECC编码的计算, 然后附加的跟随写入. 这样就导致了两个问题:
(1) 普通的无ECC的显存, 可以通过掩码的方式, 例如我可以直接传输一定大小的内容, 假设是4B of 32B, 其中将这4B(32-bit的)范围, 设定有效写入, 其他28B, 设定忽略. 但是ECC的做不到这点, 他必须将整个一定的纠错基本计算单位粒度的那么大小的，读取到L2，然后L2就地修改其中的1个或者多个字节, 才能重新计算ECC. 这样就导致了, 在一定传输和ECC粒度存在的情况下, 写入会导致读取的现象存在.
(2)不仅仅如此, 伴随写入的ECC编码, 是在这个基本计算粒度的后面跟随的, 这样哪怕你写入了1B这么小, 实际光额外追加的ECC编码就很大了(例如8B, 具体得看相应的某代的GPU的白皮书或者其他资料才能知道)。所以这个粒度的问题, 存在于对显存的读取, 更加也存在于有ECC的情况下, 对显存的写入. 反而是无ECC的显卡, 纯写入的时候可能受影响的要小一点.
更不用说本章节手册提到过的, 一些GDDR5/GDDR6显存的卡的ECC的编码还要跟随占用传输带宽和存储容量的问题(这个影响读取和写入, 本章节手册已经说过了)。
好了. 继续回到GPU显存的传输粒度这个话题上.
虽然看起来这样很坑, 但是本手册同时提到了, GPU很适合这种零散的随机性的大范围跳转的读取, 相比CPU依然具有很大的优势, 这听起来很反直觉, 是为何呢? 这是因为CPU的内存, 同样具有这个特点, 也有它的粒度. 而且这个粒度比GPU的还大(依然暂时不考虑HBM显存的情况). 依然应该是后者。
前者存在CPU同样的过大的cache大小, 和overfetch的情况, 也存在每个CPU核心, 只能应付少量的outstanding进行中的cache miss后的读取. 而后者, 作为集成的GPU, 具有GPU的特别, 可以同时有较大数量的从SM到L2的进行中的请求, 来等待取得. 从而可能有更好的性能. 我们之前在图灵系列实体显卡的测试中, 大范围的4B为单位的读取, 能提供有效的带宽是大约在显存的4%左右. 或者以32B传输为单位, 大约在30%+这个数值, 和今天手册上给出的图, 并不一致. 我们将有机会重新在新一代的卡(8.6的安培), 或者我们的Jetson产品上给出重新测试. 以便让用户知道毫无规则的访存的时候, GPU依然具有的, 比CPU好很多的优势. 这点切记. (因为实际点的代码总是可能会涉及到一个大查找表的, 此时的下标可能毫无规律)。
Okay. 我们已经理解了GPU显存的特点, 高带宽, 高延迟, 最小粒度(其中这点是和CPU共有的). 以及, 刚才上面说了, 大部分人的主要短板可能是在访存上. 那么我们下篇就继续进一步的, 说说Global Memory的相关优化.
GPU卡和Jetson上显存优化的特色
我们下面就继续进一步的, 说说Global Memory的相关优化.
要说对它的优化, 我们得先知道Global Memory是什么, 和很多人的印象里的不同, 它不一定是显存. 还可能是映射的内存. (例如zero-copy时候的手工分配的, 和退化的Unified Memory的情况). 我们主要说一下当Global Memory是显存, 和是zero-copy的情况, 而暂时忽略是退化的Unified Memory的情况。
说一下这两者时候的注意事项和优化.
首先本实践手册这里, 提到了zero-copy内存, 这种是锁定在物理页面中, 而不能被交换到磁盘上, 同时又能被GPU设备, 直接访问到的内存(映射成了global memory)。
这种内存, 具有多个特点.
先说一下实体显卡上的, 再说一下我们的Jetson上的。
在实体显卡上, 因为卡是通过PCI-E连接到CPU的, 此时的一切传输(从内存到显存), 均需要通过PCI-E进行. 而PCI-E的带宽非常有限. 通常只有16GB/s的理论带宽. 注意手册这里第一次给出了实际能达到的, 在PCI-E 3.0 x16下的传输带宽, 往往是只有12GB/s左右, 这个结果是和实际日常使用中的情况是一样的. 手册这里给出的的确是比较准确的数字. 而这个带宽, 很多时候, 只有通过锁定在页面中的缓冲区传输, 才能达到的, 为何? 主要是这个涉及了CUDA C编程指南中, 和CUDA Runtime API手册中的最前面说的一些情况. 我简单这里重复一下吧.
一个是作为PCI-E上的GPU设备, 他用自带的DMA引擎, 进行BUS Mastering的时候, 本身就只能访问物理页面(或者说, 物理页面范围). 而不能支持CPU端的虚拟内存. (否则它将需要理解CPU端的页表结构等一系列问题, 才能自动转换)。我们知道当年的某蓝色巨人的XGA显卡, 作为它的8514显卡的继承人, 具有CPU端的虚拟内存支持, 并重新实现了兼容于x86 CPU的页表和相应的CR寄存器，作为当年的该机型的用户, 超越了20年的感动依然在心中。然而其他所有后来的显卡都没有这个特点, 因此他们就只能用自己的DMA Engines访问物理内存(好吧, 某些DGX例外). 因此, 当用户要求传输一段普通的可换页内存的时候, 要么显卡驱动内部, 先将该段内存缓冲区的内容, 复制到自己内部的一个小的锁定的物理页面范围上去, 然后再从这里安全的传输; 要么就是就地尝试锁定, 然后传输. 但前者多了一次内存<->内存的传输, 后者则有锁定和解锁开销(这里面的细节可以看CUDA Runtime API手册的前面, 历年GTC也有过详细描述) 。
前者几乎将内存的有效带宽降低了一半，例如我们上次距离的那个68GB/s峰值的内存的机器(4通道DDR4-2133), 这样一倒腾, 内存带宽就实际上只有34GB/s了. 哪怕你不考虑在CPU上运行的应用程序的需要, 光这点带宽, 传输两三张卡就撑不住了，而使用这种锁定了页面的内存, 则可以就地开始传输, 节省了一半的内存的带宽, 因为过程从内存->内存->显存, 变成了内存->显存了. 还是很容易喂饱你的卡们的. 这个是很适合在实体GPU上传输的特性。
这是对于实体卡说的，对于我们的Jetson产品, 实际上是并没有独立的显存和内存的, 也不是通过PCI-E总线传输的, CPU和GPU都在SoC的内部, 共享SoC提供的内存(显存)控制器. 此时如果你照搬之前的经验, 直接来一个cudaMemcpy*()系列函数. 则实际上你无辜的在该内存(显存)内部, 倒腾了一次. 无任何意义. 当时他们提出, 一定要利用页面锁定内存, 能无传输的就地访问的特性(俗称zero-copy, 可以看成是最最简化版的Unified Memory). 取消掉这个在Jetson上的无辜传输. 然后让GPU就地访问CPU的工作数据. 这点还是非常重要的。
不过在Jetson系列产品上, 直接使用zero-copy的方式会禁用GPU的缓存(L2)的, 不如后来的更好用一点(但限制更大)的Jetson上的Unified Memory方式好(也是直接共享, 但启用L2, 但会禁用CPU-GPU同时访问). 这点等我们到了Unified Memory的时候再说。
好了．回到本章节前面的页面锁定内存的传输和作为zero-copy上的特性内容。
注意本章节这里提到了, 要在CUDA 2.2+, 无显存的那种集成显卡上, 使用zero-copy, 云云的. 而没有提到Jetson. 这是因为手册10年来没改的缘故了.
我们现在已经将近10年买不到这种无显存的卡了(以前一些笔记本上有)，而Jetson产品也已经普及。
所以我们这里做了调整, 取消了手册上说的过时的不存在的内容, 而增加了Jetson上应当考虑zero-copy就地使用的做法.
此外, 手册上介绍了如何能利用页面锁定的内存, 有效的进行计算和传输重叠的特性. 手册这里给出了一个很好用的东西:
就是当我们只有1个kernel要启动, 和1份缓冲区要使用的时候, 如果能让这个进行多次传输和计算的重叠. 本实践手册这里给出了如下建议: 即将你的kernel修改里面的线程或者block的坐标/下标映射, 将原本一次启动的kernel, 工作于一个大缓冲区上, 改成N次启动(注意每次里面的坐标的变化), 并每次传输1/N内容的缓冲区, 这样可以尽量达到传输和计算的重叠. 这个是一个很好的实践方式. 因为很多时候, 你想利用传输和计算重叠这个优化, 但是你找不到多余的kernel计算任务, 和多余的传输任务来重叠，此时, 你应当考虑本手册中的, 拆分计算规模, 和拆分成1/N每次传输的建议. 这种建议还是很好的. 当然这样做, 你需要注意坐标的偏移和变化. 不要计算错了.
实际上, 在某OpenCL对等的规范中, 在友商家的卡上, 我们可以直接在启动kernel的时候, 提供这个坐标偏移量. 从而完成类似的操作, 而不需怎么改动代码. 所以你看, 这是一个相当实用的东西, 实用到友商家已经提供了API以方便你这样做了.
此外, 手册还提供了另外一个建议, 就是直接不要传输了, 直接就地使用. 因为zero-copy可以以一定的粒度, 直接从内存跨越PCI-E到L2(然后进一步到SM中). 这样根据手册这里的说法, 能在kernel的指令级别, 实现计算和传输的重叠. 但是根据我们的实践, 大部分的使用效果实际上并不好. 我们怀疑是可能是跨越PCI-E带来了更大更难掩盖的延迟或者其他因素, 这些需要等待确定. 但是手册中不强调的另外一个用法, 实际上效果非常好. 即将zero-copy作为容纳结果数据的, 写入的缓冲区. 这种可以将结果的回传, 和kernel的计算, 也在kernel的指令级别, 进行重叠. 而无需你事后开一个异步传输在某流中. 在我们的日常实践中, 这种具有非常好的使用效果. 好到了很多人都发现了这点. 并且在arvix上发文, 用FPGA拦截了zero-copy作为写入结果的缓冲区的时候, kernel的写入, 在指令级别的重叠, 通过L2, 跨越PCI-E回传的时候的情况, 并做了分析. 我们现在大致知道, 根据此文(我不记得番号了, 但是容易找到. 用FPGA + Zero-Copy的字样, 近期文章),
手册中的这种作为直接回写操作的做法, L2会产生跨越PCI-E的32B, 64B, 128B大小的传输的.
然后本文章还分析了, 为何只能达到约12GB/s, out of 16GB/s的原因(因为PCI-E的包大小的问题的浪费, 而不是编码问题). (注意这是说的PCI-E 3.0, 而不是4.0, 后者可能大约达到25GB/s, 这是根据我们客户的反馈, 而不是实际我们的测试).
大致是分析完了显存的特性、一个令人意外的(非常分散的不合并读写的优势)、 两种global memory中的前一种, 以及它的传输上的好用的地方, 和在Jetson上的应用。
一些规避的坑和优化的要点
我们的CPU在读取的时候, 从它的内存读取到它的L2的时候(L3或者L4, 作为LLC, 很多时候是victim cache, 也就是读取的时候不经过, 只有被淘汰的数据才尽最大挽留的存放, 所以这里不提), 粒度往往是至少64B的，这样, 同样零散的分布的读取1B的数据, GPU效率是1/32, 而CPU可能只有1/64. 更加可怕的是, CPU往往会对邻近的cache块/行, 进行预读, 和预测性的预读 实际上很可能会导致, 读取1B, 传输了上百B甚至更多的情况, 此时从效率来说, GPU的1/32要远远超过了CPU. 更何况, 这个是从效率上的说法, 实际能有效提供的带宽, 要用效率乘以各自的峰值, 显存具有大得多的峰值, 此时再乘以更高的效率，就得到了在这种严重不适合GPU, 也不适合CPU的情况下, GPU的性能依然要更好的情况出现. 这点很多书上往往进行了忽略. 因为这些书教育我们, 一定要使用合并性的访问, 要使用适合GPU的访问. 从而导致了很多用户, 不敢将这种不适合GPU的访存, 进行CUDA化改写, 这是很错误的。
本实践手册的这个章节, 破除了这个迷信思想, 还是需要的. 特别的, 在我们的jetson产品上, 存储器的体系结构(hierarchy), 是缺乏一个主芯片级别的统一最后一级缓存的, 即所有的数据, 都最终要通过存储器(LPDDR4), 才能得到一致. 哪怕此时问题来说, 同样的一个渣代码, 无论用CPU还是迁移到GPU上, 访存都是很零散的, 用户你究竟是准备用自带的ARM CPU核心来读取呢? 还是准备用集成的GPU部分来读取呢?
前面的话题已经说了, 如何在Global Memory做, 以尽量取得较好的性能优势. 以及, 和传输相关的方面的话题. 但是有一点没有说, 就是8.0+计算能力新引入的, 将部分Global memory中的缓冲区, 形成一个较长时间片段内, 锁定在L2 Cache中的效果. 或者用户可以理解成, 在一定的时间范围内, 将L2的某部分设定成尽量类似L1之于shared那样的, 类似手工管理, 或者说缓慢淘汰的效果. 这个不说是因为我们还没有测试, 同时, 我们所有在售的Jetson产品都不支持这个特性. 我们可能在最后的时候, 在8.0+上进行测试, 然后重新说这个话题。
好了. 先进行今天的内容. 今天的内容是如何尽量发挥shared memory的性能. 这个其实也是老生常谈了. 要发挥shared memory的性能, 我们得知道为何我们要用shared memory, 为何它的性能是在某些特定的kernel中, 是性能影响因素，因为你既然读到这里, 如果你的kernel本身不卡在shared memory性能上, 甚至根本连shared memory都不会用到, 则自然你不继续看了, 如果你需要继续看, 则至少你已经用了shared, 或者想用, 并且想解决使用中的性能瓶颈, 或者提前避开一些坑. 所以我们就说点这些。
如同本实践手册所说, shared memory在某种意义上, 等于是手工管理的L1 cache. 这种说法, 对于来自CPU的用户来说, 听起来还是比较有吸引力的.
因为一个传统的L1 cache你只能被动的使用它, 并且预估自己的那些访存模式, 适合被L1缓冲, 从而尽量的去好好使用. 而Shared Memory作为完全用户管理的东西, 你有充分的自由可以随意使用, 任何情况下都不会像L1那样, 数据存在自动淘汰可能, 总是可以安全的存储, 高速的使用的.
但是我们作为GPU, 一个追求吞吐率的设备(上次说过的), 很多时候用户们追求近乎100%的压榨出来上面的某些单元的性能, shared也不例外。
今天就大致说了一下, 哪些是影响因素, 并再次(再N次)的给出了使用shared memory进行分块矩阵乘法和转置的例子, 用来显出使用了shared后的高速度来。
我们直接说一下一些规避的坑,和优化的要点:
第一点则是, 尽量规避shared memory上的bank conflict. 这个也是老生常谈了. 我们现在用的, 能买到的新卡, 都是4B宽的Bank. 每个Bank用户应当理解成在每个周期内, 能独立给出4B数据的独立单元，这样每个SM里面, 如果有32个Banks的话, 能给出128B/周期的性能. 这个还是很惊人的，因为对于从CPU迁移过来的老代码来说, 自家的L1 cache, 也不过常见每个周期能给出2个32B读取, 和1个32B写入这种. 也就是96B/周期. 但是CPU的核心数才多少, GPU的SM数量又多少。
一个动辄80多个SM的GPU, shared能聚合给出10TB+到20TB+/s的性能(假设频率从1Ghz~2Ghz的GPU主频). 所以很多老代码, 进行了优化, 迁移到GPU后, 第一步就是考虑尽量利用shared的这个高速特性. 从而发挥性能. 然而, 这个高速度只是理想状态, 一旦shared发生了bank conflict后, 性能会下降的. 下降的程度和你bank conflict的程度有关系. 而具体bank conflict是什么, 我们这里不讲. 因为实在是讲的太多太多次了(几十次是有了). 感兴趣的用户可以回看我们的编程手册内容, 或者回看Sisiy的阿三书. 里面都扯了好多好多。
这里主要说的一点是, 在近期的NV给出的资料中, 揭露了一个新的现象.
就是我们以前一直说Bank Conflict的时候,根据手册,都是用的warp整体(在现在你能买到的卡上), 作为bank conflict分析的, 也就是32个线程内部之间的有无bank冲突. 从而尝试优化. 但是这种手册上给出的分析方法, 和实际的使用中的profiler给出的conflict的报告, 和实际因为达到的性能, 很多时候是理论和实际结果不符合的。很多情况下, profiler给出的bank conflict数量要少很多, 性能指标也要好很多.
例如本论坛的这个例子:
https://bbs.gpuworld.cn/index.php?topic=73410.0
该例子的楼主们, 以及, 奈奈同学, 给出了自己观察到的不同于手册说明的现象. 并且进一步的挖掘出来了, NV只在GTC上给出的一个PDF资料. 该资料里有不同于手册的说法: 即: 在8B, 16B的这种非4B的访问情况下, 也就是类似float2, float4, double, double2这种访问的情况下, bank conflict的计算不是按照warp进行的, 而是分别实际上按照half-warp和1/4 warp进行的. 这点符合实际实践中的profiler的报告的性能结果. 我们今天在这里额外的从论坛揪出这个案例, 同时用NV的这个资料, 进行说明:
在特定的访存方式下, bank conflict的计算应当采用另外的范围(即1/4或者1/2的warp), 而不是从warp整体. 当读者或者用户正好使用这种访存方式的时候, 无需过度的去考虑优化Bank Conflict的问题, 因为很可能此时conflict根本就不存在.
这点需要注意了. 此外, 我们还想给出一点说明的是, 有的时候, 将shared memory作为一个高速的查找表的时候(参考我们之前编程指南手册说过的, shared memory的三大用途之一), 如果下标高度规律性的一致, 在warp内或者block内部如此, 则编译器可能会生成另外一种带有LDS.U后缀的shared读取指令, 会让实际的读取的延迟降低很多, 等效吞吐率提升很多. 该现象很容易发现, 也不报告任何的bank conflict. 但是我们目前还不知道为何会这样, 以及, 如何能让编译器触发这点. 这里的给出只是用来说明, 很多时候本实践手册中的conflict方面的相关优化并不成立, 用户应当以实际的应用中的性能分析器对相关单元的指标报告为准. 然后手册今天不出乎意料的, 继续引入了矩阵乘法/转置的内容, 用来说明shared memory在重复使用数据, 和转换不适合的低效global memory的访存为适合的高效的shared memory访问的特点.
重复使用数据就不用说了, 既然shared memory作为手工管理的L1 cache, 他自然也有cache的这种提供缓冲和高速性能的特点; 而转换不当的访存模式(例如常见的纵向坐标优先或者说大跨步等的方式), 经过shared中转了一次, 变成了恰当的模式, 则用户应该看一下. 后者是很多用户容易忽略的, 特别是对于一些案例, 数据明明只需要使用1次, 那么为何我还需要先读到shared memory中缓冲一下, 然后再从shared memory读取一下呢? 因为对于很多这种的, 哪怕你只用一次, 经过shared memory这么一倒腾, 就可以让访存模式理顺很多, 哪怕只用一次, 也是有性能优势的. 而这种优势, 在直接使用L1 — 不具有不同深度的bank的同时数据供应 — 是做不到的. 但是shared可以。
一些规避的坑和优化的要点(续）
我们需要指出的是, 任何用户现在均不应当手工去尝试进行shared memory上的”优化”, 从而能让自己的”矩阵乘法”变得更快. 因为NV自带的cublas库已经在那里了. 该库超越了大部分人的写作水平, 也包括在给你读这些的今天的我们.因为cublas可能会使用更加底层, 接近硬件的工具, 来书写, 而不是较高层次的CUDA C和PTX.任何用户都应该考虑直接使用该库. 这也是为何我们今天不说这个例子的第二个理由(第一个理由是之前说过太多次了).
越过这个例子之后, 我们再说一下常用的shared memory的, 作为warp和block内部的数据交换缓冲区这点. 这点还是对性能的发挥很重要的. 很多时候, 我们需要在协作的, warp内部的线程之间, block内部的warps之间, 进行数据交换. 来完成特定的算法实现的要求. 此时用shared非常好. 而两个常见的来自论坛的用户的其他做法, 则是不推荐的:
第一个做法是直接分配global memory上的普通缓冲区, 然后每个线程都写入自己指定的位置, 用这个来交流数据. 这个做法还是很不好的. 延迟大, 带宽低. 而且每个block中的每个warp中的每个线程, 都需要计算自己的下标, 就像普通的global中的缓冲区那样.
而使用了shared来作为内部的数据交换, 则具有延迟低, 带宽高, 以及, 更美妙的是, 每个block都有自己的同名shared中的副本, 同样的下标在不同的block中, 自然的被分开了, 从而导致你能用简单的下标来完成交换, 而无需计算完成的, 全局独立的下标. 下标是完全可以在block间的级别重复,而不发生冲突的.
第二个做法则是尝试使用local memory进行数据交换. 这个是错误的, 我们已经在论坛调试了无数这样的代码了. 这主要是针对来自CPU的用户. 甚至是一些高级的CPU用户. 他们本能的认为, 我线程还在活着, 我定义了局部数据, 哪怕是在某些类似stack上的东西, 我也可以临时性的将指向其中的指针, 给其他人(伙伴线程)使用, 只要我存在, stack上的东西就有效.
这点在CPU来说, 的确是的, 而且是在CPU上的多线程数据交换的时候的, 在每个线程都活着的时候, 一个危险而好用的技巧. (因为它规避了小缓冲区动态分配和释放, 直接一个指向stack上的临时性内容的指针就可以了) ，然而, 这点在GPU上并不成立. 在GPU上, 这样将会导致实际上为每个线程分配一种叫local memory的东西, 你的指针指向local memory的内容, 只会在本线程内部, 该指针有效.
如果你用一些技巧获取指针, 传递给另外的线程, 你会发现指针能用还是能用, 但是指向的内容是其他线程的对应指针位置的内容, 这点就非常有意思而难以debug了(其实不难, 因为你载坑一次下次就知道了).
此时依然推荐使用shared进行高效的数据交换. 这点需要注意了. 我们不想在论坛继续解决这种问题了. 我们最后补充2点, shared memory的用途, 是本章节手册没有强调的.
一点则是作为查找表. 当查找表的规模适合的时候, 也就是从128B+到几十个KB的时候, 应当考虑使用shared memory.
(过小的查找表, 例如小于128B, 你应当考虑放在warp的32个线程中的每人的1个寄存器中. 然后用shuffle进行下标查找. 这是不容易出错的方式). (而过大的, 你shared也放不了. 此时可以部分放在shared, 其他部分或者整体考虑放在global中, 或者使用其他策略)
放在shared中这点, 看起来是理所当然的. 但是实际上根据我们的经验, 大部分用户会反直觉的, 首先考虑将查找表放入constant中. 这是非常不对的. 因为后者不不能很好的支持无规律的下标访问. 等到constant的时候我们再说.
只是因为它带有”常数”的这种名字, 很多人就用它来顾名思义的进行类似常数/系数查找之类的用途, 这是不对的. 依然这里用shared才会有较好的性能.
第二点则是, 某些kernel无可避免的需要进行某种类似结果compact的操作, 或者说, 先不能一次性的生成最终结果, 而是先生成一个接近最终结果的半成品, 然后最后才能有效的排除/筛掉一部分. 这种操作往往是因为半成品本身, 需要根据前后的值进行进一步的运算, 才能去掉某些结果, 在图像处理中很常见. 此时就有多种选择了.
一种是再开一个kernel, 从global中处理. 另外一种则是回传CPU, 进行这种filtering. 此时完全可以将临时结果先写入到shared中, 然后再从shared中进行筛选. 这个是比在后续的kernel于GPU上, 或者回传后筛选都是高效的. 至少它降低了传输的大小, 我们已经知道了之前的内容说过, PCI-E传输是比较慢的, 应当尽量优化它. 这种将结果分两次写入, 第一次写入shared, 第二次再从shared中写入目标global memory中的做法, 听起来也很简单. 但是根据我司这10年来的经验, 很多用户是会自动无视这点, 或者说自动忘记这点的, 为何? 因为大部分的CUDA书都在教你, 如何有效的用shared缓冲/重用输入数据, 而几乎从来不提, 对结果的写入也可以中途经过shared stage一下,从而完成过滤/压实之类的操作.
从而让读这些书多的人形成了思维定势, 自动对结果写入使用shared这点进行了忽略, 这点还是应该要注意的. 论坛好多人已经这样思维定势了.
最后关于shared memory的则是, 从计算能力5.0开始, shared memory本身具有一定的计算能力. 例如shared本身可以比较高效的计算加法(例如你在atomicAdd的时候). 而5.0之前都是使用了类似锁定—SP加法—解锁的策略, 现在直接是让对应的存储单元(shared)完成这个计算操作了. 因为5.0+的卡是我们现在能买到的唯一的卡, 一些基本的计算, 有的时候可以让shared去完成.
例如我们之前老生常谈N次的纯粹在SP中完成的某些规约操作. 包括图像处理中的, 很经典的, 追加一个List. 可以用shared上的原子操作在shared上尝试构造一个小的. 然后整体在追加到global memory中. 这样可以某种程度的完成行/块级别的内部有序. 也能降低global memory上的, L2级别的原子操作的压力. 特别是当相当多的SM中的请求, 密集的对1个L2上的4B索引位置要求进行原子操作的时候. 因为根据我们的经验, L2上的原子操作可以可以很密集, 但一定要错开, 就想是存在某种类似shared meomry上的分片或者bank机制那样. 如果你不想去研究探讨L2上的这个机制(没错, 我们看到的L2实际上是很多独立的小L2聚合而成的), 则先在各自SM内部的缓冲区上拼接构造, 然后一次性的完成1次, 而不是等效的几十次到几百次的L2原子操作请求, 并最终写入. 还是很好的.
此外, 从目前我们能买到的新卡(例如RTX3070), 已经支持直接从global memory读取到shared memory了. 这是一个极好的特性. 是从友商AMD那里学来的特性.
从Global memory到Shared memory
上一篇里我们说到目前我们能买到的新卡(例如RTX3070), 已经支持直接从global memory读取到shared memory了. 这是一个极好的特性. 是从友商AMD那里学来的特性。我们稍微解释一下。
从大约10年前的GCN的A卡开始, A卡具有一个独家特性, 可以直接从global中加载到LDS中(相当于shared memory), 这样做有很多好处, 例如可以实现异步效果, 可以让某block在请求后台的global->shared的传输中, 主体逻辑在做一些准备或者初始化操作. 而不需要像以前那样, 必须先读到寄存器, 然后从寄存器写入到shared. 读取到寄存器本身无问题, 反正寄存器的占用只是临时的，但会导致主体逻辑卡住, 在主体逻辑一旦试图从寄存器访问到未就绪的数据的时候. 虽说上次内容, 我们都知道, 可以依靠切换warp, 让SM执行其他没有卡住的warp中的内容,但是实际上你在用老nvprof/nvvp或者新的nsight compute的时候, 在选择了PC Sampling的时候, 能看到具体的,往往前面这种一碰载入到寄存器的初始用的数据, 就卡住的情况还是家常便饭的(会显示一个很长的long scoreboard等待采样计数, 类似的东西).
而长期以来，A家则提供了异步的载入, 同时还提供了查询和等待/同步操作, 能让主体逻辑去查询, 后台的异步载入到shared memory进行到哪里了,或者在主体逻辑真的完成了所有前期工作后, 要开始使用shared了, 可以选择的进行一次等待/同步操作.幸运的是, 我们这10年来, 有两点终于得到了满足.
一点是NV终于现在提供了这个特性了, 而且异步载入指令, 选择性的计数等待(例如发出来3批传输, 等待到传输完第一批的时候), 和整体等待/同步等特性. 这样有效的降低了常见的每个warp头部的低效的”冷片段”, 有利于整体显卡性能的发挥. (所以说, 你买一张8.0+的卡还是值得的,)
二点则是, AMD从10年前引入这个特性(其实比NV做的还好, 因为还可以做轮询进度)后, 始终拒绝在自家的OpenCL中, 将该特性导出. 从而实质性的, 能让你买到支持的硬件, 但是就不让你用(有其他方式能用, 但是这里不提, 因为无关今日话题). 从而降低了NV这10年来在追赶优质硬件设计上的压力. 这是新的来自8.0+上的重要的shared memory上的特性和优化,应当注意. 如果你不喜欢现在新版本的C++风格的在CUDA C中的导出, 则你依然可以使用PTX中的传统C风格的调用方式, 手工导出特性即可. 注意对于16B读取(从每个线程的角度), 该新特性允许直接从Global memory中bypass掉L1, 直接送进shared, 避免了对L1中的内容的污染. 这点适合我们今天前面说的, 有些数据哪怕只读取1次, 但为何转换成合适的访存模型, 也可以考虑shared那条的用途. 或者其他的, 任何避免污染L1的情况的用途.
local memory你可能不知道的好处
下面我们简单的再说一下local memory.
首先要注意的是local memory并不local, 它实际上依然是一段显存. 但可能会被各级相关的缓存所缓冲。
主要用途有两点：
一点是你(读者)使用，当你需要每个线程的一段缓冲区的时候，你并不需要单独的开一个全局的大的缓冲区，然后作为参数传递给kernel, 让kernel里的每个线程找到自己对应的一部分使用。你可以直接来一个局部的大数组(不能过大!), 享受类似以前的CPU上的C风格的, stack上的定义的数组, 或者类似CPU上的alloca()的分配风格, 能自动的每人一份, 而且能自动释放, 很是方便，而且不仅仅如此, 你如果传递进来一个大缓冲区这样用, 你需要为所有的一次启动的线程分配缓冲区. 而用local memory, 则只需要保证能真正同时上到SM里执行的那些线程的数量所需要的缓冲区，举个例子说, 前者你启动了1M个线程, 每个线程需要1KB, 则你需要1GB的显存提前手工分配了.而如果你使用后者, 某GPU device实际上只能同时执行10K个这样的线程, 其他的暂时没上的, 在其他block中的线程们, 会等待下次轮批次再上, 则硬件上只需要准备/分配出来100MB的显存, 即可应付, 因为这些线程不是真的”同时”在运行中的(具体参考我们之前的编程指南手册).这点不仅仅降低了手工管理的成本, 还降低了你花钱买一张更大显存的卡的成本.特别的是在Jetson设备上, 显存(内存)容量有限, 用户应当考虑这点.但是很遗憾的, 很多人就是喜欢传递过来一个额外的数组/指针这样使用, 原因我们还未知.
此外, 使用local memory还有一个好处, 就是虽然它像global一样, 被各级缓存缓冲, 但是它有更精细的缓存控制策略, 可以允许对local memory上特定位置的访问, 标记成discard, 或者说last use(PTX手册用语). 允许cache直接将对应的cache line内容, 就地丢弃掉, 而无需必须回写下一级缓存甚至到显存. 这点作为global memory是做不到的。
此外, 今天的实践手册没有说明的是, local memory还具有强制合并访问的特性.我们都说用了local memory, 但是几乎没人讨论”local memory是否是合并的”, 既然我们今天已经知道了它也是用的显存模拟出来的, 为何不讨论这点?这是因为local memory有自动交错的特性. 例如我们定义了一个int dog[N]; 假设dog被编译器选择放置到了local memory上, warp中的每个线程都在访问同样下标的, 例如dog[K]的时候, 实际上来自32个线程的同样下标的访问会被合并成连续的地址空间上排布的一段128B的内容, 非常的合并.用户可以理解成local memory实际上总是按warp排布的, 任何int dog[N]都是内在的被存储为int _dog[N][32]这种自动交错.
从而总是自动形成了, 当下标一致的情况下, 自动合并的效果. (这点最早见于2013年的CUDA Handbook, 这是一本好书, 但是国内翻译的书质量不高，所以我们一直没推荐。也可以参考我们之前的CUDA编程指南中的内容)
因为这种自动交错/合并的存在. 对local memory中, 来自同一个warp的杂乱的下标/指针访问这种, 应当避免. 因为默认是一致的. 杂乱的访问会导致访存被拆分成多次请求, 严重降低效率.这是local memory的用途一.用途二则是, 方便编译器安排一些无法有效的放入寄存器, 例如当前阶段寄存器资源用的太多了, 或者一些访存方式(例如对寄存器试图进行下标索引—-N卡不支持这种), 不能放入.
纹理存储优势
根据之前的内容, 你已经知道, 纹理可以提供免费的值变换, 和免费的坐标变换, 以及免费的越界处理, 以及, 更加优化的访存/缓存效果. 我们主要从这4点说开.
先说一下免费的值变换. 有些算法需要将数据作为8-bit或者16-bit整数存储, 然后读取到后, 再转换为float之类的浮点数, 和其他类型进行运算. 而这个转换过程, 需要用户手工写, 哪怕是一个简单的float b = (float)a;这种. 以及, 这种转换还需要占用SFU(特殊功能单元), 注意SFU在新版本的Nsight profiler中已经简单的改名成了XU单元了. 那么此时, 无论是从转换指令本身, 需要占据额外的硬件资源; 还是从编写代码的人的角度, 他需要手写额外的代码行, 都是一种开销. 而纹理读取的时候, 可以利用上其数据路径中的自带的转换功能, 从而节省掉对SFU/XU或者人工编码成本的开销.
这样有可能带来额外的性能提升, 和对人力成本的节省.
例如我们知道, 在很多代卡的架构上, 一次SFU完成的整数到float的转换, 性能只有常规指令的1/4:
如图, 我们可以看到了7.x的卡上, 每SM每周期可以执行64条常规的float加法/乘法/乘加, 这往往构成了你的代码的运算主体;
而从8-bit或者16-bit或者其他整数类型转换成float的时候, 吞吐率就只有16条/SM/周期了, 相当于在7.X上转换本身只有常规计算的1/4的性能. 甚至这点在8.6上更加糟糕, 因为8.6的双倍速的float运算, 导致如果你读取一个普通的8-bit或者16-bit整数(u)int8/16_t, 然后进行一次手工到float的转换, 相当于大约等效8条后续的正常计算的性能被浪费掉了(某种意义上), 即转换只有1/8的效率. 此时如果你的代码SFU/XU是瓶颈, 或者因为使用SFU而导致了浪费了指令发射能力的话, 应当考虑使用texture自带的免费转换功能, 来节省对应的SFU的I2F之类的转换指令. 这样会可能带来额外的性能提升.
不过需要注意的是, 自动的转换是一个”归一化”的过程, 将会从8-bit或者16-bit的有/无符号整数范围映射到[-1.0f, 1.0f]或者[0.0f, 1.0f], 其中包括了1.0f了, 这点使用的时候应当小心. 例如考虑是等效乘以了1/255还是1/256的系数的问题(包括还是不包括1.0f右边界).
好在大部分的使用float运算的代码, 应当很容易处理这种问题. 这是使用texture的带来的可能的第一个优化上的效果.
注意第一点的值变换除了归一化读取到的值, 还有低精度的插值效果, 这个线性插值效果我们曾经已经在编程指南手册中说过了, 这里就重点说了. (虽然本手册这里强调了一下). 如果适用你的算法, 则利用硬件自动的插值的效果可以进一步节省你的手工运算量, 从而潜在的可能提升性能.
这两点都属于今天的texture带来的4点中的第一大点, 即自动/免费对读取到的值变换的好处.
第二点的好处是, 带来了自动的免费坐标变换, 即所谓归一化的坐标. 这点什么时候有好处?
例如图像处理或者神经网络的输入图像, 可以大小自动适配. 也就是说, 我一个256x256的图片, 和一个512x512的图片, 使用了自动的免费坐标归一化功能后, 后者和前者可以自动的等效缩放. 这点节省了用户单独的写一个kernel进行缩放的过程. 减少了工作量和出错可能, 也节省了一次kernel的代价.
当然, 现在用深度学习的用户可能不在乎这点, 也没法在乎, 因为他们如果使用框架的话, 能配置的只是简单的文本文件描述(例如对网络结构的描述). 不需要手写任何代码, 自然也不需要考虑这点. AI么, 会用记事本就能搞AI. 有数的.
但回到正题, 本章节说的坐标自动映射(或者等效的图像自动缩放功能), 的确节省了用户的开发成本. 此外, 和值变换不同的是, 这种坐标映射是右边界不包含的, 即一个图像(或者2D数组), 会被映射到[0.0, 1.0)的坐标范围, 手册这里的说法是, 映射到[0.0, 1.0 - 1/N], 注意)和]. 这样的映射在N是一定范围内的整数次方的时候, 或者说图像/2D数组宽度/高度是2的倍数的情况下, 可以在缩放的情况下, 依然精确表示坐标. 从而使得这个特性不仅仅适用于图像这类的数据, 也适用一定的需要严格坐标指定的普通2D数组/矩阵之类的算法/代码. 因为一定范围内的1/2^N在我们用的卡上, 是可以被精确表示的浮点数. (注意不是所有的浮点数/坐标都可以被精确表示). 这样texture就又带来了, 免费的而且一定情况下是精确的坐标变换/缩放功能. 使用它依然可以解放掉你的主代码去干其他事情. 从而可能带来无论是编程世间, 还是性能上的提升. 这是第二点.
此外, 我们往往不仅仅需要像(1)(2)点所说的那样, 无论对要读取的坐标进行变换, 还是要对读取到的值做进一步的变换处理, 在实际的2D数组/图像的读取中, 往往还需要考虑边界情况. 不考虑边界情况往往会代码你的代码行为异常, 或者出现无法预测的结果.
继续回到第三点. 我们看下纹理给我们带来的边界/越界处理都有什么好处/优势. 好处有两点:
第一点是, 在指定了一定的边界模式后, 越界不再需要考虑. 即节省了用户的代码编写工作量开销, 也消除了用户哪怕想付出努力/工作量, 却不小心遗漏导致出错的情况.
这点在今天的优化指南手册中, 正好给错过了重点.
我们知道之前在编程指南手册中, 我们和大约一起阅读过有4点边界/越界自动处理, 即自动填充0, 自动重复边界值, 卷绕和镜像模式.
而且我们当时还分别对这4种模式都画了图, 从而让你能够理解, 当时手册上只有文字描述的不好理解的尴尬. 但是今天的优化实践手册中, 只在表格中提到了后两者(卷绕/镜像). 但是实际上, 往往有用的是前两者。
我们已经无数次的在论坛上接到楼主们的求助, 诸如: “我需要在我的矩阵周围绕上一圈0, 应该怎么做” 类似这种的问题, 往往本着就问题回答问题的角度, 我们往往在论坛上给出的答案是: 重新申请一个宽度左右大2个元素, 高度也大2个元素的新矩阵/2D数组/图像, 然后将原始矩阵内容复制到中间, 然后周围一圈写入0. 或者我们给出的建议是, 每次读取都强制的走一个越界处理的code path, 即有效坐标正常读取, 越界/边界的部分, 直接范围0模拟一次读取到了一圈0的效果. 你看到, 无论是那一种, 都需要用户付出功夫. 而如果使用今天手册章节中说到的texture的自动边界/越界处理的话, 你可以免费. 我们具体说一下.
我们设定今天手册中没有说到的边界自动绕0模式, 此时, 就像论坛中很多人试图做的那样, 直接对一个纹理坐标进行读取(纹理中往往较拾取)即可, 如果没有越界, 和你的普通读取效果一样, 如果越界了, 自动返回0. 这样, 你不需要额外的处理或者if之类的判断语句, 效果却自动达到. 注意这不仅仅减少了你的编码工作量负担, 也减少了无论是多一个环绕0的kernel的执行成本, 或者是用if判断越界与否的处理的代码执行成本.
因为要知道, 绝大部分的代码, 都是要上1个或者多个if来对, x或者y坐标之类的进行有效范围判定的. 无论你是看老樊的书, 或者看阿三的书, 或者写过任何CUDA代码, 应当对这点都深有体会. 而今天, 如果纹理能适用你的数据类型/代码, 则你可以自动得到这个免费特性. 从而提升你的编码效率, 也提升了你的代码执行的性能. 这是说的第三大点的边界/越界处理中的自动返回0值的情况. 实际上, 在图像处理中, 往往还需要在边界超出的地方重复最后有效的值, 例如你在做某种梯度或者边缘之类的检测/处理之类的.
纹理读取也对这种提供了直接的免费边界/越界处理. 这就等效于你手工围绕上了一圈或者多圈边界值. 注意这个特性也很常用, 而且不用纹理用其他方式手工实现起来很麻烦. 麻烦主要在于你不知道边界需要涉及到越界出来多深(特别是对图像处理来说, 参考当年某维), 你可能需要围绕1圈, 2圈甚至更多圈, 而使用纹理的这个特性你可以免费绕上任意圈. 而没有成本. 另外则是这种需要考虑越界的方向, 往往需要考虑4个或者8个方向, 例如从左边界线, 上边界线, 右边界线, 下边界下, 或者上下左右4个顶点.
而今天, 你如果使用texture的第3大点的这种特性, 这一切都是免费的, if的多个分支可以被省略了, 从而潜在的可能提升性能. 而且主要是减少了代码编写者的成本, 和出错的可能.
注意对于第三点, 本优化实践手册说的是另外两点, 根据我们的经验和论坛上的各大楼主们的反馈, 另外两点并不常用. 我们这里就不说了, 如果你感兴趣可以看我们之前的编程指南手册, 里面说的很详细.
texture和surface
上一篇说的3大特性, 都等于在访存的同时, 还附加上一定的固定功能的运算/变换处理. 这种特性, 叫采样器特性(sampler). 而我们都知道, 采样器是在只读路径上的. 而去掉了采样器的texture在CUDA里叫做surface.
因为本优化实践手册编写的年代较早, 这里没有怎么提到surface. 我们简单的说法一下surface.
surface不具有刚才说的texture的采样器只读路径上的这些优势,但是surface具有额外的特性, 它可以写入, 而texture不能.除此之外, surface和texture还具有非采样器的另外的一个重要特性.这个重要特性是在多年前, 也包括最近一些年出的而没有动脑更新的书的经常重点强调的地方,即texture本身的缓存效果. 这是很多人至今还在坚持使用texture的重要因素. 我们来简单看下. 这个因素等于是在说存储本身, 而不是在对该存储的读取路径上的优势. 主要优势有两点, 一个是cache效果. 在某些卡上, 普通的读取不具有较好的缓存效果, 而texture读取有. 例如哪怕是到现在依然被CUDA 11.1所支持5.X硬件, 也是如此.
例如5.0的maxwell的卡, 对于普通的读取不能使用L1/read-only cache, 而texture和另外一种只读的读取方式(不维持一致性(NC)只读读取, 或者常见的__ldg()之类), 却可以充分利用. 此时使用texture或者surface读取, 就能获取此缓存效果上的优势了. 否则你的任何读取可能在此卡上都要走L2. 很亏. 这也是手册本章节说的, 具有带宽上的放大效果(注意, 本章节的其他内容这里不赞同, 因为手册可能很久没更改了, 例如手册说, 使用纹理和DRAM的直接读取具有一样的延迟啥的).另外的一种存储上的优势则是, 例如在使用cuda array的时候, 数据在显存中的排列本身, 可能是被重新排布过的,
constant和寄存器
我们继续说一下剩下的两点小内容.
第一个是内容是constant ，这个谁都知道是什么, 也知道它的优势, 例如我们都知道GPU是RISC架构,所有的数据都要单独的load操作进入寄存器后, 才能参与运算.，但是constant例外, 很多指令允许一个操作数, 作为constant的形式存在, 从而在一条指令内部, 聚合了该指令本身的计算功能, 外加对constant的读取在一起.
这里要注意两点:
(1)是7.5+的卡有单独的标量/Uniform路径, 不仅仅可以在SP的计算指令中, 集成对constant数据的读取为操作数, 从而节省了一条单独的load读取数据指令(例如常见的A = K * B + C; 这里的K就可能并不需要单独一条指令载入的). 7.5+(也就是图灵+)的卡, 其标量单元, 还可以单独在SP之外, 执行标量/constant载入指令, 进一步的提供灵活性和释放向量指令(例如可以在很提前的位置进行load, 而不是在遇到了c[Bank][offset]风格的cosntant操作数的后期时刻).
但是很遗憾的是, 目前的NV家的编译器还对标量路径代码生成支持的不好. 虽然我们知道这是竞争对手A家从10年前就有的功能(的确很多方面A卡硬件好), 而且A卡的配套软件质量非常渣, 但是这点上NV的编译器质量还是不如A卡的.
好在随着以后的CUDA Toolkit版本, 驱动版本的提升必然会逐渐的效果提升的. 总之读者现在该用constant就要用.
(2)点则是, 应当正确的使用constant, 这里的constant指的是手工放入constant中的内容. 我们在论坛常见很多楼主有很多错误/不当做法, 我举两个例子.
第一个例子是本优化实践手册这里说的, warp内的很多线程读取不同的constant数组中的元素, 这样做将完全失去constant的效果, 而且可能会起到反面作用(变慢).constant必须在warp一致的时候才能用. 其他使用将可能拖慢你的代码.这点论坛上已经有N个反面教材了.
另外一点是, 过度的使用constant, 表现为用户拼命的较近脑汁的将自己的代码中的常数, 例如1.0f, 233, 666这样的常数单独提出取出来,然后手工放入一个constant变量或者数组中. 从而认为这样可以”进一步的优化”.其实不是的. 首先说, 编译器会自动完成这个过程, 如果它认为某些数据能够从代码中自动被提取出来, 它会自动这样做, 并放入constant.
其次, constant并不是最快的, 有些数据如果合适, 可以直接嵌入在指令的内部, 作为”立即数”. 而立即数可以被保存在指令缓存(而不是任何数据缓存), 只要能取指令, 那么数据就已经免费就绪了. 所以用户手工的这样做(手工将kernel中的常数提取出来放入constant)是没有必要的, 甚至可能会起到反面的优化效果. 需要注意. 注意本实践手册是将其作为存储器分类的.一种是将寄存器作为存储器分类，一种是将其特化, 它就是寄存器, 而不将其作为通用意义上的存储器, 虽然也有register file(寄存器堆)之类的说法存在.
这里主要提到2个问题. 第一个问题是涉及到寄存器的bank conflict, 这点如同本优化指南说的,用户无法控制这个问题, 这个是编译器在生成目标代码的时候, 自动尽量规避的.这点我赞同. 同时本手册说了, 不用考虑用int4, float4, double2类似这种数据类型所可能带来的寄存器的bank conflict, 该用/不改用就用(不用). 这点可能是有点欲盖弥彰了.
因为在某代著名的3.5/3.7的时候(大Kepler), 压满显卡的峰值性能是如此的困难, 导致用户不得不考虑使用ILP(指令级别的线程内部的前后自我并行, 本优化指南后续章节会说). 而使用了ILP往往会导致使用int4/float4这种向量类型, 而根据已有的资料, 在大Kepler上这样做, 往往会导致严重的寄存器的bank conflict, 同时编译器竭尽全力还无法很好的避免, 这就很尴尬了. 所以手册虽然这里这样说了, 但是用户是否该用, 该如何用才是优化的, 请自行考虑.
好在现在随着时代的发展, K80这种卡已经逐渐的消失了. 再可预见的将来我们应当不太用担心这个问题了.毕竟, 如同人不能同时两次跨入同一条河流, NV总不能在同一个地方(指坑爹的Kepler架构)栽倒两次的. 这是第一点关于寄存器要说的.
第二点关于寄存器要说的则是, 很多代码, 并非使用寄存器越少越好, 也并非使用寄存器越多越好. 其寄存器的使用有个最佳点(甜点). 而这个甜点的值是无法确定的(和具体的kernel, 卡, 以及kernel和kernel间的组合情况有关). 我们前几天在老樊的群里看到有用户一本正经的讨论,我将寄存器从XXX个降低到了YYY个, 结果性能并没有提升, 为何(@#(@(!这个其实很正常。所以我们这里提出尽量可以考虑自动化的尝试寄存器的最佳使用点, 例如写一个脚本自动控制寄存器的用量, 用不同的用量值自动重新编译和运行评估代码, 从而能自动发现这个甜点，而不是用户自己(就像老樊的群里那样)去反复尝试, 费时费力, 可能还找不到这个最有点.
不改变代码本身如何提升性能？
因为GPU的SM是海量超线程的, 远比常见的CPU的一个物理核心的2个或者4个超线程(HT)要多的多, GPU依靠这种海量的超线程数量来提供最大可能的吞吐率(这点我们稍后说)。而这种海量的超线程, 当你每个线程的寄存器资源用的比较多的时候, 则SM上能同时驻留的线程数量就越少, 从而影响了GPU的这种海量的超线程的能力, 从而潜在的可能影响了性能的发挥，所以有的时候我们不能肆无忌惮的使用寄存器资源，而是需要通过某种手段去限制编译器生成的代码中, 对寄存器的具体使用数量。
在日常的应用中, 不改变代码本身, 而是简单的改变每个线程的寄存器资源使用数量(变多或者变少), 就有可能提升性能，所以这是一种常见的优化方式, 具体到今天的手册章节, 手册提出了两种做法:
一种做法是编译的时候, 对每个具体的.cu的CUDA源代码文件, 使用nvcc -maxrregcount=N的参数来编译。这种做法将会把此文件中的所有的kernel, 都统一限定成最多使用N个寄存器。
注意这里有需要注意的地方, 首先是这种限制是以源代码文件为单位生效的, 如果你文件中存在不止一个kernel, 则所有的kernel的限制都是一样的, 你有的时候可能不得不拆分源代码成多个文件, 从而使得每个文件里面只有1个kernel, 从而能单独的用-maxrregcount=N的参数来限定。
其次则是这种做法限定的是Regular Registers, 注意到参数中是maxrreg, 而不是maxreg了么, 中间多了一个r. 而其他种类的寄存器, 例如predicate register或者uniform register(计算能力7.5+), 则无法通过这种方式限制, 但好在一般我们也不需要限制后两种寄存器的数量。
这是手册今天告诉我们的第一种限制方法, 简单明快的限定成N这种具体值, 比较直接.
而手册中说到的另外的一种限制方式, 则是通过__launch__bounds__()来修饰kernel本身,将此行放置在kernel的最前面, 即可限制该kernel的寄存器使用数量。
注意这种方式可以每个kernel单独放置一种修饰, 甚至可以每个kernel根据编译时候的计算能力选择, 放置多种修饰. 控制性比较强，因为它不想-maxrregcount那样的是整个文件一起来的, 人家是单个kernel, 甚至单个kernel的单个计算能力编译下的效果来的, 所以可以很精细的指控。
但是坏处是, __launch__bounds__()无法直接指定一个具体的寄存器用量N, 而是间接的指定我需要1个SM上最少有XX个YY线程的Blocks, 然后编译器再自动计算一下, 这个XX个是需要限制到多少个寄存器的情况下, 才能满足约束, 类似这种的. 比较晦涩一点。
所以实际使用中, 这两种方式可以根据需要来, 一个直接的粗糙的; 一个间接的精细的. 读者们可以尝试根据实际需要使用.
这是我们今天所说的, 通过限制寄存器数量来尝试优化性能的两种具体做法.
下一篇, 我们会说一下菱形启动符号, 也就是<<<>>>这种, 和其他一些方面, 能带来的性能变化。
occupancy越高越好么？
在今天的手册上, 这些统称为对kernel的执行(环境)配置, 来调节性能.
首先是, 手册提出了occupancy的概念, 和这个概念的重要性, 以及, 观察从而能设定occupancy的3种方式.
各位读者, 只要是用CUDA的, 就一定遭遇过occupancy这个词, 俗称”SM占用率”。这是一个百分比值, 例如某kernel在某卡上运行, 取得了90%的占用率; 而某某kernel, 则在此卡上, 只有30%的占用率, 等等。
你的同学, 同事, 朋友, 总在会尝试劝告你说, 一定要提高这个占用率啊, occupancy高了才能性能好啊, 否则你就在浪费你的卡啊, 类似这种说法.
这种说法对也不对。
首先手册说了为何这种说法对:
因为我们的GPU是一个吞吐率为设计目标的处理器, 每一个晶体管都是尽量为了最大化的提供性能而存在, 而不像CPU那样, 为了延迟而设计, 很多晶体管都在为尽快做好1件事情而努力. 这点我们之前说过.
所以为了这点, GPU需要用海量的线程在SM上执行着, 当某些线程(精确的说, warp)卡住了, GPU从而能切换到另外一些线程去执行. 用这种简单的方式, 最大化的发挥性能. 等于我同时在洗衣, 做饭, 看娃多种任务同时进行, 一旦我卡在了等待洗衣机运转上, 我就可以切换去做饭, 一旦做饭正在煮着了, 我就可以去给娃喂奶. 用这种方式反复横向切换, 从而能最大化的利用我的时间(GPU的性能).
而Occupancy则代表了, GPU的SM上能驻留的线程数量(我今天在干的活的数量), 和该SM的最大能驻留的线程数量的(我累死最大能同时干的活的数量)的比值。这是occupancy的定义(实际上略微有差异, 特别是涉及到achieved occupancy的时候)。
100%的occupancy可以看成我一共能干10件事同时, 我今天就是在同时干10件事，而20%的occupancy则是我一共能干10件事今天, 但是我今天只同时干两件事。所以你看到, 一般情况下来说, 越高的occupancy(接近100%), 往往会越可能的发挥性能; 而越低的occupancy, 则往往会可能造成设备的低效运转。
这是今天手册上说, 为何为何尽量提高occupancy, 往往会提高性能的原因, 也是你的同事, 朋友, 同学往往会建议你这样提升的原因，但是事情不是那么绝对的, 有的时候, 较低的occupancy反而可能会带来更好的性能, 这点在历届GTC的演讲中都有提到过, 网上也能搜索到很多案例。
手册这里总的说法是, 因为当SM里的总资源固定的时候(想想成你家的面积好了), 较低的occupancy(想想成今天你只干2件事好了), 会给每件事带来更多的资源(想想成, 你需要一个手工画一个图, 较大的桌子可能让你干的更起劲, 从而比你同时用小桌子绘图憋屈, 同时在煮饭的总产出要好)。
下文我们会和手册一起, 对具体SM里的资源进行逐方面的分解, 看看occupancy vs 资源 vs 性能变化的具体讨论.
测量Occupancy的三种方式
一般的来说, occupancy往往有个折中点, 过高了或者过低了性能都不好. (就如同你干得过少, 或者干得过累都不好一样). 好了, 我们有了occupancy的概念, 知道了无需一味的去追逐occupancy, 就已经是一个很大的胜利了. 我们下面将具体看一下, 如何测量, 调节occupancy, 并从理论的角度看下它们可能带来的性能变化。
手册里先说了计算/测量occupancy的三个方式, 然后再说了调节一些资源的使用, 会occupancy造成怎样的变化可以反映出来。
我们先看看手册说的occupancy的测量/计算方式. 这个其实以前在编程指南手册上也有涉及, 只是可能没有今天的这样的系统一点。
一种是纯手工计算, 纯手工计算是指的人为的设定或者找到某kernel的, 寄存器使用量, shared memory使用量, block里的线程数量这三种因素/资源的使用后，通过和手册中的特定计算能力下的这三种资源的情况(该表在编程指南手册的后面有)对比, 从而手工的计算出来一个理论的occupancy.
这种方式不需要任何工具, 而且可以在你敲代码的时候, 自然的在大脑中提前形成大致的轮廓. 坏处是比较枯燥, 而且你需要比较熟悉特定计算能力的情况, 才能大致的对比出来你的kernel, 将来在该计算能力的卡上运行, 会得到一个怎样的理论occupancy。
第二种计算/测量occupancy的方式, 则是使用工具. 具体的可以分成静态的一个Excel文件(即本章节的图中的occupancy calculator的那个.xlsx文件)，在里面选好了寄存器资源, shared memory资源, block中的线程数量, 和对应的计算能力后, 该.xlsx文件中的宏, 会自动为你计算一下.
该静态的计算工具文件好处是可以免除记忆特定计算能力的参数, 而且还提供了一些高级参数(手册所没有提供的), 例如特定计算能力的某些资源是按照什么方式/粒度分配的信息. 同时不需要安装复杂的开发环境, 例如你可以将此excel文件复制到手机上打开, 或者上传到OFFICE365, 一样可以随时随地使用。坏处是你依然需要手工去测量/计算这3个基本资源, 在你的kernel下的具体使用量, 才能进行后续计算。
而同时也存在另外一种工具, 动态的分析工具, 指的是nsight或者nvprof类似这种的profiler, 它们会在你的kernel运行起来后, 自动为你抓取到这个信息, 从而免除了3个基本数据的手工取得, 也免除了后续的计算过程, 全自动。坏处则是你只能取得你所拥有的计算能力的卡(例如, 一张8.6的30HX显卡(尚未问世, 我们假定的30HX是最近的8.6计算能力)), 在此卡上实际运行时候的数据，你无法取得你所没有拥有的一张计算能力的卡上的情况. 但是手工计算和用Excel计算都是可以算出来一张不存在的卡上的情况的。
对于公司开发的情况, 例如拥有所有的要出售的产品所针对的, 市场上的所有代的计算能力的卡或者Jetson产品, 都购买回来的情况的时候(每种至少一个), 则无需担忧这种。这是第二种用工具的方式.
而第三种则比较主动一点了, 可以编程的通过相应的occupancy api (见cuda runtime api的手册, 或者我们之前的编程指南的稍微提到的部分内容), 在运行的时候, 动态的获取到我的某kernel, 在现在的某卡(例如3号卡上), 用XXX的资源配置或者线程形状, 能取得百分之多少的occupancy。
这种方式坏处是需要用户编程, 增加了额外的编码负担(和出错的可能). 好处则是, 你的代码可以在将来的卡上, 在开发的时候无论纸面或者实际的资料都没有的情况下, 在未来的某一天实际运行的时候, 代码自我分析和发现得到occupancy. 例如将来在一张30HX上, 此卡尚未问世, 我们也不知道计算能力的情况, 但是用第三种API的方式, 将来可以动态的得到, 从而潜在的能动态的(用代码)微调occupancy。
好了. 这是关于取得/测量Occupancy的三种方式, 今天我们简单的说了, 寄存器资源的限制, Occupancy的意义和高低对性能的可能影响, 以及, Occupancy的具体测量/计算方式。这三个因素其实还挺重要的，很多时候我们写代码, 当算法固定了, 实现也基本固定的情况下, 想调节性能, 只能从这3种基本不太影响现有不的代码格局的方面入手。所以关于这3方面的优化调节, 也往往排在算法—>实现—->(今天的执行/配置方面的调节)这么的一个重要顺序.
因为例如有更好的排在前面的情况, 例如一个快10倍的算法, 你应当先去考虑选择它, 而不是今天的这些”优化方面”，你很难简单的通过”优化”去将一个GPU上的应用性能继续提升10X, 但是更换算法, 你有可能。所以大家在实际使用中, 不要舍本逐末, 应当至少什么是最先考虑的. 只有当最先考虑的因素都完成后, 再进行这些介绍的经验和手册告诉你的实践操作. 就如同刚才说的某妹子, 她如果直接嫁一个100倍有钱的老公, 还管这些一天怎么干活, 怎么做事? 后面的这些将毫无意义.
我们在下次的内容中, 将会具体结合寄存器, shared memory, block形状这三种因素, 综合occupancy分析, 3因素 vs occupancy vs 性能的情况.
如何执行配置优化以及对性能调优的影响
接上一天的occupancy后面，继续说说寄存器的延迟掩盖，blocks形状和使用，shared memory的使用，以及，concurrent kernels和CUDA Context等方面，对性能调优的影响。
首先我们从寄存器的延迟掩盖开始。本小结首先讲述了，当需要使用寄存器中的数据，而该数据没有准备好的时候，从而无法取得数据喂给SM中的执行单元，从而可能导致执行的线程被卡住(stall)而不能就绪执行的状态。小结只讲述了常见的A = XXX; 这种形式的寄存器上的结果计算延迟。并用volta举例常规的计算有4个周期的延迟，在此期间内，立刻使用结果数据是不可以的，需要等待4个周期才可以。并讲述了可以临时切换到其他warps中的指令继续执行来掩盖的方式。本小结是乐观的，认为这一般不构成对性能的影响。
但是实际上，随着现在nsight compute的流行，long/short scoreboard的stall reason之类的分析指标的公开，很多操作对寄存器的结果写入，可能要超过这例子中的4个周期不少。感兴趣的读者可以参考这个链接：cloudcore：CUDA微架构与指令集（4）-指令发射与warp调度 。这里的讨论比当年scott grey在NV英文论坛的讨论要热闹一些，下面也有一些NV的国人在加入讨论。进一步扩展的读者可以参考里面的相关scoreboard的内容继续展开。
我们这里只额外说一下，使用s_xxx[idx] = d_xxx[idx]形式的，从global memory看似’一步到位’写入到shared memory的做法。实际上会被编译成中间的分步的tmp = d_xxx[idx]; s_xxx[idx] = tmp; 的经过寄存器(tmp)的分解过程，导致中间第二次写入的时候有一次对寄存器的依赖。使用8.6和8.7计算能力的人们，建议考虑新版的cuda::memcpy_async的载入方式，这种可以直接越过寄存器。
这是今天的第一小节。
第二小节讨论了block和grid的形状对性能的影响问题。这个是个喜闻乐见的讨论，在我们夏令营和冬令营的活动中，被人讨论了无数次了。小节首先澄清了，grid和block的1D还是2D还是3D的形状，从本质上并不影响性能，影响性能的只是无论1D还是到3D时候的，计算出来的每个block里的线程总数量，和blocks的总数量。
小节同时说明了，这些线程和blocks的数量（和其他资源），影响了在SM上的active warps的数量。能达到的active warps数量，才是之前的occupancy之类的很重要的原因。而active warps的数量，往往决定了延迟掩盖，和对SM各个单元的利用程度。这样性能就取决于这些单元的利用率情况，因为一旦我们买回来了一张卡，硬件的SM数量，和SM里面的执行单元配置是固定死的了，硬件本身乘以利用率，才会影响最终的性能发挥。
然后小节往下说了，该如何调整kernel启动时候的方括号里的第一个和第二个参数。大部分情况下，调优kernel，需要同时(in tandem)试验性的调整这两个参数。但每个参数也有他们自己的调整策略：
对于第一个参数(blocks数量): 基本的策略是要足够多，至少每个SM上得有1个block。同时，考虑到了1个SM上如果只有1个block的话，一旦该block中的线程们，执行了__syncthreads()进行等待同步的话，很可能导致SM上warps大部分都处于等待状态了，降低该SM的使用率。所以这个至少的1个block还需要调更多。手册的建议是，亲这边应该至少上几千个blocks每张卡。理由很简单：考虑到现在的8.6的3090的卡，有82个SM。每个SM上可以上到多达16个blocks，这样82 * 16等于差不多1000。几千个差不多能将一张卡上个几批次。手册说到，我们要面向未来考虑，将来的卡更强。所以数量不能保守。
阅读到这里，我们应当结合实际一点。因为随着block对资源的使用不同（例如shared memory), 一个批次能上多少个blocks，对于固定的卡，随着kernel的不同是不同的。建议读者使用nsight compute, 观察里面特定kernel的waves数量指标，该指标说明了某kernel的blocks需要分成几个wave（批次），才能上完。
以及，对于某些因为算法的角度的限制，不能有效扩大blocks数量的情况下，针对本章节讨论到的，因为__syncthreads()而导致1给block中的warps在SM上整体stall的问题。可以考虑使用细粒度的部分同步手段。也就是使用cuda::barrier（需要计算能力7.0+），进行1个block中的部分线程进行同步。这样当部分线程在wait()或者arrive_and_wait()进行同步的话。该block中的其他不参与barrier同步的线程依然有机会执行，继续利用SM上的执行单元。
以及，新版本的上一部分手册(CUDA Programming Guide), 现在已经正式引入了很多C++风格的东西了。上一段说到的asynchronous barrier, 在当年我们阅读编程指南的时候，没有涉及。建议读者重新阅读相关章节。
然后继续回到<<<>>>的第二个参数，也就是block中的线程数量的优化考虑。手册这里主要考虑了你不能用过小的blocks，例如只有32个线程的block. 因为SM往往还有例如16个block/SM的硬限制。使用过小的block往往会导致SM上去的总warps数量不足，可能会影响性能。手册这里建议的方式是，至少上64个线程的block，然后逐步调整block中的线程数量, 找到特定kernel的最佳性能点。这个逐步调整，可以从128或者256个线程起步。
手册继续说，调整到适可而止就行了，没必要追求极限。例如通过调整前两个参数，让SM能上到66%的occupancy，和能上到100%的occupancy，可能并不会对性能起到太显著的影响。因为调整的目的是追求性能，而不是单纯追求指标。为了得到过高的occupancy，有的时候你只能降低寄存器数量之类的，从而导致使用了过多的local memory, 反而影响性能。
而另外一方面，因为除了我们之前说过的TLP（例如依靠切换warps）来充分利用硬件的执行单元，还存在ILP的方式，也就是线程内部的前后指令本身的并行性，来提高效率。手册这里指出了，只要内部的ILP程度足够，哪怕较低的occupancy也是足够的。对于这个问题，我们建议读者继续扩展阅读经典文章：《Better performance at lower occupancy》（链接： http://dmacssite.github.io/materials/volkov10-GTC.pdf ），该文章描述了哪怕很低的occupancy，也可以通过ILP取得优异性能的方式。虽然这个文章较老，但是依然非常经典。
另外的，我们夏天搞夏令营活动的时候，客串出场的樊博士，也在他的实践中(GPUMD项目），指出了这点，例如在他的《Efficient molecular dynamics simulations with many-body potentials on GPU》中，老樊写道：“哪怕使用float的时候只有50%的occupancy；或者使用double的时候只能到25%的occupancy。性能也相当不错”。（arvix: https://arxiv.org/abs/1610.03343 ), 感兴趣的读者也可以扩展阅读。
这两篇文章都分别有12年和5年的历史了，但是里面的思想，是正确和不过时的。
函数和指令使用的选择和优化
今天的主要内容是<优化指南>手册里面，对一些函数和指令使用的选择和优化。大致分为普通的计算函数/指令，和访存相关的方面。
我们先从计算函数/指令开始。
首先上去的小节，是关于整数除法和求余操作的优化写法。当除法A / B, 和求余A % B的时候，如果B是2的整数次方，也就是B = 2^N的时候，前者A / B可以直接写成移位操作A >> N；后者A % B, 可以直接写成逻辑与操作A & (N - 1)。 无论是移位操作，还是逻辑与操作，都是单周期的指令，远比老老实实的除法和求余快得多。
手册本节指出了，当B是编译时刻的2^N形式的常数的时候，编译器会自动发现这一点，同时自动为你进行这个优化。但是如果B不能在编译时刻确定，例如作为一个参数，B传递给了kernel，此时为了避免进行昂贵的除法和求余，可以考虑手工将B转换成指数值，然后手工进行移位和逻辑与操作。例如原本要传递进来256, 现在可以传递进来8（也就是log2(256)), 然后直接A >> 8和A && (8 - 1)即可，从而规避了昂贵的代码产生。
这是第一小节。第二小节则依然是说的整数，主要涉及到在使用下标和循环控制变量的时候，对有符号整数和无符号整数的选择。并讨论了C语言默认为有符号整数时候，编写代码的人如果偷懒不写上unsiged字样，则在循环控制变量和下标计算上，将生成较为劣化的代码。
小节说明了，这是因为无符号整数的溢出和累加都很方便，而有符号的则需要处理溢出的特殊情况，需要占用额外的指令。
我们这里以前忽略过这点，今天我们用计算能力8.6上的指令生成，分别测试了默认情况，和标注了unsiged字样的整数，在这两种情况下带来的优势——我们给读者测试了对于常见的形如p[i * 8] = 0，当i是int和unsigned int时候的，单语句的代码生成的效果对比，用来验证一下手册的这个优化的说法：
在i是无符号整数的时候，p[i * 8]编译生成了2条指令的序列。
指令（1）：用LEA指令计算p的低32位地址累加i左移3位
指令（2）：如果有进位溢出，p高32位+1
我们的GPU是32位机，只能每次进行32位整数运算，对于这p[i * 8]形式的64-bit最终地址计算，这已经是最优的代码生成了。
而在i是常规有符号的整数的时候，却编译生成3条指令的序列，多了一条：
(1)单独计算i * 8的值
(2)整数加法, 并得到是否溢出的标志
(3)根据溢出标志，执行32位符号扩展的LEA.HI.X.SX32指令。
你看，在使用下标的时候，在int i的定义身上，简单的加上unsigned的无符号标注，就能得到性能优化。
类似的，根据手册本小节的说法，当下标在循环里面的时候，编译器还可以对unsigned的下标，进行更强的替换处理(strength reduction，参考: https://en.wikipedia.org/wiki/Strength_reduction ), 例如我们在一个for(i)循环中的p[i * 8]的使用，发现了每次i的递增，乘以8被reduced到每次加8，和地址的计算等方面的指令生成，也有类似的优化效果。所以你付出的代价只是将声明变量的时候，添加一个unsigned标记，就可以得到显著的好处。这点值得考虑。
以及，本小节实际上说的是：对于循环变量尽量使用有符号的整数，理由是，无符号的行为是精确定义的，有符号没有精确定义溢出行为，所以编译器有更多的操作（优化）空间，但是我们编译测试发现是反的，建议读者们自己实验决定究竟是什么情况。
—刚才例子中的无符号情况的生成结果(cuobjdump), 一共两条指令。LEA和加法(8.6上用FP32 path的IMAD.X的A + 0 * B + 进位指令，模拟了A + 进位加法）。
—刚才例子中，有符号情况的生成结果, 一共三条指令（移位用FP32路径的IMAD.SHL模拟替换）。
也就是移位，第32位加法，高32位符号扩展加法。这三条。
（关于7.5和8.6上，用FP32路径上的IMAD/IMUL的指令模拟常规INT32指令，达到port平衡，是另外一个话题。这两个计算能力都能超过64指令/周期/SM的INT32的指令峰值上限，因为这种模拟替换和平衡，用nsight很容易发现这点）
好了。两个小节的整数指令方面的优化选择说完了，我们下面继续今天的主要内容，关于float方面的优化选择。
首先说的是，计算1/ � ，时候的做法（你想法将X放到根号里面，我不会），本小节指出了，有单独的单精度和双精度的rsqrtf()、rsqrt(), 来直接完成求根号X，然后再求倒数的一体化运算。如果可能，尽量使用这个。会带来更好的效率和精度。编译器并不总是能将1.0 / sqrt()的写法，转换成对应的一体化函数版本的。
然后下一小节手册从上面两个相似名字的数学运算函数（结尾带有f和不带有它）开始，说了容易不小心将float写成double，并生成了double运算代码，导致速度降低很多的情况。主要有这两点：
(1)读者写代码的时候，如果不小心，使用1.0，而不是1.0f这样的常数，根据C的规则，含有这个常数的式子，将在运算过程中，提升到double进行运算，式子算完后，再转换回来成float进行赋值给lvalue对象。这样有了来回转换double<->float的指令开销，也有了慢速的double指令计算的开销。
（2）CUDA编译器实际上是一个C++编译器，在math_functions.h之类的头文件里面，有C++风格的重载。例如sqrt()函数，有double sqrt(double)的版本的，也有float sqrt(float)的版本的。如果用户不小心，在式子里面给出了double的中间结果作为参数，同时函数结尾没有显式的写出f()结尾，那么因为重载的同名函数存在，将实际上使用的是慢速的double版本的。也有生成慢速的代码。
所以我们读者应当尽量小心注意使用浮点常数和函数后面f结尾，避免生成慢速的代码（double总是要慢的，而且会占用更多的资源），特别是在家用卡上（8.6的家用卡走double路径将只有1/64的速度）。
我们的老朋友樊博士，对于忘记了f后缀写上，而导致代码变慢了很多很多，具有惨痛的教训。并在冬令营/夏令营上，给我们深刻的说过这点。
然后这小节还提了在进行概率统计之类的运算的时候，如果要使用正态分布的误差函数，特别要注意这点。因为erfcf()这个函数（注意f结尾），在单精度的时候特别快。例如我们在计算N(0, 0.5)的正态分布的2个西格玛内的概率的时候，使用float p = 1.0f - erfcf(1.0f / 0.707f);类似这种写法（注意好多f结尾），将特别快。
最后这小节还提到了，不仅仅我们浮点数有这种情况，8-bit和16-bit的整数，在直接在我们GPU上使用的时候，通常情况（不考虑深度学习时候的多个打包在一起的运算），都需要转换成32-bit整数，才能进行运算。这是因为我们的N卡，在进行整数计算的时候，是严格的32-bit机器，不像x86 CPU那样能就地干8-bit和16-bit指令。这样不小心就会导致额外的代价产生。
总之，适当的写法，和数据类型的使用，能避免转换的代价，和昂贵代码路径的生成。读者还是需要注意这里的优化的。
下一节则谈论了对我们读者们喜闻乐见的powf()、pow()（分别是float和double版本，如上面说过的）的通用幂函数运算时候的优化，主要针对了几种特殊的指数值，可以不用通用的幂运算完成：
如图，通用的指数运算的快速替换写法。
我们这里简单的举几个例子就好：
计算x的1/6次方，可以先计算一次x的平方根倒数，再计算一次立方根倒数，这样就得到1/6次方的值，而无需使用昂贵的pow之类的函数。再例如：x的2/3次方，可以先求一个立方根，然后再求一次平方，这样就快速得到了2/3次方。
注意这个快速替换表格里的公式，很多都使用了特殊的GPU上专用的函数，例如rsqrt, rcbrt（二次方和三次方根的倒数），而不是标准的C库(libm)，在CPU上我们能见到的sqrt、cbrt(二次方和三次方根），如果我们读者从以前的代码编写经验来，可能喜欢使用嵌套两次立方根，得到1/9次方的值，我们不推荐读者这样来。因为特殊的rsqrtf()这种，可能在实现上具有更好的精度和性能。例如我们之前的章节知道过，SFU这种喜欢提供平方根的倒数的这种快速的接近，可能有助于性能的提升。总之我们建议保留此表格，直接使用里面的写法，而不是读者使用更熟悉的替换形式，为了能够保证足够的精度和性能。
注意事项
今天的主要内容将讲述三个方面。一个是Global Memory访存时候的优化注意事项，另外则是循环或者条件语句导致的分支时候优化注意事项，和几个消除warp内分支的案例。
我们先从今天的对Global Memory的优化开始：首先，优化指南手册说，要尽量减少对Global Memory的使用。因为uncached的global访问，将有巨大的延迟。应当考虑尽量使用shared memory代替，如果可能的话。同时手册给出了一个s_buffer[index] = g_buffer[index];的简单从global载入shared memory的例子分析。
先说下这里的uncached的意义，在NV的多个文档中，uncached是指通过L2 cache进行的访存，例如对显存或者映射过来的内存，而不是完全没有任何cache，这点需要注意用词。
其次则是对于这个例子来说，手册指出了：看似只有一个简单的等号，就赋值了，但这里实际上是右边进行了从global memory ===> 寄存器的载入，然后再从左边寄存器===> shared memory的写入，这样右侧的载入将导致很大的延迟，影响性能。这种代码实际上经常在我们的代码中出现，例如1个warp或者block，从global memory载入一个矩阵，然后在shared memory重新映射下标，进行转置；或者是基本上大部分的kernel开头，将一些公用常见数据，从global进行载入到shared memory，方便后续的使用。那么这两种情况应当具体怎么优化呢？
首先，无论是哪种情况，手册说了，可以完全不用手工处理，依赖硬件自动优化即可。因为GPU的SM硬件本身，存在warp之间的调度，当一个warp的从global的载入，需要等待而卡住的时候，另外一个warp可以被调度执行，从而可以互相掩盖了。也就是常说的TLP延迟掩盖。但这种情况下想要有较好的效果，你可能需要有较多的warp发出进行中的global memory访存请求才可以。例如论坛最近的这个例子：求问几个有关优化方面的问题 。我们的会员已经明显的意识到了这一点，并试图用launch_bounds来自动化的提升驻留的活跃warps数量，从而自动进行这点优化(虽然他的效果不太好）。
此外， 在计算能力8.0+的卡上，可以手工使用memcpy_async进行异步的global memory —> shared memory的载入，从而直接越过global —-> register —-> shared的中间的寄存器的依赖。这种优化很多时候有正面效果，但是不一定总会发生。例如对于刚才的2种情况来说：
如果是kernel开头就需要载入一些数据到shared memory, 必须等Shared memory中的数据就绪了，才能作为初始值参与运算的话，这种直接载入（用等号），和手写memcpy_async异步载入并无本质区别，因为你都会立刻卡住在初始数据的准备上。
但如果你是kernel运算的中途，需要载入部分数据到shared memory, 则memcpy_async的异步载入就很有用了。像是流程: 载入1—->数据1计算—-> 载入2—->数据2继续参与运算—->载入3——>数据3继续参与运算的流程。这样的话，可以按照另外一本手册的流程，稍微变形成为：载入1, 异步载入2—->数据1计算, 异步载入3—->等待之前的异步载入2就绪并参与计算—->如此重叠。
这样的话，根据你重叠的程度不同，如果有一份重叠，像是刚才箭头的那样，则等于你提升了两倍的驻留warps，或者说提升了两倍的occupancy的延迟掩盖效果。而如果你还有余力（手册上有重叠的程度的做法）提升overlap的async load的程度的话，例如有2X的重叠，则等于提升到了三倍的occupancy。效果将很显著。
特别是对于那些occupancy已经提升到了极限了，或者是我们刚才的那个论坛的链接，因为某些原因无法继续提升occupancy了，则用这种方式掩盖global memory的延迟，将是另外一条路的选择，它有效的提升了等效occupancy，或者从另外一个角度说，提升了warp内部的ILP并行。
这是今天的第一点，对于global memory的高延迟的掩盖优化。
今天的第二个方面则是重头戏，也是频繁在我们论坛上出现的内容。实际上，刚才那个论坛链接，也涉及到了下面的内容。即循环和条件控制指令，以及warp边界上的分支，和warp内部的分支，对性能的影响。
手册首先说，if/while/for之类的语句，会生成进行边界判断和流程控制指令，而这些指令本身会影响性能（因为要执行的指令变多了）。这种情况下，对于循环，可以考虑用#pragma unroll将循环展开。展开后，单次循环所需要的边界判断、循环控制变量的增加或者减少运算、跳转之类的操作，就一定程度的消失了，从而执行的更多的循环体的计算本身，提升性能。这是一点。
另外一点则对于循环和判断语句来说，这些附带的流程控制指令，还可能会导致warp边界上和warp内部的分支，也会影响性能。
手册说了，warp边界上的分支就是以32为单位（目前的warpSize是32），不同的warp在执行了流控指令后，跳到不同的位置来执行，这种情况的分支，具有较弱的性能影响，因为可能只有对于I-Cache指令缓存之类的取指令的代价增加。而手册继续说了，warp内部的分支，也就是我们喜（深）闻（恶）乐（痛）见（绝）的divergent branch, 具有比较严重的性能影响。这种分支，是在每32个线程的内部位置，不在32的整数倍的边界上，会导致warp内部不能100%的效率执行代码。因为SIMT的目前的CUDA的执行情况，例如有2个分支，分别在1个warp内部，有13个线程，和19个线程要执行两段不同的代码。只能scheduler分别以13/32的效率（41%）执行一段，和以19/32（59%）的效率执行一段。浪费了执行单元的峰值性能。
例如刚才论坛的例子，这里的profiler报告的实际有效执行效率才13.5/32 = 42%。
如果是那种卡计算的kernel（本例不是），一张3090秒变成了3070，对于warp内的分支(divergent branch)，先不要着急, 我们稍后说下如何解决这个优化上的重头戏。然后手册继续往说，除了warp边界的分支，和warp内的分支，手册还指出了第三种情况导致的性能损失，就是在7.0+的卡上的独立线程调度(Indepent Thread Scheduling)导致的性能损失。
我们都知道，常用的图灵和安培这两代卡，都有独立线程调度，和因此引入的新的syncwarp()或者XXX_sync()之类的新函数，例如我们以前用的shfl, 变成了__shfl_sync(). 这种独立线程调度，有它的好处，一个最大的好处就是可以应对渣代码，和应对直接从CPU上移植过来的线程之间的锁之类的东西。
渣代码这里指的是目前我们好多群里都比较常见的，SM和MEM任何一个都无法用满的情况（例如论坛的这个链接里的情况），看上去不卡计算，也没有卡在访存瓶颈上（这种往往卡在了延迟等待上）。对于这种代码，7.x和8.x的独立线程调度，允许以比warp单位更细的粒度，进行线程调度执行。例如说：
1 | if (xxxx) |
这种代码在7.0+的卡上，因为独立线程调度的存在，当if条件有效的上面的那个路径，卡在了等待x的值读取回来的时候，可以让warp中的其他线程切换到else部分先执行着f2和f3，而不至于因为f1卡在访存等待上，导致其他分支不能执行。这样总的延迟就从 很长的访存—->f1 -> f2和f3，变成了很长的访存（暂停) —-> 切换到f2和f3执行——>（差不多之前的延迟完成了)，再f1继续执行。这样就有效形成了warp内部的不同分支不同的重叠执行。f2和f3部分的执行，掩盖在了另外一个分支的延迟中了。
我们用数字继续说明，假设从宏观的SM角度看，如果传统上：读取长延迟期间SM 0%有效使用率，10X周期；执行F1期间，因为warp不满，40%的效率，5X周期；F2和F3期间，也是warp不满，60%的效率，8X周期。则没有独立线程调度能力的时候，总宏观上的SM执行成本：10X周期的SM 0% -> 5X周期的40% SM -> 8X周期的60% SM效率。总共用了23X周期，SM有29.5%的平均效率。则在7.0+的卡上，变成了：10X周期的读取等待（期间SM以8X的周期，60的效率执行了另外的F2和F3操作；和2X的空闲），5X的40%效率的F1操作。这样独立线程调度情况下，一共只需要15X周期的时间成本，减少了8/23; 同时期间的平均效率，SM也提升到了45.3%，提升了20%。
所以这种卡上，虽然分支依然是divergent branch的warp内分支，但是总执行时间成本，和平均SM使用效率，都提升了。也就是说，独立线程调度，不能自动解决warp内分支的问题。但在举例的这种情况下，可以减少延迟的等待，和期间的平均执行效率。那么看起来独立线程调度的确是一个好事，但优化指南手册为何说它可能会导致性能降低呢？（刚才的例子明明是提升了啊！）
那是因为，独立线程调度，不会自动汇合不同的分支，从而往后原本正常的warp内完全合并无分支的代码，再继续享受着独立调度的好（恶）处（果）。我们还是用刚才的例子好了，实际的代码中，再宏观点看，往往是这样的：
1 | if (xxxx) |
我们刚才说了，在if和else里头，我们享受到了福利了。那么后续的zzzzz可不是这种情况了。因为独立线程调度不会自动愈合，后面的zzzzz将会依然处于独立调度阶段，也就是说：
原本走if分支的那些线程(40%)，会以40%的warp中的线程有效状态，执行一遍zzzzz；然后走else分支的那些线程（剩下的60%），会继续再按照独立线程调度，走一遍zzzzz！这样的话，后续的就完全没有福利了，反而是性能下降了。
此时优化指南手册告诉我们，我们需要使用例如一条syncwarp(), 让warp中的40%和60%的两部分线程们，进行合并。这样合并以后，就会用100%的warp内（完全无分支）效率，一共执行1次后续的zzzzz操作即可。这是我们在较新的卡上写代码需要注意的，可以时不时的加一个syncwarp(), 往往有助于性能提升。以及，因为syncwarp()本身会编译出来一条WARPSYNC指令，占了1条指令发射机会。如果你认为不能太频繁的使用过多的话，也可以考虑带有sync结尾的其他函数，例如说：__shfl_sync(), 它正常的作用执行一次warp内的不同线程间的数据交换（执行成本类似一次无bank conflict的最好情况下的Shared访存，见GTC上的多次演示）。同时它还能给你带来一次warp的sync效果，却不需要任何额外执行，也可以考虑这点。
这是今天手册的第二部分，也就是warp间分支、warp内分支、较新的卡上的独立调度引入的后续隐形的分支的优化注意事项。
那我们继续今天的第三部分，也是最后一部分，关于之前的重点需要优化的：warp内部的分支（divergent branch）的优化处理方式。
还记得刚才的3090秒变3070的例子么？divergent branch很多时候非常伤害性能的。论坛的这个例子，楼主也在苦苦追求减少warp内部分支的方法。我们将大概说明3个常用的小技巧，来处理warp内的divergent branch。两种分别应对warp内有选择性的执行一部分代码，其他线程等待。和一种对应warp内分支有常见的两种代码路径，warp需要部分线程选择性的执行一部分，和其他线程再选择性的执行另外一部分的情况。而其他更复杂的分支，往往是这三种情况的组合，就不说了。
先说一下第一种情况，我举个例子。常见的GPU上，有N个线程，每个线程处理1个任务，一共处理完N个任务，也就是最常见的CUDA的分而治之的典型做法。在这种典型代码中，这N个任务往往处理的成本不同。有的任务可能只需要50%的处理步骤就能完成，有的任务可能却需要90%的时间，而有的任务则特别快，10%的时间就能完成。这种代码跑起来，会性能warp中的空洞。例如所有线程的平均是100%的时间，某warp内有19个线程，只用了10%的就完成了，剩下13人却还需要很长的时间，那么等于有大部分的时间，warp都只有13/32的执行效率，剩下的19/32都变成了空洞袖手旁观。
这种情况有两种常见的处理方法，一种叫warp compaction, 一种叫task reload/lane bubble fill, 搜索相应的文章能看到大量的具体技巧的应用效果演示，有图有文字。
简单的说一下前者就是，如果我有一个256线程的block，也就是里面有8个warp，处理一批256个任务。其中在代码的执行1/3位置，预估有一些人会提前结束（例如每个warp中大约会有40%好了）；在代码的2/3位置，每个warp大部分的人会结束（例如大约warp中大约有90%）的人会结束；而剩余的1/3代码执行量，则为warp中的10%的顽固分子。
那么整个流程看起来，前1/3每个warp都有100%的执行效率，中间1/3，每个warp只有60%的有效执行效率，后1/3，每个warp只有10%的可怜效率了。平均起来全程只有56%的总体效率。这下子，一张3090，可能变得连3070都不如了，怎么办？
这个时候可以考虑compact一下warps，例如说，在1/3位置压缩一下，从分布在8个warps中的60%有效线程，和40%的空洞，压缩到前面，变成了前4-5个warp基本全满载，后面3-4个warps无任务了，退出。这个时候效率立刻就又恢复到了几乎100%了。
类似的，我们在90%的位置也压缩一下，密集任务到1个warp里，剩下warps也退出。这样从执行多个warps中的不满lanes的代码，变成了只有较少数量的warps执行满载的代码。提升了效率。注意这个技巧需要前序和(prefix sum, 或者叫scan）操作。能否有效写对scan操作对于很多CUDA用户来说，是个问题。
而另外一个技巧则不需要prefix sum操作（虽然这个操作在NV的博客上已经出现了无数次了），比较简单，但需要每个任务是内部有多次循环，循环量不同的情况。例如我们有N个任务，其中40%的任务需要迭代100次才能完成，50%的任务却需要迭代500次，而10%的顽固的，需要迭代1000次才能搞定。
如果正常的写代码，那种1个线程对应1个任务的，随便抽取出来一个warp来说，平均来说：
整体一共需要1600迭代步，其中前100步是100%效率的；中间500步是60%效率；最后1000步只有可怜10%效率了。warp的整体迭代过程，只有31%的效率！现在3090可以变成3060了！怎么办？这种代码的提升warp中的空泡，可以考虑使用重新加载任务的方式来完成。
例如说：
1 | while(true) |
这种技巧不需要任何scan之类的操作，只需要你将线程的数量缩小到任务的1/M，这样每个线程平均执行了M倍的任务量，整体空泡将第1个任务后面的M-1个任务给填充起来。如果任务的时间步是随机分布的话，则这种方式具有较好的效果（平均M个任务后，warp里的每个人总时间都差不多了）。
上面的两种技巧，无论是从blocks中将多个含有空泡的warp，压缩起来；还是重新装填任务，将空泡填满。都能有效的提升warp的执行效率。像是我们论坛的这个例子的warp里的没有被predicated off掉的才30%多的线程的情况，就可以提升大约3X的性能。当然，这两种优化技巧，都需要付出额外的代价，不适合那种非常非常小的空泡/分支/任务，因为此时，你填充的”优化”代码所付出的执行成本，超过了你原本的细小空洞了，无意义。不能为了优化而优化，或者为了追求profiler报告的数字好看，硬上，毕竟小优化怡情，大优化伤身，强上优化灰飞烟灭。
然后我们继续看下另外一种情况，如果代码中不规则的夹杂了可选的代码路径，怎么办，例如这个：
1 | ....//正常处理 |
如果条件1和条件2在每个warp中对于每个线程，都有50%的概率进去的话，同时这种可选的额外步骤，占据总执行量的50%的话，那么整体执行效率将只有75%.
此时就不好使用刚才的重新装填任务来填充warp里的lanes空泡的方法了，不过可以考虑compact一下warps，但是这里有个问题，我们还需要恢复到正常的执行状态，来执行中间的正常处理过程（因为中间的这些正常处理过程，本来就是100%的warp无分支的效率）。此时你可以将之前的方法1进行变种，compact后，分配到纯空泡的那些warps/线程不能退出，需要在来一个__syncthreads()之类的等待（等待期间不会占用额外的SM里的SP或者其他单元的处理能力）。这样简单变种后，代码整体变成了：
(1)blocks中的所有warps都无分支满载效率
(2)少数warps满载或者接近满载的效率，剩下warps不占用任何执行单元资源。
(3)同步后恢复各自身份，继续回到(1)的情况。
注意这个方法有两倍的block内部数据交换的成本（因为压缩warps空泡的时候，线程间交换了一次数据；恢复身份的时候又交换了回来数据），和最开始介绍的约压缩越小的那种方式的每次压缩只有1次的成本要高的，是否整体合算，读者自行决定（或者搜索了相关的文章后，看他们文章的例子里的数据）。以及，实际上，如果读者能转过来弯，不怕数据的下标映射之类的混乱的话，实际上第二次交换可以省略，但需要较多的脑力成本（你自己想一下）。我们可能会在下一次冬令营讲述完divergent branch后，介绍这个优化方式，并出一道block内部交换数据成本较高的考试题，来尽量诱导大家不进行二次交换。
最后要说的则是上面这三种方法的扩展开来，对于常见的代码中的：
1 | if (....) |
最后如果你的kernel会是这种代码，总是充斥了这种两路分支的话，如果结合上面的方式，在if前重拍成为两组任务，需要两组前序和的计数序列。但是具体怎么做，这里就不说了，很容易能扩展得到。实际上可以简易的证明，对于一个有K个warps的block，总是可以得到至少K - 1的重拍任务后的warps，和最多只有1个的有divergent branch的warp。
这样我们就结束了今天的内容。注意divergent branch总是一个优化重点。以及，注意以上的所有优化方式都有一个前提：优化引入的额外操作的成本，要小于将warp内的空泡lanes填充后的收益。否则，优化就是白忙乎。注意，优化的方式不能万能的，得根据实际问题才能知道是否如此。所以有的时候，干活的人一顿操作猛如虎，最后没有收益，也不要失望。