这一系列文章面向CUDA开发者来解读《CUDA C
Best Practices Guide》 （CUDA C最佳实践指南）
大家可以访问：
https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html 来阅读原文。
这是一本很经典的手册。
CUDA优化的冷知识|什么是APOD开发模型？
CUDA优化的冷知识2| 老板对不起
CUDA优化的冷知识 3 |男人跟女人的区别
CUDA优化的冷知识 4 | 打工人的时间是如何计算的
CUDA优化的冷知识 5 | 似是而非的计时方法
好了. 你已经会了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, 进行计时, 正确的逻辑是:
(1) cudaEventRecord(start, stream); //在流中发布计时命令, 要求记录start时间 (2) K<<<....stream>>>(); //在流中发布kernel K (3) cudaEventRecord(end, stream); //在流中发布计时end时间 (4) 同步
其中第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计时就没有大问题了.
