如何在CUDA中将稀疏直方图转换为密集直方图？

Question

我正在使用原始 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] 之后我可以原子添加sindex

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

这种技术行得通吗？还有比这更好的技术吗？

Answer 1

将稀疏直方图转换为密集直方图是一种分散操作。如果稀疏直方图由s_index[S_N] 和s_hist[S_N] 组成，那么首先我们创建一个由全零组成的密集直方图d_hist[N]（也许您可以从主机代码中执行此操作）。然后我们用d_hist[s_index[i]] = s_hist[i]; 填充密集直方图这可以并行完成，并使用与稀疏直方图中有效索引一样多的线程（i

对于每个线程块执行单独直方图的情况可能没有意义，但您可能也对thrust scatter 感兴趣。

Answer 2

好吧，我想最简单的方法是找出哪个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，除此之外，移动和检查部件只是恒定时间操作，如果有人有不同的想法这个，请分享。