CUDA性能优化—-warp深度解析 – 樂不思蜀的日志

把持流申请有特殊教育必要通常在各式各样的训练文风中找到。,GPU支援习俗、C作风的显式把持流体系结构,比如,假如…else,for,附加物。

CPU有复杂的五金器具设计可以罚款的做支管预测,指前面提到的事物是预测敷用药将要走的路程支管。假如预测是优美的的,因而CPU只会消费很小。与CPU比拟,GPU不一样的支管预测这么复杂。

这执意成绩哪里。,因缠住就是同任一warp正中鹄的thread只好落实同卵的的指导原则,假如这些线索不期而遇把持流申请有特殊教育必要,,假如你进入另任一支管,异样的时代,要缺点在落实的支管,其他的的树枝都被受监护人了。,对功能有很大冲撞。这类成绩执意warp divergence。

当心,warp divergence成绩只会发作在就是同任一warp中。下图展现了warp divergence成绩:

成功最适度功能,就必要防止就是同任一warp在不一样的落实路程。有很多办法可以防止同样成绩。,去的的情境,比如,万一有两个支管,支管的断定使适应是线索惟一的ID的平价。,内核有或起作用列举如下):

经过化验一下子看到两个核有或起作用在功能上是切近的。,你霉臭远超过预期的为什么二者都是同卵的的。,说起来,这是因本人的行为准则去复杂。,可以预测,CUDA编辑者将自动手枪帮忙最佳化本人的行为准则。。略谈GPU支管预测,这边,称为预测器变量的事物将被设置为1或0。,缠住支管将被落实,只仅仅预测变量是1。,支管将被落实。当使适应缺少本人界限值时,编辑者会将任一支管指导原则掉换为预测指导原则。)从此处,如今回到自动手枪最佳化成绩,一节较长的行为准则就可能性会造成warp 使发散成绩。

4、Resource 分区(资源分区)

任一warp的context包孕以下三把正式送入精神病院:

  1. Program counter
  2. Register
  3. Shared memory

再次重申,在同卵的的落实左右包装不消费切换。,因在十足warp的性命期内,SM处置的每个warp的落实context都是“on-chip”的。

每个SM在登记簿金中都有任一32位登记簿集。 file中,和通过作弊预先安排好结果的等同的共享 memory,这些资源都是用线索分段的。,因资源对公众不完全开放的,因而,假如线索的等同更多,过后每个线索使全神贯注更少的资源。,在另一方面,假如线索的等同较不重要的。,每个线索使全神贯注更多的资源。,这必要争辩本人的必要停止均衡。。

资源限局限到处钐中寓居的BLCOK等同,不一样GPU,register和shared 内存的等同也不一样。,就像费米和开普勒环形山建筑风格的分别。假如缺少十足的资源,内核启动衰退。下图是CAL知识参量的比拟。:

当块成功十足的资源时,就变得

active block

。block正中鹄的warp就称为

active warp

。active warp又可以被分为上面三类:

  1. Selected warp
  2. Stalled warp
  3. Eligible warp

在钐中warp调解器每个整套城市主动语态 warp送去落实,任一被选正中鹄的warp称为Selected warp,没被选中,只预备落实的东西叫做Eligible warp,不预备落实的东西叫做Stalled warp。warp套装落实必要填写上面两个使适应:

  1. 32个CUDA 激励是收费的
  2. 容易指导原则的缠住参量都已预备愿意。

