请问大家,我需要三个一维block,每个block含1000个thread。三个block的执行不同的计算。那么我对任务进行划分,可不可以让线程的索引不用 block索引 + thread索引来区分任务?比如,我需要block3里的计算数据,不用block的索引,就可以选出block3里的1000个thread啊?还有,对划分的任务进行处理时,除了用“if(相应线程索引){ ………… }”对3个block进行任务划分外,还有其他方法么?
谢谢啊!
因为CUDA是SIMT的形式,所有的threads跑的都是同一段代码,所以如果要分开执行不能的内容分支,就得知道这一堆线程谁是谁,而判断线程谁是谁就得使用线程的索引(当然也可以部分使用)。同时线程索引也可以用来辅助寻址。
至于您说的3个block需要执行不同的任务,那么假定3个block也是一维x方向排布的,那么您可以if blockIdx.x==1,2,3这样直接用block的索引分支,而不用具体到某一个线程。
因为,所有同一个grid里面的线程都是跑的同一个kernel,所以kernel里面必须包含所有分支的代码,而不能只包含一部分。所以,我一时还没想到有什么其他的方法。不过有个变通的方法就是开3个小的kernel,分别对应以前的一个block,然后用kernel并发的方法实现。不过可能更麻烦一些,效果也不见得好。
祝您编码愉快。
第一个问题,你可以把数据存在shared memory 中。这样,就可以每个block利用threadId访问自己的数据,不需要获取自己的blockId,而block之间又不会有冲突。
对于第二个问题,如果你指的“任务的划分”中的“任务”是一连串的顺序执行的指令序列,而“划分”是指让不同的线程跑不同的指令序列,那么除了使用条件分支语句逻辑上没有其他可能。
就上面的问题,假如我要使用第2个block里的线程,那么我这么定义:if(blockId.x == 2)
{ i = threadIx.x; ……………… } 是不是就意味着i的索引值是0——999呢?而不是在整个Grid里的
1000——1999?
在请问一下,如何才能给属于不同的block的shared memory分配数据啊?
[
上面的问题,threadIdx是block内线程索引,所以1000个线程索引从0到999
第二个问题,shared memory只对block可见,一个block是无法看到另一个block的shared,所以需要用线程的全局索引来获取不同的数据也就是必须用blockIdx.x*blockDim.x+threadIdx.x了
global void ice (…)
{
if blockIdx.x==0
{ // for the first block
int i=threadIdx.x; //i 的取值范围是 0~999,这里是对应第一个block里面的线程,这些线程总的索引是0~999
do something
}
if blockIdx.x==1
{ // for the second block
int i=threadIdx.x; //i 的取值范围是 0~999,这里是对应第二个block里面的线程,这些线程总的索引是1000~1999
do something
}
if blockIdx.x==2
{ // for the third block
int i=threadIdx.x; //i 的取值范围是 0~999,这里是对应第三个block里面的线程,这些线程总的索引是2000~2999
do something
}
return ;
}
dim3 blockdim (3);
dim3 threaddim (1000);
int main()
{
…
ice<<<blockdim,threaddim>>>(…);
…
}
大致是这个样子。
您可以用blockIdx.x区分是哪个block,从而执行不同的任务。在block内部,只需要区分1000个线程,只用threadIdx.x就可以,当然为了寻址方便等,您也可以使用完整的索引。
祝您编码愉快。
再请问一下,如果这样写,可不可以给不同的shared分配数据啊,还是上面的3个block。
if( blockIx.x == 1 ){ 给shared 赋值 } …………if( blockIx.x == 3 ){ 给shared赋值 }
谢谢您如此详尽的回复! 就您刚刚写的程序,假如我第1个block里的第5个thread要使用第3个
block里的第88个thread的计算值,那么我应该怎么把block3里的thread88的值传给block1里的
thread5呢?可以直接把值附过去么?
可以的。
block之间只能通过global memory通信了,但是不像同一个block内部可以用__syncthreads(),block之间的先后执行顺序是没有保证的,所以比较难于处理。
谢谢!也祝您愉快。
客气客气,如果有具体需求,可以随时贴上来具体问题具体分析。
欢迎莅临cuda zone~
再请问一下,我如何才能将我的compute和sm默值10改成30呢?
早好!您说block间的执行顺序没有保证。如果我这3个block在计算时需要相互传输数据,那
么是不是只有让三个block执行串行计算,才能保证他们正确地传输数据啊?有办法让他们三个
block并行计算么?
block本身就是并行计算的,而且你不能控制哪个先哪个后,也不能控制block里面哪个warp先,哪个warp后。所不同的是block内部可以__syncthreads(),让所有的threads都停在同一个地方,然后再继续,但block之间没有这种机制。一般地,要各个block停在同一个地方,只能结束当前kernel,然后发射下一个kernel继续剩下的工作。
我觉得可以如此考虑,某些GPU只有一个SM,block是加载在SM上的,而你的block比较大的话,一个SM上也只能加载一个block。此时如果block0需要block2结果才能继续,而block2又无法加载进去,程序就无法运行了。对于有多个SM的GPU,你也无法保证哪个block先算出结果。似乎有人手工做过在global里面加标记的方法实现多个block的同步,但是这种做法似乎不为官方所推荐(似乎是在不同的GPU上可能会锁死,比如一个SM的GPU),而且我也不知道其实现细节如何。
我觉得您可以将您的具体的原始的问题描述一下,看看有没有什么变通的解决方法,如果按照您给出的要在多个block之间传递数据,似乎不好实现。
感谢您早上莅临论坛抢占沙发,以上个人观点仅供参考,祝您编码愉快。
下面就是我需要用GPU处理问题的原始部分和一些描述,请您帮我看看block间怎样传输数据才好,同时实现两个block并行计算。
for ( e = 0; e <= TimeStop; e++ ) //大循环开始
{
t = e*dt;
Ein[Isource] = sin(t); //part1:产生数据,后面部分循环使用(这部分我打算再用1个block实现)。
for ( j = IncidentStart; j <= IncidentEnd - 1; j++ )
{
Hin[j] = Hin[j] - 0.5*( Ein[j+1] - Ein[j]);
}
for ( n = IncidentStart +1; n <= IncidentEnd - 1; n++ )
{
Ein[n] = Ein[n] - 0.5*( Hin[n] - Hin[n-1]);
}
//上面部分的Ein、Hin大小都为1000。每一次小循环之后得到它两的一次计算值。下面的小循环都要用的上面
Hin[Itmin-1]和Ein[Itmin]计算值。我打算用1个block实现上面的计算,下面的计算也用含1000个thread的block。
//part2:
//这部分循环的Ex、Hx也是1000的大小,我打算用用第2个block来实现,但这个block里的Ex[Imin] 要用到上面的
Hin[Itmin-1]计算值,而Hy[Itmin-1]要用到上面Ein[Itmin]计算值。每一次大循环,两个block都是相关的。
for ( i = Itmin + 1; i <= Itmax -1; i++ )
{
Ex[i] = FE1[ Ob[i] ][ Ob[i] ] * Ex[i] - FE2[ Ob[i] ][ Ob[i] ]*( Hy[i] - Hy[i-1] );
}
Ex[Itmax] = Ex[Itmax] - FE*( Hy[Itmax] - Hy[Itmax-1] );
Ex[Imin] = Ex[Itmin] - FE*( Hy[Itmin] - ( Hy[Itmin-1] + Hin[Itmin-1] ) );
for ( i = Imin + 1; i <= Itmin - 1; i++ )
{
Ex[i] = Ex[i] - FE*( Hy[i] - Hy[i-1]);
}
for ( i = Itmax + 1; i <= Imax - 1; i++ )
{
Ex[i] = Ex[i] - FE*( Hy[i] - Hy[i-1] );
}
for ( i = Itmin; i <= Itmax - 1; i++)
{
Hy[i] = FH1[Ob[i]][Ob[i+1]]Hy[i] - FH2[Ob[i]][Ob[i+1]](Ex[i+1] - Ex[i]);
}
Hy[Itmin-1] = Hy[Itmin -1] - FH*( Ex[Itmin] - Ex[Itmin-1] - Ein[Itmin]);
Hy[Itmax] = Hy[Itmax] - FH*( Ex[Itmax+1] - Ex[Itmax] );
for ( i = Imin; i <= Itmin -2; i++ )
{
Hy[i] = Hy[i] - FH*( Ex[i+1] - Ex[i] );
}
for ( i = Itmax + 1; i <= Imax - 1; i++)
{
Hy[i] = Hy[i] - FH*( Ex[i+1] - Ex[i]);
}
} //大循环结束
那么,在这种情况下,我应该怎么来实现两个block的并行呢?还有,这程序里要用到好多数据(针对某些线程),那么
我应该怎样才能给某一个thread分配register memory和local memory呢(其他线程就没有此数据,这样好不占用他们的空
间,以留出空间保存其他数据)? 非常感谢!
如无看错,您实现的是电磁场时域有限差分算法。
一般而言,要将源更新,电场更新,磁场更新分别写成三个kernel,依次循环调用。(当然,电场更新或者磁场更新也可以分别写成3个kernel,每个kernel算一个方向)
如此不仅计算规模可以轻易扩大(您的实现方法受到block大小限制,无法发挥出GPU的计算能力),而且可以解决您的一切关于block间数据依赖的烦恼。
快来试试看吧,亲~
另外,发较多的代码尽量用代码模式,否则容易丢失一些部分,而且会部分变成斜体,看起来不舒服。
ice版主,真是太感谢您的回复了,看了您这几句话,我真有醍醐灌顶之感啊!而且您眼力真好!难道您也做过FDTD?