1

在内核计算之后，我产生了不同的值，范围从 0 到 6399，它们存储在共享内存中。我有 24336 个块，因此有 24336 个__shared__大小为 256 的数组实例。每个块数组都以没有特定顺序的计算值填充。

我想要的是计算所有这些块共享内存中有多少次某个值，并且该值应该是另一个数组（驻留在全局内存中）的索引，其各自的值将是它的次数出现了。

在一个带有 2 个块和一个__shared__ int array1[3]

__device__ array2我可能有 ：

为了blockIdx.x=0

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

如何尽可能地并行完成？

编辑：

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

在我的例子中，我使用了一个包含 68000000 个整数的数组，其 bin 范围从 0 到 6399。它工作得很好，而且我得到了加速，所以我忘记了将所有 bin 存储在共享内存中并从那里计算 bin 数量的最初想法，但我对执行时间不太满意，我想尝试其他方法。

我想知道是否有人知道如何回到我最初的想法以及我应该使用什么技术（即独家扫描等）。我记得有一个 stackoverflower 同事对此发表了一个答案，但我认为他很快就删除了他的帖子，而我没有时间彻底查看它。

4

2 回答 2

2

我理解这个问题与直方图相同。存在许多好的解决方案，可以通过谷歌搜索，包括Nvidia 自己的示例。

特别值得一提的是这些代码：

于 2013-02-24T20:19:08.627 回答
1

据我所知，这只是建立一个直方图。虽然这种方法可能不是最快的（除非您使用的是 Kepler K20），但您可以在内核末尾做一些相对简单的事情（假设您的共享数组 1 的大小为 256 个元素并且您在一维中启动至少 256 个线程线程块）：

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

（假设原子函数的计算能力为 1.1 或更高）

于 2013-02-24T20:21:59.290 回答