具有单分支的 CUDA 内核比没有分支的内核运行速度快 1.5 倍

Question

我在过滤器内核上有一个奇怪的性能倒置，不管有没有分支。带分支的内核比不带分支的内核运行速度快约 1.5 倍。

基本上我需要对一堆辐射射线进行排序，然后应用交互内核。由于附带的数据很多，所以不能多次使用thrust::sort_by_key()之类的东西。

算法思路：

1. 为所有可能的交互类型（共五种）运行循环
2. 在每个循环中，warp 线程都会为其交互类型投票
3. 循环完成后，每个 warp 线程都知道另一个具有相同交互类型的线程
4. 线程选举他们的领导者（每个交互类型）
5. Leader 使用 atomicAdd 更新交互偏移表
6. 每个线程将其数据写入相应的偏移量

我使用了这篇 Nvidia 帖子https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/中描述的技术

我的第一个内核在循环中包含一个分支，运行时间约为 5 毫秒：

int active;
int leader;
int warp_progress;
for (int i = 0; i != hit_interaction_count; ++i)
{
  if (i == decision)
  {
    active = __ballot(1);
    leader = __ffs(active) - 1;
    warp_progress = __popc(active);
  }
}

我的第二个内核使用两个元素的查找表，不使用分支并运行约 8 毫秒：

int active = 0;
for (int i = 0; i != hit_interaction_count; ++i)
{
  const int masks[2] = { 0, ~0 };
  int mask = masks[i == decision];
  active |= (mask & __ballot(mask));
}
int leader = __ffs(active) - 1;
int warp_progress = __popc(active);

公共部分：

int warp_offset;
if (lane_id() == leader)
  warp_offset = atomicAdd(&interactions_offsets[decision], warp_progress);
warp_offset = warp_broadcast(warp_offset, leader);
...copy data here...

怎么可能？有没有办法实现这样的过滤器内核，使其运行速度比分支更快？

UPD：完整的源代码可以在 filter_kernel cuda_equation/radiance_cuda.cu https://bitbucket.org/radiosity/engine/src

找到

Answer 1

我认为这是CPU程序员大脑变形。在 CPU 上，由于消除了分支和分支预测错误，我预计性能会有所提升。

但在 GPU 上没有分支预测，也没有惩罚，所以只有指令计数很重要。

首先我需要将代码重写为简单的代码。

有分支：

int active;
for (int i = 0; i != hit_interaction_count; ++i)
    if (i == decision)
        active = __ballot(1);

没有分支：

int active = 0;
for (int i = 0; i != hit_interaction_count; ++i)
{
  int mask = 0 - (i == decision);
  active |= (mask & __ballot(mask));
}

在第一个版本中，有 ~3 个操作：compare、if 和 __ballot()。 在第二个版本中，有大约 5 个操作：compare、make mask、__ballot()、& 和 |=。 公共代码中大约有 15 个操作。

两个循环都运行 5 个周期。首先是 35 个操作，其次是 45 个操作。这种计算可以解释性能下降。