optimization - 加快许多箱/几个箱的 CUDA 原子计算

Question

我正在尝试优化我在 CUDA 中的直方图计算。它为我提供了相对于相应 OpenMP CPU 计算的出色加速。但是，我怀疑(根据直觉)大多数像素都落入了几个桶中。为了论证的缘故，假设我们有 256 个像素落入让我们说，两个桶。

最简单的方法是这样做似乎是

将变量加载到共享内存中

如果需要，对 unsigned char 等进行矢量化加载。

在共享内存中进行原子添加

对全局进行合并写入。

像这样的东西:

__global__ void shmem_atomics_reducer(int *data, int *count){
  uint tid = blockIdx.x*blockDim.x + threadIdx.x;

  __shared__ int block_reduced[NUM_THREADS_PER_BLOCK];
  block_reduced[threadIdx.x] = 0;

  __syncthreads();

    atomicAdd(&block_reduced[data[tid]],1);
  __syncthreads();

  for(int i=threadIdx.x; i<NUM_BINS; i+=NUM_BINS)
    atomicAdd(&count[i],block_reduced[i]);

}

当我们减少 bin 的数量时，这个内核的性能(自然地)下降，从 32 个 bin 的大约 45 GB/s 降低到 1 个 bin 的 10 GB/s 左右。争用和共享内存库冲突作为原因给出。我不知道是否有任何方法可以以任何重要的方式删除这些计算中的任何一个。

我还一直在尝试来自 parallelforall 博客的另一个(美丽的)想法，其中涉及使用 __ballot 来获取扭曲结果，然后使用 __popc() 来减少扭曲级别。

__global__ void ballot_popc_reducer(int *data, int *count ){
  uint tid = blockIdx.x*blockDim.x + threadIdx.x;
  uint warp_id = threadIdx.x >> 5;

  //need lane_ids since we are going warp level
  uint lane_id = threadIdx.x%32;

  //for ballot
  uint warp_set_bits=0;

   //to store warp level sum
  __shared__ uint warp_reduced_count[NUM_WARPS_PER_BLOCK];
   //shared data
  __shared__ uint s_data[NUM_THREADS_PER_BLOCK];

 //load shared data - could store to registers
  s_data[threadIdx.x] = data[tid];

  __syncthreads();


//suspicious loop - I think we need more parallelism
  for(int i=0; i<NUM_BINS; i++){
      warp_set_bits = __ballot(s_data[threadIdx.x]==i);

      if(lane_id==0){
        warp_reduced_count[warp_id] = __popc(warp_set_bits);
      }

     __syncthreads();

      //do warp level reduce 
      //could use shfl, but it does not change the overall picture
      if(warp_id==0){
        int t = threadIdx.x;
        for(int j = NUM_WARPS_PER_BLOCK/2; j>0; j>>=1){
          if(t<j) warp_reduced_count[t] += warp_reduced_count[t+j];
          __syncthreads();
        }
      }                                                                                                                                                                                                                                                                



      __syncthreads();


      if(threadIdx.x==0){
        atomicAdd(&count[i],warp_reduced_count[0]);
        }  

    }                                                                                                                                                                                                                                             

  }

这为单个 bin 案例(1 个 bin 为 35-40 GB/s，而 10-15 GB/s 使用原子)，但是当我们增加 bin 数量时性能会急剧下降。当我们使用 32 个 bin 运行时，性能下降到大约 5 GB/s。原因可能是因为单线程循环遍历所有 bin，要求 NUM_BINS 循环的并行化。

我尝试了几种并行化 NUM_BINS 循环的方法，但似乎都没有正常工作。例如，可以(非常不雅地)操纵内核为每个 bin 创建一些 block 。这似乎表现相同，可能是因为我们将再次遭受多个 block 尝试从全局内存中读取的争用。另外，程序很笨拙。同样，在 y 方向上对 bin 进行并行化也会产生同样令人沮丧的结果。

我尝试的另一个想法是动态并行，为每个 bin 启动一个内核。这是灾难性的缓慢，可能是由于子内核没有真正的计算工作和启动开销。

最有希望的方法似乎是 - 来自 Nicholas Wilt 的 article

使用这些所谓的私有(private)化直方图，其中包含共享内存中每个线程的 bin，这在表面上对 shmem 的使用非常重要(我们在 Maxwell 上每个 SM 只有 48 kB)。

也许有人可以对这个问题有所了解？我觉得应该去改变算法而不是使用直方图，使用不那么频繁的东西。否则，我想我们只使用原子版本。

编辑:我的问题的上下文是计算用于模式分类的概率密度函数。我们可以通过使用非参数方法(例如 Parzen Windows 或 Kernel Density Estimation)来计算近似直方图(更准确地说是 pdf)。然而，这并没有克服维度问题，因为我们需要对每个 bin 的所有数据点求和，当 bin 的数量变大时，这会变得很昂贵。见这里:Parzen

Answer 1

我在使用聚类时遇到了类似的挑战，但最终，最好的解决方案是使用扫描模式对处理进行分组。所以，我不认为它对你有用。既然你要求这方面的一些经验，我会和你分享我的。

问题

在您的第一个代码中，我猜想通过减少箱数来处理低性能与扭曲停顿有关，因为您对每个评估数据执行的处理很少。当 bin 数量增加时，该内核的处理和全局内存负载(数据信息)之间的关系也会增加。您可以通过 Nsight 的性能分析中的“问题效率”实验非常轻松地检查这一点。可能你得到的循环率很低，至少有一个清晰的经线(经线问题效率)。


由于我无法将可识​​别扭曲的数量提高到接近 95%，因此我放弃了这种方法，因为在某些情况下它会变得更糟(内存依赖性使我 90% 的处理周期停滞不前。


如果垃圾箱的数量不是很大，则洗牌和减少投票非常有用。如果它太大，则应该为每个 bin 过滤器激活少量线程。所以你最终可能会出现很多代码分歧，这对于并行处理来说是非常不可取的。您可以尝试对分歧进行分组以消除分支并获得良好的控制流，因此整个扭曲/ block 呈现类似的处理，但跨 block 的机会很多。



一个可行的解决方案

我不知道在哪里，但是我看到了针对您的问题的非常好的解决方案。你试过this one ?

您也可以使用 vectorized load并尝试类似的方法，但我不确定它会在多大程度上提高您的性能:

__global__ hist(int4 *data, int *count, int N, int rem, unsigned int init) {

__shared__ unsigned int sBins[N_OF_BINS]; // you may want to declare this one dinamically
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (threadIdx.x < N_OF_BINS) sBins[threadIdx.x] = 0; 

for (int i = 0; i < N; i+= warpSize) {
    atomicAdd(&sBins[data[i + init].w], 1);
    atomicAdd(&sBins[data[i + init].x], 1);
    atomicAdd(&sBins[data[i + init].y], 1);
    atomicAdd(&sBins[data[i + init].z], 1);
}

//process remaining elements if the data is not multiple of 4
// using recast and a additional control
for (int i = 0; i < rem; i++) {
    atomicAdd(&sBins[reinterpret_cast<int*>(data)[N * 4 + init + i]], 1);
} 
//update your histogram data here
}