8.x 之前版本相当于 CUDA 中的 __reduce_max_sync()

Question

cuda-memcheck

在执行以下操作的代码中检测到竞争条件：

condition = /*different in each thread*/;
shared int owner[nWarps];
/* ... owner[i] is initialized to blockDim.x+1 */
if(condition) {
    owner[threadIdx.x/32] = threadIdx.x;
}

基本上，这段代码根据某些条件计算每个扭曲的所有者线程。对于某些扭曲，可能没有所有者，但对于某些扭曲，所有者的数量可能超过 1，然后会发生竞争条件，因为多个线程将值分配给同一共享内存区域。

尝试文档后，我认为我需要的可以通过以下方式完成：

const uint32_t mask = __ballot_sync(0xffffffff, condition);
if(mask != 0) {
    const unsigned max_owner = __reduce_max_sync(mask, threadIdx.x);
    if(threadIdx.x == max_owner) {
        // at most 1 thread assigns here per warp
        owner[threadIdx.x/32] = max_owner;
    }
}

但是，我的尝试有两个问题：

1. 我真的不需要找到最大线程 - 如果存在带有

   condition==true

的线程，为每个经纱选择任意 1 个线程就足够了
2. 需要CUDA计算能力8.x，而我需要支持5.2计算能力的设备

您能帮我解决以上问题吗？

Answer 1

以下功能似乎可以解决问题：

void SetOwnerThread(int* dest, const bool condition) {
  const uint32_t mask = __ballot_sync(0xffffffff, condition);
  if(!mask) {
    return;
  }
  const uint32_t lowest_bit = mask & -mask;
  const uint32_t my_bit = (1 << (threadIdx.x & 31));
  if(lowest_bit == my_bit) {
    dest = threadIdx.x;
  }
}

Answer 2

不，这不是正确的方法。 您应该使用

__ffs``.  

__ffs` 来告诉您哪个是第一个设置位，其中 1 是 LSB，32 是 MSB。 0 表示没有设置位。

const uint32_t min_owner_plus_1 = __ffs(__ballot_sync(0xffffffff, condition));
if (min_owner_plus_1 != 0) {
  const unsigned min_owner = min_owner_plus_1 - 1;
  if (threadIdx.x == min_owner) {
    // at most 1 thread assigns here per warp
    owner[threadIdx.x/32] = min_owner;
  }
}

This will only use a single warp sync operation instead of 2.