之前发过一个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版主! 新年快乐!