parallel-processing - スレッド同時チェックの最適化

Question

スレッドを使用してバイト配列をチェックするデバイス関数があり、各スレッドは特定の値について配列内の異なるバイトをチェックし、bool true または false を返します。

すべてのチェックが true を返したかどうかを効率的に判断するにはどうすればよいですか?

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 リダクションを最適化できます。

編集：

OK、gamerx がコメントで質問しているので。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];
}