Эквивалент __reduce_max_sync() до 8.x в CUDA

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;
    }
}

Однако у моей попытки есть 2 проблемы:

1. Мне действительно не нужно находить максимальный поток - достаточно выбрать любой 1 поток для каждой деформации, если есть поток с condition==true
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;
  }
}