一个样本求平均值问题

加速计算 CUDA
1

现在我有一个数据立方体，大概是806050的数据量。即组成一个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;

}

2

你的W ，BLOCK_SIZE 多少都不知道

3

(Sum+j)=(Sum+j)+sum[j];
这句中所有线程都往Sum+j中写数据。会有写冲突。

4

LZ您好，大致分析了您的代码和叙述，意见如下：

您给出的信息不足，大致有：
1：您没有叙述您要求的是什么和什么的平均值，是每点50个数值的平均值？每一层8060个数的平均值？还是其他什么。这将极大地影响观看您帖子的人获知您意图的速度。在此，我猜测您是需要计算每一层8060个数据的平均值。

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上快。

大致分析如上，祝您编码愉快~

5

谢谢您，您的回答很详细，我要求的确实是80×60这么多样本的均值，其中每个样本是个50维的数据。既然我的程序出现这么多问题，请问我该怎么去实现这个功能呢？

6

最后那个求Sum的除只能拿到global函数外面啦，这就要从显存复制到主存来除了，接下来的其他的global函数还要用到这个结果，又要拷回显存，太浪费时间了。这是个问题呀

7

1） 如果楼主坚持自己写内核，请遵从4#的建议，并可以尝试在必要时候使用原子操作。
2） 如果你仅仅是想实现这个问题，又觉得自己写内核很痛苦，个人建议你去看看cublas库中的cublasasum函数。可以直接对80*60个数据求和并直接写回CPU内存，然后再在cpu上算一下除法。这样可以省很多精力。
3） 对于这种求均值问题的计算量+数据传输，放到GPU上不一定划算。

8

LZ您好，实现您的功能可以有多种方法：

1：您可以继续使用您之前的算法思想，修正2#，3#，4#指出的在实现上的问题即可，依然可以得到正确的结果（不考虑计算效率）。

2：您也可以使用规约的方法，并在规约到最后一步的时候，加上除法，直接获得最后的正确结果。

3：您也可以考虑7#指出的，使用现有的线性代数库的函数，求和，并在后处理的时候，作个除法。

总之，实现您意图的方法还是很多的。

祝您编码顺利~

9

LZ您好，这种情况，您也可以把数据留在global memory里面不回拷，再次启动一个kernel专门执行除法。

或者您也可以修改前面的求和部分，将除法一起完成即可。

祝您好运~