能不能对特定线程同步?

__syncthreads()是对一个boock中的所有线程做同步,是否存在其他的函数对选定的线程做同步?或者一些做法?
谢谢!

对“选定的一些线程同步”是要实现什么行为呢?
如果LZ熟悉汇编的话,不妨看下PTX手册(cuda toolkit自带),里面关于“bar”的用法,也许能满足您的要求,当然,也许不能。

祝您编码愉快~

谢谢你的回答。
比如一个block有32个线程,用前16个执行操作A,然后想保证前16个线程都执行完操作A,好接下来执行操作B;同样后16个线程执行操作C,也想保证都执行完后执行操作D。
这样的情况,如果在两个不同的block中,只需要在AB间和CD间加__syncthreads()即可,但是我现在需要在同一个block中实现。

如果您要使用一个block 32个线程的话,那么这32个线程是在一个warp里面的,你可以利用一个warp内部线程是同时完成的特性。

global void ds32(…)
{ if (前16个线程)
do A;
else
do C;
if(前16个线程)
do B;
else
do D;

}

这样估计可以,因为一个warp内部,线程是同时完成的,第一个if的时候,前面16个线程在do A,而后面16个线程在do C,逻辑上他们是同时完成的。然后才进入下一轮if,前16个线程do B,后面16个线程do D,当然此时前面的A,C都是完成的了。

如此可以实现您叙述的要求。

但是,必须指出,一般而言一个block 只有32个线程的话,会导致occupancy太低,无法发挥效能。

假如您一个block不是32个线程,而是比如说512个线程,要使前256个线程先干A,再干B;后面256个线程先干C,再干D。

那么可以如下方式使用__syncthreads()

global void ds512 (…)
{ tid=threadIdx.x;
if (tid<256)
do A;
else
do C;
__syncthreads();
if (tid<256)
do B;
else
do D;
__syncthreads();


}

这样可以保证前256个线程先do A,然后do B;后256个线程先do C,然后do D。

以及同时,__syncthreads()只能用于同步一个block内部的线程,如果您有两个block,是不能用__syncthreads()来维护两个block之间的等待/先后关系的。当然如果您叙述中的意思是block 0完成do A,do B,用__syncthreads()来维护完成A之后再完成B,(block 1,C,D类似)这是可以的。

祝您编码愉快~

谢谢斑竹。您回答的已经很细致、全面。不过我的问题应该还涉及到递归的问题,我大概想到有个for循环(内嵌_syncthreads())来解决。我先好好想想,有问题再问。

万分感谢!

好的,不客气,您可以先自己揣摩一下。

祝您编码愉快~

另外,需要强调的是,warp是一个整体,只能对warp同步而不能对某个特定线程同步,这是对您标题问题的直接回答。上面几楼的实现实际上使用了分支。