優化線程同步檢查

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的設備，你將需要實現和減少（讀者練習，因爲它是有據可查）。對於計算能力1.3的特殊情況，您可以使用CUDA頭文件中提供的內部函數，使用warp投票指令優化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]; 
}