如何在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，如果我可以计算哪些直方图bin是非零的(这个操作是并行的，因为每个线程被分配到每个bin，我可以自动添加每个线程生成的值)，让我们说我可以启动一个新的共享索引变量sindex=0，每次线程想要将值存储到d_hist[]时，它可以从sindex中获取最新的值并将其值存储到d_hist[sindex]=hist[treadIdx.x]，之后我可以atomicadd sindex

现在只有一个问题，将有一个竞争条件，以获得sindex的值，所以我可能不得不设置一个标志，可以锁定或解锁时，一个线程添加任何值d_hist(但我认为这里可能有死锁的情况)

这个技术会工作吗?还有比这更好的技术吗?

Answer 1

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

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

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加，除此之外，移动和检查部分只是常数时间操作，如果有人有不同的想法，请分享。