cuda合并访问的要求_CUDA优化的冷知识 4 | 打工人的时间是如何计算的

这一系列文章面向CUDA开发者来解读《CUDA C Best Practices Guide》 （CUDA C最佳实践指南）

大家可以访问：

https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html 来阅读原文。

这是一本很经典的手册。

GPUS开发者：【CUDA优化的冷知识】什么是APOD开发模型？​zhuanlan.zhihu.com

GPUS开发者：【CUDA优化的冷知识2】 老板对不起​zhuanlan.zhihu.com

GPUS开发者：CUDA优化的冷知识 3 |男人跟女人的区别​zhuanlan.zhihu.com

今天主要说两点, 一点是如何正确的计算一段操作所用的时间。这里的一段操作是指的, 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如果被反复调用的话, 是会越来越慢的"。

我们看下该楼主的具体做法:

start = clock();
DeModuate <<<BLOCK_NUM, THREAD_NUM >>> (....);
end = 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)即可.

回到该楼主的帖子, 我们很遗憾的看到, 该楼主在我们给出了两次回答和解决方式建议后, 即分别要求楼主用第三方工具验证他的计时错误的前提(这样他可以自行发现他的错误, 从而增长经验), 和直接给出了建议(即明确的告诉了他哪里是理解错了后), 他均无视了我们. 并继续在后续的跟帖中, 给出他自行认为的理解. 这点我们是感觉非常可惜的.

实际上人是互相尊重的, 特别是在作为提问者, 你更加应该尊重回答者给出的信息的. 无视这一点, 并取得"面子上"的好处, 是无益于事情的. 我们在这里今天严肃的提出这一点。

是希望其他的客户或者非客户, 在论坛提出了问题后, 在看到论坛给出的解答后, 不要为了"面子", 带着有色眼镜, 从而实质上的无益于楼主们在论坛的经验的获取, 和以后遭遇相似问题时候的快速解决.

(反过来, 如果你尊重了论坛, 则你本次能反思得到了经验, 得到技术上的成长; 下次遇到后还能快速回忆场景, 快速解决, 节省干活时间, 增加在老板心中好的评价).

然后我们继续说一下该例子, 楼主的正当做法应该是:
(1)CPU记录开始时间
(2)CPU给GPU派活
(3) [CPU等待GPU完成该活]
(4)CPU记录结束时间

我们在这里插入了步骤3, 也是手册上今天的CPU计时内容章节, 所推荐的做法(cudaDeviceSynchronize()同步等待, 或者其他任何等效的同步方式). 只有加上了该等待, 你的开始到结束的时间差, 才是真正的干活时间.

下一篇我们将继续讲如何正确的计时. 因为这话题的确很重要了,