现在我有一个数据立方体,大概是8060 50的数据量。即组成一个80*60的面阵,每个点有50个数据组成。
下面是我写的求这4800条数据的样本平均的程序,可是发现结果不正确,希望有大神可以指出问题在哪?
global void kernel_average(float *Lbg,float *Sum)//Lbg是存储数据的数组,按行顺序存储,已读入显存。Sum是一个50的数组,保存样本均值。
{
int bx = blockIdx.x;
int by = blockIdx.y;
int tx = threadIdx.x;
int ty = threadIdx.y;
int c = W * BLOCK_SIZE * by + BLOCK_SIZE * bx;
int i,j;
float sum[50];
for( i=0;i<50;i++)
{
sum[i]=0.0f;
sum[i]=Lbg[(c+W*ty+tx)*LEN+i];
}
for( j=0;j<50;j++)
{
*(Sum+j)=*(Sum+j)+sum[j];
}
__syncthreads();
for(j=0;j<50;j++)
*(Sum+j)/=4800;
}
(Sum+j)= (Sum+j)+sum[j];
这句中所有线程都往Sum+j中写数据。会有写冲突。
LZ您好,大致分析了您的代码和叙述,意见如下:
您给出的信息不足,大致有:
1:您没有叙述您要求的是什么和什么的平均值,是每点50个数值的平均值?每一层8060个数的平均值?还是其他什么。这将极大地影响观看您帖子的人获知您意图的速度。在此,我猜测您是需要计算每一层80 60个数据的平均值。
2:如同2#所说,您没有给出一些宏/变量的值,也没有给出您调用kernel的参数,所以我假定您上述参数都是使用正确的,合理的,以线程参数等完成寻址是正确的。如果您实际使用有问题在这些地方,请您根据实际情况予以排查。
下面根据您的代码和上述猜测,说一下观察出来的问题:
1:您在kernel里面给每个线程分配了一个数组sum[50]并使用第一个循环对其赋值,之后,使用“(Sum+j)= (Sum+j)+sum[j]”打算完成所有线程j位置的局部值sum[j]对全局的Sum的累加。正如3#指出的,这个累加是错误的。这里必须使用原子操作才能保证正确性。
2:第二个循环之后您使用了__syncthreads(),这个线程同步只能在一个block之内起作用,block之间的顺序和完成时间是不被保证的。也就是说所,即便前面使用了原子操作,您这里加上__syncthreads()是无法保证所有的线程都累加完毕的。
3:您在kernel的最后使用一个循环,对Sum的值除以4800。这里必须指出,即便上一步真的能完成所有线程的同步(实际上只能完成一个block的同步),您这样写,每个线程都会去把Sum除一遍,而且这里又是非原子操作,所以最终Sum被除了多少遍将成为一个无法得知的值,但可以肯定,这个行为和您所需要的行为不相符合,您只需要除一次。
综上,您的代码存在大量问题,希望您能理解多线程行为的规律和一些基本原则,修改您的代码,而不是照搬已有的串行实现。
此外,上述分析只针对您当前实现思路,给出保证逻辑正确的建议,无关性能分析,如果您需要性能更高一些,建议使用规约算法;以及,您的数据量其实是比较小的,或许您折腾完GPU之后会发现,还没有直接在CPU上快。
大致分析如上,祝您编码愉快~
LZ您好,大致分析了您的代码和叙述,意见如下:
您给出的信息不足,大致有:
谢谢您,您的回答很详细,我要求的确实是80×60这么多样本的均值,其中每个样本是个50维的数据。既然我的程序出现这么多问题,请问我该怎么去实现这个功能呢?
最后那个求Sum的除只能拿到global函数外面啦,这就要从显存复制到主存来除了,接下来的其他的global函数还要用到这个结果,又要拷回显存,太浪费时间了。这是个问题呀
1) 如果楼主坚持自己写内核,请遵从4#的建议,并可以尝试在必要时候使用原子操作。
2) 如果你仅仅是想实现这个问题,又觉得自己写内核很痛苦,个人建议你去看看cublas库中的cublasasum函数。可以直接对80*60个数据求和并直接写回CPU内存,然后再在cpu上算一下除法。这样可以省很多精力。
3) 对于这种求均值问题的计算量+数据传输,放到GPU上不一定划算。
LZ您好,实现您的功能可以有多种方法:
1:您可以继续使用您之前的算法思想,修正2#,3#,4#指出的在实现上的问题即可,依然可以得到正确的结果(不考虑计算效率)。
2:您也可以使用规约的方法,并在规约到最后一步的时候,加上除法,直接获得最后的正确结果。
3:您也可以考虑7#指出的,使用现有的线性代数库的函数,求和,并在后处理的时候,作个除法。
总之,实现您意图的方法还是很多的。
祝您编码顺利~
LZ您好,这种情况,您也可以把数据留在global memory里面不回拷,再次启动一个kernel专门执行除法。
或者您也可以修改前面的求和部分,将除法一起完成即可。
祝您好运~