前往小程序,Get更优阅读体验!
立即前往
首页
学习
活动
专区
工具
TVP
发布
社区首页 >专栏 >CUDA优化的冷知识 4 | 打工人的时间是如何计算的

CUDA优化的冷知识 4 | 打工人的时间是如何计算的

作者头像
GPUS Lady
发布2021-01-06 14:46:18
8050
发布2021-01-06 14:46:18
举报
文章被收录于专栏:GPUS开发者

这一系列文章面向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 |男人跟女人的区别

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

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

代码语言:javascript
复制
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()同步等待, 或者其他任何等效的同步方式). 只有加上了该等待, 你的开始到结束的时间差, 才是真正的干活时间.

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

本文参与 腾讯云自媒体同步曝光计划,分享自微信公众号。
原始发表:2020-12-25,如有侵权请联系 cloudcommunity@tencent.com 删除

本文分享自 GPUS开发者 微信公众号,前往查看

如有侵权,请联系 cloudcommunity@tencent.com 删除。

本文参与 腾讯云自媒体同步曝光计划  ,欢迎热爱写作的你一起参与!

评论
登录后参与评论
0 条评论
热度
最新
推荐阅读
目录
  • Best Practices Guide》 (CUDA C最佳实践指南)
相关产品与服务
GPU 云服务器
GPU 云服务器(Cloud GPU Service,GPU)是提供 GPU 算力的弹性计算服务,具有超强的并行计算能力,作为 IaaS 层的尖兵利器,服务于生成式AI,自动驾驶,深度学习训练、科学计算、图形图像处理、视频编解码等场景。腾讯云随时提供触手可得的算力,有效缓解您的计算压力,提升业务效率与竞争力。
领券
问题归档专栏文章快讯文章归档关键词归档开发者手册归档开发者手册 Section 归档