cuda中if-else的效率问题

xxjjs 2009-07-22 07:18:38

加精

据说cuda中if-else的效率很不好，
还据说cuda中if-else的执行时间相当于各分支展开的总执行时间
但我做了一个测试，很不解ing

测试平台：GTX280 （27SM）
cuda设置：27×8 blocks，每blocks 512 threads，每threads运行128循环，循环中包括以下代码，共计运行27*8*512*128约1400万次

注：代码中的什么c[0]c1[1]j等大家就当看笑话，是垃圾数据



			if (j >= 0) 

				{ 

					if (c[1] < c[3])                                                        

					{

						c[0] = 1 - c[0];

						c[1] = c[2];

					}                                                               

					else                                                              

					{

						c[2] = c[3];

					}                                                               

				}                                                                   

				else                                                                  

				{ 

					j += c[0];

					if (j >= 0) 

					{ 

						if (c[1] < c[3])                                                        

						{

							c[2] = c[1];

						}                                                               

						else                                                              

						{

							c[0] = 1 - c[0];

							c[3] = c[2];

						}                                                               

					}     

				}

这样的代码执行约需136ms

然后我将所有if-else注释掉成为：



					{

						c[0] = 1 - c[0];

						c[1] = c[2];

					}                                                               

					{

						c[2] = c[3];

					}                                                               

				}                                                                   

				{ 

					j += c[0];

					{ 

						{

							c[2] = c[1];

						}                                                               

						{

							c[0] = 1 - c[0];

							c[3] = c[2];

						}                                                               

					}     

				}

那么如果cuda是执行所有分支的话，现在的速度会大大加快，因为至少少了最后选择结果的步骤，可是这样测试了一下，速度反而降低为

问题：

1. 这是怎么回事？cuda中ifelse如何处理？
2. if-else多的话有什么提高效率的途径？上述代码在CPU上跑同样1400万遍单线程也只要344ms，有个4-cores的就比cuda快了

...全文

2650 17 打赏收藏转发到动态举报

写回复

用AI写文章

17 条回复

切换为时间正序

请发表友善的回复…

发表回复

henrynus 2010-01-27

打赏
举报

每个stream multi-pocessor 在同一时刻执行一个warp中的32个threads，执行的内容是将一个instruction 分配到8个threads，连续四次。所以如果一个warp中没有实际的branch，就不会浪费时间。如果分配的合理，那么整个process将不会有分支，因为每32个threads都在一个方向执行。

hello_hi_hi 2010-01-24

打赏
举报

指令数不一样，第一例中整个过程最多6条指令，最少为3条指令，而第二个例子中必须执行7条指令，这肯定有差别，如果你要做比较，那么最好用指令数一样的。

彭令鹏 2009-12-30

打赏
举报

呵呵，按照你给出的谓词优化定义——“@p cmd.p成立执行cmd,否则,如同执行nop.”
那楼上的代码应该是没怎么谓词优化，因为代码中的第二种情况就是谓词优化的相似情况^_^（所有指令都执行了遍）
你说的第2）种情况估计是主要原因~
[Quote=引用 3 楼 l7331014 的回复:]
引用楼主 xxjjs 的帖子:
1. 这是怎么回事？cuda中ifelse如何处理？

1)
在cuda中有谓词执行语句!
LZ程序中,分支执行的指令很少,被nvcc优化成谓词执行了吧.
2)
cuda中if-else的效率很不好，
是指如果在一个warp中的线程执行不同的路径.
如是相同的路径.(估计LZ的程序就是这种情况),效率并不差!
[/Quote]

ILoveJUV 2009-08-25

打赏
举报

[Quote=引用 8 楼 cyrosly 的回复:]
简单点，举例说在线程ID用在条件中时使用无分歧代码大概有如下形式：
if(threadIdx.x==const num){ ... }//这里的const num只是说在当前是，但可以是是个变量，比如一~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~(这个为什么是无分歧的呢，只有特定的thread取值才为true，其他的都为false，这样理解对吗？？)
个在每次循环都计算的变量，但每遍中对WARP内的所有线程参考不变
or
if(threadIdx.x <sharedcc[warpid]){ ... }
or
if(threadIdx.x <threadIdx.x/WARP_SIZE){ ... }
[/Quote]

YHL27 2009-07-28

打赏
举报

