0

我正在使用原始 CUDA 内核实现一个算法，其中每个线程块都需要该线程块的可用数据的密集直方图，现在的问题是我是否必须从头开始计算密集直方图？（如果我已经拥有使用共享内存实现的稀疏直方图，是否值得计算密集直方图）

我想出了这个转换的想法，我会尝试用例子来详细说明我的想法（temp和hist都在共享内存中）

   0,1,2,3,4,5,6... //array indexes
   4,3,0,2,1,0,5... //contents of hist[]
   0,0,2,0,0,5,0... //contents of temp[] if(hist[x]>0)temp[x]=x;
   for_every_element     //this is sequential part :(
       if(temp[x]>0)
         shift elements from index x to 256
   4,3,2,1,0,5... //pass 1 of the for loop
   4,3,2,1,5...   //pass 2 of the for loop
                  //this goes on until all the 0s are compacted

现在我知道上面本质上是顺序的，但是可以使用恒定时间（并且并行）进行移位，因为threads_per_block已经设置为256，所以移位不是主要问题，主要问题是如何改进它（或任何欢迎其他建议）。

编辑：我正在考虑另一个想法，如下假设threads_per_block=256我是否可以计算哪些直方图箱是非零的（此操作是并行的，因为每个线程都分配给每个箱，我可以原子添加每个线程生成的值）假设我可以启动一个新的共享索引变量sindex=0，每次线程想要将值存储到d_hist[]其中时，可以从 sindex 获取最新值并将其值存储到d_hist[sindex]=hist[treadIdx.x]之后我可以 atomicAdd sindex

现在只有一个问题，获取 sindex 的值将存在竞争条件，因此我可能必须设置一个标志，当线程向其中添加任何值时可以锁定或解锁d_hist（但我认为可以在这里陷入僵局）

这种技术会奏效吗？还有比这更好的技术吗？

4

2 回答 2

1

将稀疏直方图转换为密集直方图是一种分散操作。如果稀疏直方图由s_index[S_N]和组成s_hist[S_N]，那么首先我们创建一个d_hist[N]由全零组成的密集直方图（也许你可以从主机代码中做到这一点）。然后我们填充密集直方图，d_hist[s_index[i]] = s_hist[i]; 这可以并行完成，并使用与稀疏直方图中有效索引一样多的线程（i < S_N）。假设您的直方图已排序，您将根据稀疏直方图索引的分布获得任何可能的合并好处。

对于每个线程块都在做单独的直方图的情况，这可能没有意义，但您可能也对推力散射感兴趣。

于 2013-04-12T15:18:44.920 回答
0

好吧，我想最简单的方法是找出哪个bins>0和之后，然后可以进行排他扫描（例如为了计算目标索引sum_array[]），然后将所有bins>0移动到d_hist[sum_array[threadIdx.x]-1]=s_hist[threadIdx.x]

 0,1,2,3,4,5,6... //s_indexes[]

 4,3,0,2,1,0,5... //contents of s_hist[]
 1,1,0,1,1,0,1... //all bins which are > 0 = sum_array[]

 1,2,2,3,4,4,5... //inclusive_scan of summ_array[]

 //after the moving part
 0,1,3,4,6... //s_indexes[]
 4,3,2,1,5... //d_hist[]
 0,1,2,3,4... //d_indexes[]

我倾向于使用这种模式的原因是因为它需要 log_base_2(256) 时间来计算 sum_array plus，除此之外，移动和检查部分只是恒定时间操作，如果有人有不同的想法，请分享。

于 2013-04-13T04:24:01.323 回答