DAY26:阅读性能优化策略

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第26天,我们今天开始讲解性能,希望在接下来的74天里,您可以学习到原汁原味的CUDA,同时能养成英文阅读的习惯。

本文共计304字,阅读时间5分钟

注意:最近涉及到的基础概念很多,所以我们备注的内容也非常详细,希望各位学员认真阅读

5. Performance Guidelines

5.1. Overall Performance Optimization Strategies

Performance optimization revolves around three basic strategies:

  • Maximize parallel execution to achieve maximum utilization;
  • Optimize memory usage to achieve maximum memory throughput;
  • Optimize instruction usage to achieve maximum instruction throughput.

Which strategies will yield the best performance gain for a particular portion of an application depends on the performance limiters for that portion; optimizing instruction usage of a kernel that is mostly limited by memory accesses will not yield any significant performance gain, for example. Optimization efforts should therefore be constantly directed by measuring and monitoring the performance limiters, for example using the CUDA profiler. Also, comparing the floating-point operation throughput or memory throughput - whichever makes more sense - of a particular kernel to the corresponding peak theoretical throughput of the device indicates how much room for improvement there is for the kernel.

5.2. Maximize Utilization

To maximize utilization the application should be structured in a way that it exposes as much parallelism as possible and efficiently maps this parallelism to the various components of the system to keep them busy most of the time.

5.2.1. Application Level

At a high level, the application should maximize parallel execution between the host, the devices, and the bus connecting the host to the devices, by using asynchronous functions calls and streams as described in Asynchronous Concurrent Execution. It should assign to each processor the type of work it does best: serial workloads to the host; parallel workloads to the devices.

For the parallel workloads, at points in the algorithm where parallelism is broken because some threads need to synchronize in order to share data with each other, there are two cases: Either these threads belong to the same block, in which case they should use __syncthreads() and share data through shared memory within the same kernel invocation, or they belong to different blocks, in which case they must share data through global memory using two separate kernel invocations, one for writing to and one for reading from global memory. The second case is much less optimal since it adds the overhead of extra kernel invocations and global memory traffic. Its occurrence should therefore be minimized by mapping the algorithm to the CUDA programming model in such a way that the computations that require inter-thread communication are performed within a single thread block as much as possible.

5.2.2. Device Level

At a lower level, the application should maximize parallel execution between the multiprocessors of a device.

Multiple kernels can execute concurrently on a device, so maximum utilization can also be achieved by using streams to enable enough kernels to execute concurrently as described in Asynchronous Concurrent Execution.

本文备注/经验分享:

