计算能力 6 上 __match_any_sync 的替代方案是什么？答案

【问题标题】：What's the alternative for __match_any_sync on compute capability 6?计算能力 6 上 __match_any_sync 的替代方案是什么？
【发布时间】：2020-05-09 18:26:01
【问题描述】：

在 cuda 示例中，使用了e.g. here、~~__match_all_sync~~ __match_any_sync。

这是一个示例，其中一个 warp 被分成多个（一个或多个）组，每个组都跟踪自己的原子计数器。

// increment the value at ptr by 1 and return the old value
__device__ int atomicAggInc(int* ptr) {
    int pred;
    //const auto mask = __match_all_sync(__activemask(), ptr, &pred); //error, should be any_sync, not all_sync
    const auto mask = __match_any_sync(__activemask(), ptr, &pred);
    const auto leader = __ffs(mask) - 1;  // select a leader
    int res;
    const auto lane_id = ThreadId() % warpSize;
    if (lane_id == leader) {                 // leader does the update
        res = atomicAdd(ptr, __popc(mask));
    }
    res = __shfl_sync(mask, res, leader);    // get leader’s old value
    return res + __popc(mask & ((1 << lane_id) - 1)); //compute old value
}

这里的__match_any_sync 将warp 中的线程分成具有相同ptr 值的组，这样每个组都可以自动更新自己的ptr 而不会妨碍其他线程。

我知道 nvcc 编译器（从 cuda 9 开始）会自动在后台进行这种优化，但这只是 __match_any_sync 的机制@

有没有办法实现这种预计算能力 7？

【问题讨论】：

标签： cuda gpu-warp

【解决方案1】：

编辑： 博客文章现已修改为反映 __match_any_sync() 而不是 __match_all_sync()，因此应忽略下面对此效果的任何评论。编辑下面的答案以反映这一点。

根据您的陈述：

这只是关于 __match_any_sync 的机制

我们将专注于替换__match_any_sync 本身，而不是重写atomicAggInc 函数的任何其他形式。因此，我们必须提供一个掩码，该掩码与 __match_any_sync() 在 cc7.0 或更高架构上返回的值相同。

我相信这将需要一个循环，它广播 ptr 值，在最坏的情况下，对 warp 中的每个线程进行一次迭代（因为每个线程都可以有一个唯一的 ptr 值）并测试哪些线程具有相同的值。根据每个线程中的实际ptr 值，我们可以通过多种方式“优化”此函数的此循环，以便可能将行程计数从 32 减少到某个较小的值，但在我看来，这种优化引入了相当大的复杂性，这使得最坏情况的处理时间更长（这是早期退出优化的典型特征）。所以我将演示一个没有这种优化的相当简单的方法。

另一个考虑是在warp没有收敛的情况下怎么办？为此，我们可以使用__activemask() 来识别这种情况。

这是一个有效的例子：

$ cat t1646.cu
#include <iostream>
#include <stdio.h>

// increment the value at ptr by 1 and return the old value
__device__ int atomicAggInc(int* ptr) {
    int mask;
#if __CUDA_ARCH__ >= 700
    mask = __match_any_sync(__activemask(), (unsigned long long)ptr);
#else
    unsigned tmask = __activemask();
    for (int i = 0; i < warpSize; i++){
#ifdef USE_OPT
      if ((1U<<i) & tmask){
#endif
        unsigned long long tptr = __shfl_sync(tmask, (unsigned long long)ptr, i);
        unsigned my_mask = __ballot_sync(tmask, (tptr == (unsigned long long)ptr));
        if (i == (threadIdx.x & (warpSize-1))) mask = my_mask;}
#ifdef USE_OPT
      }
#endif
#endif
    int leader = __ffs(mask) - 1;  // select a leader
    int res;
    unsigned lane_id = threadIdx.x % warpSize;
    if (lane_id == leader) {                 // leader does the update
        res = atomicAdd(ptr, __popc(mask));
    }
    res = __shfl_sync(mask, res, leader);    // get leader’s old value
    return res + __popc(mask & ((1 << lane_id) - 1)); //compute old value
}



__global__ void k(int *d){

  int *ptr = d + threadIdx.x/4;
  if ((threadIdx.x >= 16) && (threadIdx.x < 32))
    atomicAggInc(ptr);
}

const int ds = 32;
int main(){

  int *d_d, *h_d;
  h_d = new int[ds];
  cudaMalloc(&d_d, ds*sizeof(d_d[0]));
  cudaMemset(d_d, 0, ds*sizeof(d_d[0]));
  k<<<1,ds>>>(d_d);
  cudaMemcpy(h_d, d_d, ds*sizeof(d_d[0]), cudaMemcpyDeviceToHost);
  for (int i = 0; i < ds; i++)
    std::cout << h_d[i] << " ";
  std::cout << std::endl;
}
$ nvcc -o t1646 t1646.cu -DUSE_OPT
$ cuda-memcheck ./t1646
========= CUDA-MEMCHECK
0 0 0 0 4 4 4 4 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0
========= ERROR SUMMARY: 0 errors
$

（CentOS 7，CUDA 10.1.243，设备 0 为 Tesla V100，设备 1 为 cc3.5 设备）。

我为翘曲发散的情况添加了一个可选优化（即tmask 不是0xFFFFFFFF）。这可以通过定义USE_OPT来选择。

【讨论】：