Оптимизация одновременной проверки с потоком

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 устройств, вам необходимо будет осуществить и сокращение (упражнение для читателя, так как он хорошо документирован). Для специального случая вычислительной способности 1.3 вы можете оптимизировать сокращение AND с помощью команд голосования по варпу, используя встроенную функцию __all(), предоставленную в заголовках CUDA.

редактировать:

ОК, так как 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]; 
}