switch效率高，古怪if-else实用性强过

kuangquansheng 2009-07-27

打赏
举报

关注中。顶一下。

无心人_过过小日子 2009-07-26

打赏
举报

[Quote=引用 10 楼 darkstorm2111203 的回复:]
谓词执行是几个分支编译成一块，然后用谓词决定线程是否执行某条指令么？
[/Quote]

@p cmd.p成立执行cmd,否则,如同执行nop.

darkstorm2111203 2009-07-26

打赏
举报

谓词执行是几个分支编译成一块，然后用谓词决定线程是否执行某条指令么？

darkstorm2111203 2009-07-26

打赏
举报

这个程序的执行时间和j是什么有比较大的关系

分支应该会比两个分支路径都执行要好一点的
有时人品好可以全走一个分支

Cyrosly 2009-07-26

打赏
举报

[Quote=引用 7 楼 l7331014 的回复:]
引用 5 楼 weilaishijie2008 的回复:
• N条指令路径→1/N throughput
 只需要考虑同一warp即可，不同warp的不同的指令路径不具相关性
 G80上使用指令预测技术加速指令执行
即使没有硬件分支预测，但是指令预测+编译器优化我想在一定程度上可能减…

准确说是编译器(程序员)负责转移的正确性.
其实,在ptx中有一类无分歧转移指令,即一旦转移,warp中所有线程不再执行临外一条路径.
[/Quote]

无分歧指令并不只是在PTX中，是否会产生分歧分之路径要看同一个WARP内的条件码是否共享还是数据独立。

简单点，举例说在线程ID用在条件中时使用无分歧代码大概有如下形式：
if(threadIdx.x==const num){ ... }//这里的const num只是说在当前是，但可以是是个变量，比如一个在每次循环都计算的变量，但每遍中对WARP内的所有线程参考不变
or
if(threadIdx.x<sharedcc[warpid]){ ... }
or
if(threadIdx.x<threadIdx.x/WARP_SIZE){ ... }

weilaishijie2008 2009-07-23

打赏
举报

可能是我上面说的不太准确。这个是清华大学的讲义内容：
控制流(Control Flow)
 同一warp内的分支语句可能执行不同的指令路径
 不同指令路径的线程只能顺序执行
• 每次执行warp中一条可能的路径
• N条指令路径→1/N throughput
 只需要考虑同一warp即可，不同warp的不同的指令路径不具相关性
 G80上使用指令预测技术加速指令执行
即使没有硬件分支预测，但是指令预测+编译器优化我想在一定程度上可能减少一些由于分支带来的影响吧。

无心人_过过小日子 2009-07-23

打赏
举报

[Quote=引用 5 楼 weilaishijie2008 的回复:]
• N条指令路径→1/N throughput
 只需要考虑同一warp即可，不同warp的不同的指令路径不具相关性
 G80上使用指令预测技术加速指令执行
即使没有硬件分支预测，但是指令预测+编译器优化我想在一定程度上可能减…
[/Quote]

准确说是编译器(程序员)负责转移的正确性.
其实,在ptx中有一类无分歧转移指令,即一旦转移,warp中所有线程不再执行临外一条路径.

无心人_过过小日子 2009-07-23

打赏
举报

[Quote=引用 4 楼 xxjjs 的回复:]

不是据说GPU不进行分支预测吗？看很多文档如是说
[/Quote]

不进行分支预测.

xxjjs 2009-07-22

打赏
举报

[Quote=引用 2 楼 weilaishijie2008 的回复:]
问题1：cuda中分支展开执行，是串行执行过程。带有if-else程序虽然有分支但是分支内语句要比没if-else程序少，这有可能使时间相差不大，并且现在的分支都有预判优化。

[/Quote]

不是据说GPU不进行分支预测吗？看很多文档如是说

无心人_过过小日子 2009-07-22

打赏
举报

[Quote=引用楼主 xxjjs 的帖子:]
1. 这是怎么回事？cuda中ifelse如何处理？
[/Quote]

1)
在cuda中有谓词执行语句!
LZ程序中,分支执行的指令很少,被nvcc优化成谓词执行了吧.
2)
cuda中if-else的效率很不好，
是指如果在一个warp中的线程执行不同的路径.
如是相同的路径.(估计LZ的程序就是这种情况),效率并不差!

xxjjs 2009-07-22