这些都是老生常谈的一些东西: 首先上来说, 为了最大程度的发挥应用的性能,该应用应当尽量并行化, 而且将这种并行化映射到系统里面可用的各个部件(GPU计算,CPU计算,PCI-E传输,内存读写,显存读写,磁盘....),并同时让它们同时尽量可能的忙碌。注意NV这里同时提到了非GPU的部分(例如系统的内存/CPU/磁盘/甚至网络),(下一段也是这样提到的) 因为这些非GPU的部分,也对GPU应用的性能起到影响,甚至影响很大。 不过我们一般常说的是,如何在GPU上尽量利用GPU的各个部件,能分别利用起来它们,并让他们充分忙碌。 这里提到了使用流,其实之前手册里轻描淡写的暗示过一些,比如之前我们说过的,论坛也很多客户问过的, 使用3个流,将3个kernel(可以是同一个kernel,或者是不同kernel)的3次处理流水线话,形成: 流1:从内存传输到显存 ---> kernel计算 ---> 结果从显存传输回来到内存 流2: 从内存传输到显存 ---> kernel计算 ---> 结果从显存传输回来到内存 流3 从内存传输到显存 ---> kernel计算 ---> 结果从显存传输回来到内存 然后这样首先说, (1)这里利用了显卡的至少3个独立功能:正向传输,kernel计算,反向传输 (2)可以同时利用它们流1流2流3。 这样就符合本章节说的,首先能将你的程序并行化映射到多个部件上(这里至少有3个部件),其次能尽量同时使用它们,让它们同时忙碌起来。 所以你看,之前论坛常见的这种处理,就暗合了本章节的含义。所以往往能提高性能(至少kernel在进行中,下一个kernel的数据就开始准备了,缩小了延迟。甚至本kernel的结果正在回传不等完成,下一个kernel就已经开始了,也缩小了延迟。很多项目,减少时间就是金钱。这是很常见的一个优化手段。 这里提到,为了能并行的利用这种,往往需要使用多个流,这是因为CUDA的流是严格串行的,使用多个流才有可能得到GPU部件之间的并行关系。 需要说明的是,像是这样的同时双向传输,缩小时间,需要卡具有双向传输能力了,也就是Copy Engines >= 2 以前这个是专业卡专有的,家用卡只有1个copy engine,想能这样提高性能,必须上专业卡。但是从后期的maxwell+开始,这个特性已经放开了。后期的maxwell已经具有2个copy engines了。可以充分利用手册本章节的优化说明。而现在的专业卡,或者准专业卡(Titan-V)例如,甚至具有7个或者更多的copy engines, (从EE来的用户们可以将这个看成DMA Controller的channels),不仅仅可以双向,还可以充分服务较小的传输大小 + 较多较频繁的传输频次。

手册的本章节还说,如果将算法改写到GPU的过程中,如果线程之间要通信,或者要同步,尽量将范围缩小到block的内部,因为block内部可以代价较低的通过bar.sync操作(__syncthreads())来局部同步,完成通信。但如果实在这样改写不了,必须要全局同步的话,就只能结束前一个kernel,重新开下一个新kernel。这种代价较高,应当考虑在改写的过程中,尽量避免。(这也是为何论坛天天有人问,如何能全局同步,我们总是回答,重开kernel是唯一的全局同步方式。或者回答:前一个kernel死亡是唯一的全局同步方式的原因)。 需要说明的是,9.0+提供了协作式启动,可以在一定小规模的情况下,和其他的一些手册不说明的情况下(例如TCC驱动在Windows上),直接能完成全局同步。不需要重开kernel。这里的小规模的情况下是指至少能同时将需要同步的kernel里面的blocks都同时驻留在SM们上。

在pre 9.0的时代,这种规模的全局同步也可以变通的实现。例如在确保了所有的blocks都能同时驻留在SM上的时候(还记得啥叫驻留吧,前天刚说的)。可以通过维持一个global memory上的计数器,通过原子+1(当1个block执行到同步点后),然后反复通过volatile读取+自旋,等待kernel里的其他blocks都抵达此点。这种在满足了blocks能同时驻留的情况下,可以不开新kenrel,完成全局同步。 考虑到以前我们经常对客户说的一个优化,例如前些日子论坛某图像处理的:有一个问题,她是这么一个处理流程:用库读取图片文件到内存 ---> 传输内存图片内容到显存 ---> kernel进行图像处理 ....这里我们给客户分析到,尽量要并行一个: Host线程1:用库读取图片文件到内存 ---> 传输内存图片内容到显存 ---> kernel进行图像处理 ....Host Host线程2:用Host上的同步对象等待线程1完成前一章图片读取,并开始它的的传输到显存后。线程2用库读取图片文件2到内存 ---> 传输内存图片内容到显存 ---> kernel进行图像处理 .... 这样普通的大容量的机械硬盘,传输速度较慢(例如只有200MB/s),但在前一个Host线程开始传输给显存后,立刻就地再次开始读取磁盘,这样让磁盘开始充分忙碌。有利于防止GPU端kernel的启动之间形成气泡,在时间轴上阻隔开。然后用户这样处理后,重新查看运行时间,显著的缩短了。并查看GPU的时间轴(用profiler),kernel和kenrel启动之间的间隔明显缩短。这是一个典型利用。这个例子让应用能充分利用系统上的各个部件,而不仅仅是GPU设备内部的部件,却缩短整体该GPU应用的时间,减少了GPU设备上的kernel启动间隔,很有意义。 此外,本章节的最后那段,还说了如何充分利用设备上的SM本身,在现在所有被支持的卡上(3.0+),SM是可以同时执行来自多个kernel的warps和blocks的。如果有的时候,用户改写一个算法,因为算法本身的限制(或者用户本身能力的限制),就是无法充分改写成很多线程来做一次kernel启动,那么没有关系,用户可以保留小规模的kernel启动(例如都无法同时压满20个SM,例如在GTX1080上),那么用户可以选择同时再启动一个或者多个小规模的kernel,无论是本kernel,还是本应用中的其他处理过程的小kernel,现在的卡都具有并发kernel能力(concurrent kernel execution),这样可以在1个小kernel无法提供足够多的并行线程(parallel threads)规模的情况下,用来自其他kernel的线程充数上去,提高性能。这个特性也叫“任务并行”。

注意:虽然CUDA不像OpenCL,能让我们启动只有1个线程的kernel这样来利用设备(因为CUDA在N卡上不具有良好的串行执行能力),但我们依然可以通过组合多个少量并行程度的小kernel,来提高整体并行性能。这点其实很重要的。 NV自带的samples中,也有一个并发kernel的例子。原本是8个小kernel,拼凑到一起后,执行之间只有1/8.这个例子可见一斑。

好了。本章节说的各种组件,级别上的并行化的考虑都是很基础的,用户如果不能自己买较好的硬件配置,必须努力记住它们。越差的硬件需要的用户努力越多! 举个例子说,有个重要的优化是cudaMemcpyAsync*(),在正确了使用了它(需要两个条件:真正的page-locked的内存,和正确使用流)会提高传输速度。同时会减少内存到内存(你没看错,中间有个staging的过程,虽然是最后内存到显存的,但中间会导致内存到内存)的读写。如果老板买了好机器,双路+4通道都插好。那么在传输的时候,该客户使用普通的无async的cudaMemcpy,可能就和买了差机器,只有1个通道内存的客户,使用cudaMemcpyAsync的成绩一样, 客户完全减少了学习和使用异步传输的成本。而甚至买了差机器的用户,就算用来cudaMemcpyAsync,也性能很差,特别是多卡的时候(瓶颈卡内存到内存自己了)

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

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

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-06-05

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏程序员八阿哥

王老板Python面试(3): 一个初级python web后端开发工程师的面试总结

之前一直在做C++的MFC软件界面开发工作。公司为某不景气的国企研究所。(喏,我的工作经验很水:1是方向不对;2是行业有偏差)。

682
来自专栏牛客网

测试开发:面试真题+学习准备一

岗位:测试开发 我的秋招已经接近尾声了,目前拿到了滴滴,58,链家,电信it研发中心的offer,在等头条、腾讯和百度的结果。秋招面了的公司除了这7家还有搜狐和...

4317
来自专栏CDA数据分析师

Python告诉你:单词软件火了,但真的有那么多人在背单词吗?

0x00 前言 你想知道背单词软件有大概多少人注册第一天都没有背完嘛? 你想知道背单词软件这么火,这么多人在使用,真的有多少人真的在背诵嘛? 别急,Python...

1817
来自专栏吉浦迅科技

DAY62:阅读Glossary

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第62天,我们正在讲解CUDA C语法,希望在接下来的38天里,您可以学...

813
来自专栏牛客网

前端工程师:电信专业转前端是如何拿到阿里、腾讯offer的?

1.个人情况 ● 211本科 985硕士 电信专业 女生 ● 16年3月开始学习前端 ● 16年7月开始实习,共五家实习经历(不是特别厉害的厂) ● 秋招拿到两...

3296
来自专栏牛客网

深信服提前批面经

4讲一下hdfs的写入过程,发请求给namenode,返回一个地址进行写入,写入完告诉namenode,namenode完成副本备份。

1083
来自专栏牛客网

便利蜂内推电话面

昨天,突然接到了便利蜂内推成功约面的电话,要我今天去面试,这是春招头一个面试电话呀,很是激动,这得感谢牛客那个便利蜂发了内推邮箱的不认识的小哥哥的帖子吧。 结果...

3638
来自专栏大数据挖掘DT机器学习

python利用结巴分词做新闻地图

应用语言学的期末Pre花了差不多一个月零零碎碎的时间完成了。最初的打算爬取网易、新浪、腾讯的国内新闻,再通过提取关键词,比较这三个网站社会新闻报道的内容的倾...

3544
来自专栏带你撸出一手好代码

细说10月24号为什么是程序员的节日?

今天是10曰24日,有人把这个日子定为程序员的节日,因为1024这个数字和程序员密切相关。 下面我就为大家解密,1024跟程序员有什么关系,程序员写程序又到底是...

2694
来自专栏牛客网

19届前端实习生面经

1170

扫码关注云+社区