假设如下原程序：
for (i=0; i<in+1; i++){
b(i)=f(i,c1,c2,…,cn);
}
for (i=0; i<in; i++){
a(i)=b(i+1)-b(i);
}
其中in>32。
将其改造成GPU计算，假设32threads/block。
那么第一个block的32个线程执行a(0),a(1),…,a(31)的计算，一 一对应。
但是需要b(0),b(1),…,b(31),b(32)的数据，b(i)的数据有33个，比32个thread多一个。
可以如下方案解决：在32个线程一 一对应地计算完b(0),b(1),…,b(31)之后，让threadIdx.x=31的线程单独计算b(32)。缺点是程序变长了。
看似问题不大，但是如果b(i)=f(i,c1,c2,…,cn)中的f(i,c1,c2,…,cn)特别复杂：包括很长的公式，调用其他参数，调用函数等等。就会使得程序增加将近一倍的行数。并且如果a(i)的表达式为a(i)=b(i-1)+b(i)+b(i+1)时，问题更严重。
有没有好办法啊？！
LZ您好，大致看了您的叙述，您的方法我大致解释和建议如下：
1：假定in是比较大的一个数，以及不是32的整倍数。这其实没什么的，这情况很常见，您只需启动大于in个线程，然后在kernel里面用if判断一下，让编号超过in的线程直接return不干活即可。
2：您对block的安排是不合适的，32个线程太少了，一般会导致GPU效率低到无法接受的程度。您可以启动多个线程，按照1：的说法计算即可，即便是真的只有33个线程，1：的方法也比您的方法好。所以请不要局限在一个block里面只有32个线程，然后想办法折腾。
3：根据您的代码，您的a(i)的值实际上是依赖于b()的，那么您必须等b()完全算好以后，才能开始计算a(),一般来说，这需要两次invoke kernel，第一次用来计算b()，第二次用来计算a()。在第一个kernel结束的时候，能保证b()完全计算好。
此时两个kernel，线程数量完全可以不一致，亦无纠结。
以及如果您一定要一个block搞定，那么请在b()计算语句后面加上__syncthreads(),以保证b()被更新完成之后再计算a();以及如果您真的打算只用一个warp来处理，那么可以不要__syncthreads(),因为一个warp总是同时完成的。
但是这两个方法都是不推荐的，因为如果数据量如此少，可能CPU搞定更快。
大致如此，请您参考。
祝您编码顺利~
太感谢版主了，这是我第一次发帖，居然这么快就有如此详细、专业的回复！
其实我把我要处理的问题简化了，实际上是一个二维问题，在y方向上还有8个网格，
因此实际上是a(i,j)=b(i-1,j)+b(i,j)+b(i+1,j)的问题
因此一个block里面有32*8=256个threads，并且有多个blocks
我想将b()放在shared memory里，并在一个kernel内算完，因为重启kernel的话要访问global memory
我打算采用版主建议的多启动一些线程的办法
优点是代码简洁，缺点是浪费了一些计算资源，这或许也是不可避免的
LZ您好：
1：每个block256线程，这个从线程数上说一般是合适的，以及您有多个blocks，预计GPU的占用率还不错。
2：是的，重启kernel的话免不了要把前面的数据读进来，如果这个确实影响您的效率，那么可以考虑一个kernel搞定。
3：您还需要考虑shared memory容量是否够用，以及是否会因为shared memory使用较多导致SM上的resident block数量下降。以及，fermi或者kepler的话，您可以把shared memory调整到每SM 48KB，如果您可用的shared memory数量更为充沛一些。
4：注意到您在更新a()的时候使用了之前计算出来的b()。以及您打算将b()放在shared memory中，考虑到shared memory不能跨block访问（即A block申请的shared memory不能被B block访问），以及各block执行的进度是不保证先后顺序的（即，是乱序的），以及，除了kernel结束以外，无法同步全部的block。
所以，您的每个block可能需要考虑把更新a()所需的b()的一些边界值自己计算完毕，每个block需要重复计算一些边界值，以免依赖于其他的blocks。
也就是a(i,j)=b(i-1,j)+b(i,j)+b(i+1,j)这个更新中，假定i是从1到30，那么需要把b()对应的i=0和i=31的两种情况也计算进来，这个是“边界值”，自己计算完毕以后，a()的更新就不受制于人了。
这样，你可以在每个block里面先生成包括边界值在内的b()的值，然后__syncthreads()一下，保证本block所有生成b()的工作都完成，之后再生成a()。
5：以及，如果有需要就多上一些线程好了，影响并不是很大，不要太纠结。
大致如此，供您参考。
祝您编码顺利~~
以及非常荣幸为LZ服务，祝LZ好运！
楼主通过阅读您的代码，看上去您的b(i)是中间结果，但是为了能在多个线程间交互，不得暂时写入，然后再次读取。
以及，因为您的N+1个线程会完成N份有效数据a(i)，那么作出如下建议：
（1）如果不修改算法，可以考虑较大的block(假设N个线程), 这样，效率N/(N+1)较大。较好。
例如您32个线程，可以完成31个数据(96%效率), 而256个线程，可以完成255个数据(99.6%效率)
（2）对于线程的形状安排，我在ICE的建议上稍微作出修改，不建议32N+1完成32N份数据这种。而建议32N完成32N-1份数据这种。因为虽然在计算b(i)和a(i)的过程中，他们要求的数据量差了1份，导致必然无法同时满载，但建议选择让前者。举个例子说：
宁愿32个线程计算b(i), 然后31个线程计算a(i), 而
不要33个线程计算b(i), 然后32个线程计算a(i),
理由是：
(a)b和a的计算量是不同的，前者可能是个复杂的函数过程，而后者只是简单的减法。所以尽量在前者满载，后者浪费1/32基本无所谓。
(b)即使计算量一样，在小block规模下(例如33和32个线程的block),前者33个线程的安排，50%的计算能力将在第1步被浪费(因为第二个warp只有1个线程)，第2步满载。而后者在第一步和第二步基本都满载。
（3）更进一步的，建议：
每个线程计算b(2i)和b(2i + 1), 即每个线程计算临近的2个b(i),和2次a(i)的那个减法，而将总的线程数减少一半。这样：
（3-1）对于1.x和2.x, 可以有一半的操作b(奇数下标) - b(偶数下标)可以直接寄存器中执行减法。
同时另外一半（偶数下标的b()-奇数下标的)可以继续走shared memory。
这样前1次减法直接寄存器中减法了，基本无代价。而后1次和原来相同。
（3-2）对于3.x, 可以前一半直接寄存器中执行直接减法。同时后一半(偶数下标减奇数下标的b()), 可以线程间来个shuffle down(__shfl_down()), 这样一步即可交换完数据，这样这后半部分原本需要2步shared memory操作，只需要一步。
这样(3-2)总体比你单独的实现，只有大约25%的总体shared memory的代价(shuffle算1/2的shared memory操作).
以及，(3)中的建议还提高了(1)和(2)原始情况中的等价效率，减少了浪费。例如256个线程的block, 利用3种的计算安排，可以大致相当于原来的512个线程的数据浪费（都浪费1个）。
这是对ICE的2#的补充。
以及，您可以尝试将我的上文扩展到您的3#中的具体用例。
以及，如果您的计算是卡在了计算b上，而不是卡在数据的交换上，那么也可以只考虑ICE和我的(1)(2)说法。因为(3)对于不是瓶颈的地方的优化基本无太大意义。