用户
 找回密码
 立即注册
Shuai_Shuai 该用户已被删除
发表于 2013-6-10 09:52:43
3222324
我的kernel函数需要在执行的过程中让所有的thread同步,但使用“__syncthreads();”语句只能做到block内部的同步,请问如何做到block之间的同步呢?

目前想到的办法,就是把kernel函数在每次需要同步的地方断开,分成一个个小kernel,由主机端分别调用。
但这样做就必须先把已经加载到shared memory中的数据存回global memory,调用下一个小kernel时再重新加载回去,似乎做了很多重复劳动,请问还有什么更有效地办法吗?


另外弱弱的问一下,这个论坛的搜索功能在哪?
预计这个问题应该有人问过,想直接搜一下,可没找到搜索的选项……
使用道具 举报 回复
发表于 2013-6-10 10:43:28
block間的同步只能藉由kernel的啟動、結束來同步。

block間是無法同步的,因為每個block不是同時執行的,所以不能同步。

除非除非是只有很少的block才能同步,請google這篇論文

Inter-Block GPU Communication via Fast Barrier Synchronization

但能多少個block可以同步,這跟硬體的架構是有關係的

論壇文章編號6740對於這個問題有一些討論
使用道具 举报 回复 支持 反对
发表于 2013-6-10 13:39:22
LZ您好:

2# iHakka网友正解!一并感谢iHakka网友回馈论坛!

以及1#提到必须重新读入shared memory的问题,确实是这样的,必须重新读回,除非您能修改算法什么的。

论坛以前一直没有搜索功能,后来好像添加了,但是我也一时找不到地方了。
建议您使用google搜索代替,比如搜索 TDR,那么在google里面输入 TDR site:cudazone.nvidia.cn,这样临时应急好了。

同时也希望论坛支持团队能更好地完善论坛功能,造福网友。

祝您一切顺利~
使用道具 举报 回复 支持 反对
发表于 2013-10-27 19:12:13
正好需要这个问题的答案,感谢楼上几位朋友.
使用道具 举报 回复 支持 反对
发表于 2013-10-28 09:34:55
ice 发表于 2013-6-10 13:39
LZ您好:

2# iHakka网友正解!一并感谢iHakka网友回馈论坛!

版主,我借用这个楼问一下:
楼上兄弟说:  block間是無法同步的,因為每個block不是同時執行的,所以不能同步。
那么,有没有这种极端的情况: 一共有3个block, GPU先执行 block2 -> block1 -> block0.
还是说, GPU一定是按block0 -> block1 -> block2 执行, 只不过有可能不是同时执行.
多谢了,我最近在研究如何划分block,搜索时发现这个帖子.
使用道具 举报 回复 支持 反对
发表于 2013-10-28 10:03:25
dairlom 发表于 2013-10-28 09:34
版主,我借用这个楼问一下:
楼上兄弟说:  block間是無法同步的,因為每個block不是同時執行的,所以不能同 ...

楼主您好,这个无法保证。

如果您至少有3个SM, 每个SM上都有1个blocks,那么此时,显然无法也不应该保证任何执行顺序。
(因为原因很简单,如果保证了严格的0->1->2的block顺序执行,那么显然同一时刻只有1个SM在工作,另外2个必须空闲。这显然是不合理的)

您觉得呢?
使用道具 举报 回复 支持 反对
发表于 2013-10-28 10:31:48
玫瑰幻想 发表于 2013-10-28 10:03
楼主您好,这个无法保证。

如果您至少有3个SM, 每个SM上都有1个blocks,那么此时,显然无法也不应该保证 ...

嗯,有些感悟,多谢玫瑰...
使用道具 举报 回复 支持 反对
发表于 2013-10-28 10:39:22
玫瑰幻想 发表于 2013-10-28 10:03
楼主您好,这个无法保证。

如果您至少有3个SM, 每个SM上都有1个blocks,那么此时,显然无法也不应该保证 ...

追问一下:
是这样,如果我打算给GPU分配N个 blocks, 但是这个N很大很大...超过了GPU能一次性同时执行的blocks数量, 我想这种情况下, GPU肯定只能分批次执行这 N 个blocks(这样说对吧?). 假如GPU把这N 个blocks 分成了2个组, N1(block0 - blockm), N2(blockm+1 - blockN).
我问题是: GPU在执行时, 是随机执行(N1,N2)组呢? 还是按照block编号的顺序, 先执行N1, 再来N2?
多谢啊...
使用道具 举报 回复 支持 反对
发表于 2013-10-28 10:48:09
本帖最后由 横扫千军 于 2013-10-28 11:03 编辑
dairlom 发表于 2013-10-28 10:39
追问一下:
是这样,如果我打算给GPU分配N个 blocks, 但是这个N很大很大...超过了GPU能一次性同时执行的blo ...

楼主您好,目前是这样的:

首先上去要平铺满所有的SM, 编号为N的block分别在编号为N的SM上,这样如果一圈下来,SM上还能继续上BLOCK, 则2N的在N号SM上,2N+1的在N+1号SM上,依次类推。

当执行中,有个SM上的某block结束了,则给此SM继续上block,但是此时编号是什么规律,我没有做过测试(猜测是和上面的一样?),NVIDIA也没有公开过资料,就不能说了。

以及,任何情况下,这个行为都是无保证的,这个测试是当年在某卡上做的,不能保证将来此行为不改变,因为要用靠Block在SM上的顺序来使用任何的部分blocks间同步的措施都是不安全的,因此强烈建议您不要这么用。

唯一安全的方式是重开个kernel,并且总是建议这样。
感谢您的来访。

EDIT:修正了几个错别字。
使用道具 举报 回复 支持 反对
发表于 2013-10-28 14:25:50
玫瑰幻想 发表于 2013-10-28 10:48
楼主您好,目前是这样的:

首先上去要平铺满所有的SM, 编号为N的block分别在编号为N的SM上,这样如果一圈 ...

好的,我就是需要Blocks之间妥同工作.
您的讲解清晰明了,受益良多,玫瑰V5~~
使用道具 举报 回复 支持 反对
123下一页
发新帖
您需要登录后才可以回帖 登录 | 立即注册