前8.x等效于CUDA中的__reduce_max_sync()

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
2. 它需要CUDA计算能力8.x,而我需要支持5.2计算能力的设备

你能帮我解决以上问题吗?