parallel-processing - 优化线程同时检查

Question

我有一个设备函数，它使用线程检查字节数组，每个线程检查数组中不同字节的某个值并返回 bool true 或 false。

我怎样才能有效地确定所有支票是否都返回真或其他？

Answer 1

// returns true if predicate is true for all threads in a block
__device__ bool unanimous(bool predicate) { ... }

__device__ bool all_the_same(unsigned char* bytes, unsigned char value, int n) {
    return unanimous(bytes[threadIdx.x] == value);
}

的实现unanimous()取决于硬件的计算能力。对于计算能力 2.0 或更高版本的设备，这是微不足道的：

__device__ bool unanimous(bool predicate) { return __syncthreads_and(predicate); }

对于计算能力 1.0 和 1.1 的设备，您需要实现 AND 缩减（读者练习，因为它有详细的文档）。__all()对于计算能力 1.3 的特殊情况，您可以使用CUDA 标头中提供的内在函数，使用扭曲投票指令优化 AND 减少。

编辑：

好的，因为玩家在评论中询问。在 sm_13 硬件上，您可以执行此操作。

// returns true if predicate is true for all threads in a block
// note: supports maximum of 1024 threads in block as written
__device__ bool unanimous(bool predicate) {
    __shared__ bool warp_votes[32];
    if (threadIdx.x < warpSize) warp_votes[threadIdx.x] = true;
    warp_votes[threadIdx.x / warpSize] = __all(pred);
    __syncthreads();
    if (threadIdx.x < warpSize) warp_votes[0] = __all(warp_votes[threadIdx.x];
    __syncthreads();
    return warp_votes[0];
}