warp内一条线程如何启动其他线程

  • 8 replies
  • 875 views
warp内一条线程如何启动其他线程
« 于: 十二月 15, 2019, 09:50:50 am »
CUDA10.0,2080ti,32个thread处理一组任务,任务的主体流程由第一条thread来走,计算密集的步骤,剩下的31条thread一同参与来解决。
即:
if(threadIdx.x==0){
function();
}
问题是:在function里如何调用剩下的31个thread,参与运算?
还是只能分开写?

Re: warp内一条线程如何启动其他线程
« 回复 #1 于: 十二月 16, 2019, 01:37:04 pm »
你可以思考你的计算是否可以通过warp primitives或者cooperative groups实现。参见CUDA C Programming Guide 附录B和C。这里包括线程束洗牌函数、选举函数和匹配函数等。另外,最传统的做法就是用共享内存了。

Re: warp内一条线程如何启动其他线程
« 回复 #2 于: 十二月 16, 2019, 01:40:14 pm »
再补充一下,如果不用动态并行,你应该不能“在一个线程调用其它线程计算”,只能是在一个线程使用共享内存,或者让一个warp内的线程合作计算。关于动态并行,参考手册的附录D。

Re: warp内一条线程如何启动其他线程
« 回复 #3 于: 十二月 16, 2019, 02:23:55 pm »
CUDA10.0,2080ti,32个thread处理一组任务,任务的主体流程由第一条thread来走,计算密集的步骤,剩下的31条thread一同参与来解决。
即:
if(threadIdx.x==0){
function();
}
问题是:在function里如何调用剩下的31个thread,参与运算?
还是只能分开写?

楼主你可以直接将思路颠倒过来即可:
your_function()
{
   {无任何if约束, 所有线程执行}
   if (约束条件) {某些线程执行}
}

而不是反过来,上去默认存在一个约束条件(例如你套在你的最外层的那个),然后试图抵消掉这个条件,这是做不到的。但更换一个等价写法,就可以了。
« 最后编辑时间: 十二月 16, 2019, 02:25:10 pm 作者 屠戮人神 »

Re: warp内一条线程如何启动其他线程
« 回复 #4 于: 十二月 17, 2019, 03:57:40 pm »
楼主你可以直接将思路颠倒过来即可:
your_function()
{
   {无任何if约束, 所有线程执行}
   if (约束条件) {某些线程执行}
}

而不是反过来,上去默认存在一个约束条件(例如你套在你的最外层的那个),然后试图抵消掉这个条件,这是做不到的。但更换一个等价写法,就可以了。
似乎并不可以,我的原函数比较冗杂,可以理解成这种结构:

子函数1()
{
xxxx
子函数2();//这里需要一个warp的32条thread一同工作
xxxx
}
function()
{
xxxx;
if(threadIdx.x==0)
{
子函数1();//这里为了避免warp内分支,只用threadIdx.x==0
}
xxxx;
}

Re: warp内一条线程如何启动其他线程
« 回复 #5 于: 十二月 17, 2019, 04:03:46 pm »
似乎并不可以,我的原函数比较冗杂,可以理解成这种结构:

子函数1()
{
xxxx
子函数2();//这里需要一个warp的32条thread一同工作
xxxx
}
function()
{
xxxx;
if(threadIdx.x==0)
{
子函数1();//这里为了避免warp内分支,只用threadIdx.x==0
}
xxxx;
}

不对显而易见的事情做出解答。

Re: warp内一条线程如何启动其他线程
« 回复 #6 于: 十二月 17, 2019, 04:12:19 pm »
不对显而易见的事情做出解答。
这逻辑实现不了的对吧(不存在,从thread.x=0的线程中调用整个warp剩下31条thread的编程方法),现在能想到的解决办法是把N层的子函数展开,在同一层实现,再加thread分支判断,这是不是唯一的实现方式?

因为这个工程量比较大,想问清楚了再动手。

Re: warp内一条线程如何启动其他线程
« 回复 #7 于: 十二月 17, 2019, 04:37:42 pm »
这逻辑实现不了的对吧(不存在,从thread.x=0的线程中调用整个warp剩下31条thread的编程方法),现在能想到的解决办法是把N层的子函数展开,在同一层实现,再加thread分支判断,这是不是唯一的实现方式?

因为这个工程量比较大,想问清楚了再动手。

你完全可以实现的。你可以将限制条件逐层传递的,这是可以完全做到,而且显而易见的。

in level 1:
int cond0 = ...;
....
if (cond0) do something; //not warp uniform
call_level_2(cond0); //warp uniform, passing cond0 along the call.

in level 2:
int cond1 = threadIdx.x < 123; //warp uniform
do something; //warp uniform
if (cond1 && cond0) do something; //utilizing our new cond1 while respecting the old cond0
call_level_3(cond1, cond0) ; //or call_level_3(cond1 && cond0)

in level 3:
do any thing you want; //warp uniform
if (exp && cond1 && cond0) //utilizing our new exp condition, while repecting the old cond1 and cond0

你看,这种情况下,可以在任意调用级别扩展到全warp一致执行,同时还能随时收缩到你想要的所有限定条件下,只需要你改变一下写法,将之前的所有调用层递的条件跟随传递,这样既可随时伸缩。

Re: warp内一条线程如何启动其他线程
« 回复 #8 于: 十二月 17, 2019, 04:46:37 pm »
你完全可以实现的。你可以将限制条件逐层传递的,这是可以完全做到,而且显而易见的。

in level 1:
int cond0 = ...;
....
if (cond0) do something; //not warp uniform
call_level_2(cond0); //warp uniform, passing cond0 along the call.

in level 2:
int cond1 = threadIdx.x < 123; //warp uniform
do something; //warp uniform
if (cond1 && cond0) do something; //utilizing our new cond1 while respecting the old cond0
call_level_3(cond1, cond0) ; //or call_level_3(cond1 && cond0)

in level 3:
do any thing you want; //warp uniform
if (exp && cond1 && cond0) //utilizing our new exp condition, while repecting the old cond1 and cond0

你看,这种情况下,可以在任意调用级别扩展到全warp一致执行,同时还能随时收缩到你想要的所有限定条件下,只需要你改变一下写法,将之前的所有调用层递的条件跟随传递,这样既可随时伸缩。
明白了,感谢~