比如,开普勒环形山建筑风格GPU在什么竞选运动时期 warp数量只好少于或整个含义64个。selected warp数量只好缺少或整个含义4个(因scheduler有4个?无法断定,至若4无论太少,别撕咬,内核启动前,将会有任一预热推拿,可运用CUADFRE变卖。假如任一warp闭塞了,调解器将选择任一合格的 warp预备去落实。

CUDA PR中应注重计算资源的分派:这些资源限局限了竞选运动。 warp的等同。从此处,本人只好攫取五金器具的某些边界。,为了最大限地应用GPU,本人只好最大限地运用积极作用。 warp的数量。

5、Latency 遮住(推延遮住)

从开端到完毕消费的仪表 整套高处指导原则潜育期。。当每个流通都契合使适应时 warp被调解时,计算资源将充分应用,因为此,本人就可以将每个指导原则的latency遮住于issue其它warp的指导原则的奔流中。

与CPU训练比拟,latency 遮住朝一个方向的GPU去重要。CPU 内核被设计为最低限度一到两个线索推延。,只GPU中线索的等同并不一样的两个这么复杂。。

当关涉指导原则推延时,指导原则可分为以下两种:

  1. Arithmetic instruction
  2. Memory instruction

望文生义,Arithmetic  instruction 推延是算术运算的开端和完毕间距。。备选的是装货或蓄电的开端和完毕间距。。

二者都的潜育期约为:

  1. 10-20 cycle for arithmetic operations
  2. 400-800 cycles for global memory accesses

上面的情节是任一复杂的落实奔流,当warp0闭塞时,落实其他的的warp,当warp变为eligible时重行落实。

 

您可能性远超过预期的方法评价竞选运动 warps 皮数 latency。Little’s 法度可以供应有理的报价:

用于算术 operations来说,一致性可以用任一遮住来体现。  Arithmetic 潜育期推拿次数。下表显示费米和开普勒环形山建筑风格师的相干记载。,这是(A) + b * c)作为推拿的任一样板。不一样算术指导原则,生产能力(吞吞吐吐)也不一样。。

这边的生产能力限制为每个SM每整套的推拿数。。鉴于每个warp落实恒等的种指导原则,从此处每个warp对应32个推拿。因而,费米,每个SM必要640/32=20个warp来留在心中计算资源的充分应用。这中间它。,arithmetic operations的一致性可以表达为推拿的数量或许warp的数量。

二者都暗中的相干也对应于放PA的两种方法。:

  1. Instruction-level 一致度(ILP):在恒等的线索中有更多孤独指导原则
  2. 线索级 Parallelism (TLP):更多同时契合使适应 threads

为了记得 operations,一致性可以体现为每个整套的八位位组数。。

因记得 生产能力前后是GB/SEC的单位。,本人必要率先停止中肯的的替换。。您可以经过以下阐明检查知识的内存。 frequency:

$ nvidia-smi -a -q -d CLOCK | fgrep -A 3 “Max Clocks” | fgrep “Memory”

以费米建筑风格为例,其memory 频率可能性是,开普勒环形山列岛。过后替换奔流是:

乘以同样92可以记下74在上的。,这边的数字是十足知识的,缺点每任一SM。

用这些记载,本人可以做某些计算。,以费米建筑风格为例,万一每个线索的税收是获名次任一浮点法数(4)。 八位位组)因为大局的记载类型 内存进展到SM停止计算,你必要大概18500个线索,也执意579个warp来遮住缠住的memory latency。

费米有16米,因而每个SM必要579/16=36个warp来遮住memory latency。

6、拨款(拨款)

当任一warp闭塞了,SM将落实另任一资历 warp。抱负的情境是,一向确保激励被使全神贯注。拨款是每个SM的主动语态 warp占最大warp数量的攀登:

本人可以运用cuda库有或起作用的办法来获取warp最大数量:

cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop, int 安顿)

过后用maxThreadsPerMultiProcessor成功假设值。

网格和块的分配额基准:

  • 确保块正中鹄的线索数是32的连锁商店。
  • 防止块太小:每个BLCOK至多有128或256个螺旋环
  • 争辩内核所需资源调解块
  • 确保块的等同极大于SM的等同
  • 做更多的试验来找出最适度的分配额

Occupancy专注于每个在钐中可以一致的thread或许warp的数量。不管怎样,拨款缺点惟一的的体现指示,当拨款走到必然值时,的比较级的最佳化可能性不再无效。,静止的数不清的其他的指示必要调解。。

7、使时间互相一致(使时间互相一致)

使时间互相一致是一致训练正中鹄的常见成绩。。在CUDA中,使时间互相一致有两种方法:

  1. System-level:等候负责人和知识的整个任务填写
  2. Block-level:等候知识正中鹄的缠住线索落实到本人点。

因CUDA API和负责人行为准则是异步的,cudaDeviceSynchronize可以用来中断CPU等候CUDA正中鹄的推拿填写:

cudaError_t cudaDeviceSynchronize(void);

因块正中鹄的线索落实次缺点通过作弊预先安排好结果的的,CUDA供应了任一在块中使时间互相一致线索的有或起作用。。

发表评论

电子邮件地址不会被公开。 必填项已用*标注