为图像直方图优化金属计算着色器

Question

我有一个金属着色器，可以像这样计算图像直方图：

#define CHANNEL_SIZE (256)
typedef atomic_uint HistoBuffer[CHANNEL_SIZE];

kernel void
computeHisto(texture2d<half, access::read> sourceTexture [[ texture(0) ]],
             device HistoBuffer &histo [[buffer(0)]],
             uint2  grid [[thread_position_in_grid]]) {
  if (grid.x >= sourceTexture.get_width() || grid.y >= sourceTexture.get_height()) { return; }

  half gray  = sourceTexture.read(grid).r;
  uint grayvalue = uint(gray * (CHANNEL_SIZE - 1));
  atomic_fetch_add_explicit(&histo[grayvalue], 1, memory_order_relaxed);
}

这可以按预期工作，但所需时间过长（> 1ms）。我现在尝试通过减少原子操作的数量来优化此操作。我想出了以下改进的代码。这个想法是计算每个线程组的局部直方图，然后将它们原子地添加到全局历史缓冲区中。

kernel void
computeHisto_fast(texture2d<half, access::read> sourceTexture [[ texture(0) ]],
             device HistoBuffer &histo [[buffer(0)]],
             uint2 t_pos_grid [[thread_position_in_grid]],
             uint2 tg_pos_grid [[ threadgroup_position_in_grid ]],
             uint2 t_pos_tg [[ thread_position_in_threadgroup]],
             uint  t_idx_tg [[ thread_index_in_threadgroup ]],
             uint2 t_per_tg [[ threads_per_threadgroup ]]
             )
{
  threadgroup uint localhisto[CHANNEL_SIZE] = { 0 };

  if (t_pos_grid.x >= sourceTexture.get_width() || t_pos_grid.y >= sourceTexture.get_height()) { return; }

  half gray  = sourceTexture.read(t_pos_grid).r;
  uint grayvalue = uint(gray * (CHANNEL_SIZE - 1));
  localhisto[grayvalue]++;

  // wait for all threads in threadgroup to finish
  threadgroup_barrier(mem_flags::mem_none);

  // copy the thread group result atomically into global histo buffer
  if(t_idx_tg == 0) {
    for(uint i=0;i<CHANNEL_SIZE;i++) {
      atomic_fetch_add_explicit(&histo[i], localhisto[i], memory_order_relaxed);
    }
  }
}

有两个问题：

1. 改进后的例程与第一个例程相比不会产生相同的结果，我目前不明白为什么？
2. 运行时间没有改善。实际上，它的运行时间是未优化版本的4倍。根据调试器的for循环是问题。但是我不明白这一点，因为原子操作的数量减少了3个数量级，即线程组的大小，这里（32x32）= 1024。

谁能解释我在这里做错了什么？谢谢

编辑：2019-12-22：根据Matthijs的回答，我也将本地直方图更改为原子操作，如下所示：

threadgroup atomic_uint localhisto[CHANNEL_SIZE] = {0};

half gray  = sourceTexture.read(t_pos_grid).r;
uint grayvalue = uint(gray * (CHANNEL_SIZE - 1));
atomic_fetch_add_explicit(&localhisto[grayvalue], 1, memory_order_relaxed);

但是结果与上面的参考实现中的结果不同。必须存在另一个严重的概念错误???

Answer 1

您仍然需要在线程组内存上使用原子操作，因为它仍被多个线程共享。这应该比您的第一个版本更快，因为相同锁的争用较少。