请问大家，我需要三个一维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？