2013-02-24 80 views
1

内核计算后,我生成了不同的值,范围从0到6399,存储在共享内存中。我有24336个块,因此大小为256的__shared__数组的24336个实例。每个块数组都以没有特定顺序的方式填充计算值。Cuda:使用共享内存计算元素

我想要的是统计在所有这些块共享内存中有多少次存在某个值,并且该值应该是另一个数组的索引(驻留在全局内存中),并且其相应的值将是数字它出现过的时间。

在具有2个块的变形较短例和__shared__ int array1[3]

__device__ array2我可能有:

对于​​

array1[0]=10; 
array1[1]=20; 
array1[2]=30; 

而在blockIdx.x=1

array1[0]=30; 
array1[1]=0; 
array1[2]=10; 

结果应该是

array2[0]=1; //value 0 has appeared one time 
array2[10]=2; //value 10 has appeared two times 
array2[20]=1; //value 20 has appeared one time 
array2[30]=2; //value 30 has appeared two times 

这怎么可能尽可能地并行完成?

编辑

从跟随我的问题,我发现了很多的帮助,我的问题的答案。 尤其是一个代码,它可以生成任何类型的直方图,并将任何数量的垃圾箱和包含垃圾箱的数组作为输入。 https://devtalk.nvidia.com/default/topic/511531/code-general-purpose-histogram/ 我忘了我最初的计划,并创建了一个__global__数组,并在那里存储了所有的bin。

在我的情况下,我使用了一个68000000整数的数组,范围从0到6399。它工作得很好,我得到了加速,所以我忘记了我最初的想法,将所有分箱存储在共享内存中,垃圾箱,但我对执行时间并不满意,我想尝试其他方法。

我想知道是否有人对如何回到我最初的想法以及我应该使用什么技术(即独家扫描等)有任何想法。 我记得一个同伴stackoverflower已经发布了一个答案,但他删除了他的帖子,我想很快,没有我有时间仔细查看。

回答

1

尽可能靠近我可以告诉这只是建立一个直方图。虽然这种方法可能不会是最快的(除非你可能在开普勒K20上),你可以在你的内核结束时做一些相对简单的事情(假设你的共享array1是256个元素,并且你在1D中启动至少256个线程threadblock):

if (threadIdx.x < 256) 
    atomicAdd(&(array2[array1[threadIdx.x]]), 1); 

(假定计算能力1.1或更好atomic function