之前发过一个kernel各种错误的帖子，终于发现了几个错误。由于代码太长太乱，我也就不整个发上来了。
错误1：对SMEM的读写搞错了偏移量。
错误2：错误地以为使用threadfence就可以确保block间的同步。事实上，threadfence只是确保当前block中的thread对内存的写入对其他block可见而已，并不含有任何同步的含义。这也是threadfence系列与synthreads的根本不同。
错误2是我在看5.0的Programming Guide中讲解threadfence的例子时发现的。基于这个例子（代码如下），我还想请教大家一个问题。
先放代码：
__device__ unsigned int count = 0;
__shared__ bool isLastBlockDone; //改为__device__
__global__ void sum(const float* array, unsigned int N, float* result)
{
// Each block sums a subset of the input array
float partialSum = calculatePartialSum(array, N);
if (threadIdx.x == 0) {
// Thread 0 of each block stores the partial sum
// to global memory
result[blockIdx.x] = partialSum;
// Thread 0 makes sure its result is visible to
// all other threads
__threadfence();
// Thread 0 of each block signals that it is done
unsigned int value = atomicInc(&count, gridDim.x);
// Thread 0 of each block determines if its block is
// the last block to be done
isLastBlockDone = (value == (gridDim.x - 1));
//这里加入__threadfence();
}
// Synchronize to make sure that each thread reads
// the correct value of isLastBlockDone
__syncthreads();
//加入 do { rubbish } while (!isLastBlockDone)
if (isLastBlockDone) {
// The last block sums the partial sums
// stored in result[0 .. gridDim.x-1]
float totalSum = calculateTotalSum(result);
if (threadIdx.x == 0) {
// Thread 0 of last block stores total sum
// to global memory and resets count so that
// next kernel call works properly
result[0] = totalSum;
count = 0; //这句话多余了吧？
}
}
}
这段代码中使用了threadfence和atomicinc来确保所有的block都贡献了自己计算的部分和，并在最后写完的block中进行总和的计算。
我的问题是：如果每个block的部分和不仅仅用来计算总和，而是找到部分和最大的block，并让所有的block都知道，那么这该如何操作呢？ 难不成要给先完成的block弄个do–while循环做点无用功等着？初步想法也已经标在上面的代码中了。
谢谢大家！预祝大家新年快乐！
楼主不错。
很认真的发表了自己的心得，甚至说，在使用CUDA过程中各样的“心路历程”。挺好！感谢楼主！
关于您的“如果每个block的部分和不仅仅用来计算总和，而是找到部分和最大的block，并让所有的block都知道，那么这该如何操作呢？”这个问题，我表示无法就地实现。
但是可以通过在下一个kernel里统计出，那个block的结果是最大值。
没有办法在运行的时候，就地得出，并让其他block知道的。
这个很容易理解：不是所有的blocks都同时在运行的，假设你有10000个blocks，但同时可能你的卡上，只有64个blocks在同时存在。那么此时，你如何能得知没有运行的剩下的9000多个blocks里的结果值不是最大的呢？（果断你不能)。以及，如果前面1000个blocks都运行完毕了，你在后续的blocks里发现了最大值，又如何能通知前面已经死去的blocks呢？（依然果断你不能）
所以你只能在下一次统计了。
BZ您好，对于这个疑惑我今天试了一下。用一个在GMEM上的volatile全局变量value做为开关信号，每个block先向一个GMEM[NumofBlock]写入，然后使用原子操作对这个value+1，最后做while空循环直到value==NumofBlocks-1（这样做是为了让先做完的block等着最后做完的，所以value定义成了volatile型）。这时所有的block都读取GMEM[NumofBlock]，然后在block内部各自进行比较。目前来看这样做所得的结果是正确的，但考虑到GPU的复杂性又让我不知道这样是否特别可靠。BZ大人能帮我分析一下么？十分感谢啊！！！
另外，另开kernel的方法对我这次的应用来说不太适用，我是想着尽量减少device与host之间的传递，重开kernel的方法太奢侈了啊
你这样做不行。
但是为何你今日测试“成功”了呢？因为你的blocks数太少。
当你的blocks稍微一多，例如你的卡可以上100个blocks, 但是你的kernel需要200个，
那么显然先上去的100个，在等待你的atomicAdd对你的value到200的while过程中。会一直死循环。
(楼主想想为何！）
因为这个想法不可以（更详细的说法见2#, 已经说过了), 所以你不能考虑所有的blocks都同步。
而且，在我2#给出了详细的解释为何试图这样来全局同步不行的情况下，你还继续这么问！太让我伤心了。请立刻返回看2#!
还有其他小问题，但在大前提都不成立的情况下，不为你指出更多了。
祝好！
重开kernel, 不一定比你神马空循环要代价高（先不说后者是否靠谱）。
您的拒绝接受建议的权利是您自有的，我只是一个服务员。
我只能尽我最大的努力来服务您。
您对我重开kernel让blocks同步的说法拒绝后，您还可以寻求其他版主/会员/NVIDIA工程师的其他方法来同步您的多个blocks。
但是我只能尽力而为。请立刻咨询其他成员是否还有同步所有的blocks的其他方法。
LZ您好，横扫千军斑竹已经详细说明了各种情况，我赞同他的说法，并且愿意以一句话简单地总结下：
要实现所有block的同步，只能采用结束kernel的方法。
（其原因在于，只有死亡是人人平等的。）
大致如上，供您参考，祝您春节愉快~
非常抱歉让您误会了，我没有任何拒绝接受建议的意思。其实恰恰相反，我在这个论坛得到了你们几位BZ非常大的帮助，这也正是我发帖求助的目的啊。我当时回复的时候是赶着要去上课，所以2#的内容没有看仔细，这个确实是要怪我自己。
让我来详细解释一下吧，我这次的程序是要对很大的一块数据进行4维空间的查表匹配。我现在的做法是把待匹配的数据都读到SMEM中，各个block分管不同的参数空间来进行查找。所以我一直都觉得能在一个kernel里搞定就尽量不要多开或者重开（因为每次kernel发射时都要重新读取GMEM并写入SMEM）。
基于此，我严格控制了kernel中所用到的block数量，使之不大于MP的数目（这样应该就不会导致您在2#中所说的情况）。每个block中也分配不多于512个threads，以充分利用所有的register。
我的问题是：采用一个全局volatile变量做为指示信号来同步所有的blocks，对于我现在的做法可行么？
再次感谢！
BZ您好，我之前没来得及详细介绍我的思路，还请原谅。
如果确定blocks数目小于所有MP上能驻留的总数。
那么这样做可以。
只是可以，无论效率。
此外，你还要小心。别估计的乐观了，导致死循环挂了。有的时候，理论能驻留的blocks总数不是你认为的那么些。所以要小心，最好的办法就是不要。
也就是说这样做仍然不见得比重开kernel有优势么？
这个让我确实很费解了
别混淆了。认真看吧。我虽然是拿钱干活的，但也不是能被你如此的无视。
你引用的楼层你读过么？？你懂中文么？？你看了么？
引用：
“此外，你还要小心。别估计的乐观了，导致死循环挂了。有的时候，理论能驻留的blocks总数不是你认为的那么些。所以要小心，最好的办法就是不要。”
请问你反问我的和重开kernel的效率问题，和你引用批判我的文字有关系么？？
有么！！！！！！
考虑到你连续在1个主题中对我的2个认真回复给予轻视。
我拒绝在14d内回答你的任何问题以及反问。
建议延请高明。
冤枉啊。。。。我说我当时没看到第二页的内容您信么？
不管怎样，十分感谢您的帮助！ 新年快乐！
LZ您好，请淡定。
其实仔细地看看横扫版主的 各楼的回复，然后自己试验一下，您一定能有更深刻的认识。
（至于他说14days 之内不理你，据我对他的了解，我可以悄声告诉你，他肯定会提前忘记的，只要你仔细看他写的每句话。嘘，别说是我说的啊，闪了。）
:lol嘿嘿。
我就是打算试试重开kernel的做法啊。多谢ice版主！ 新年快乐！