Acerca de la función de votación warp

Question

La guía de programación CUDA introdujo el concepto de función de votación warp, "_ all", " _any" y "__ballot".

Mi pregunta es: ¿qué aplicaciones usarán estas 3 funciones?

Answer 1

__ballot se utiliza en CUDA-histogram y en la biblioteca CUDA NPP para la generación rápida de máscaras de bits, y combinándolo con __popc intrínseca para hacer una implementación muy eficiente de reducción booleano.

__all y __any se usó en la reducción antes de la introducción de __ballot, aunque no puedo pensar en ningún otro uso de ellos.

Answer 2

El prototipo de __ballot es la siguiente

unsigned int __ballot(int predicate); 

Si predicate es distinto de cero __ballot, devuelve un valor con el conjunto N ésimo bit, donde N es el índice de hilo.

Combinado con atomicOr y __popc, se puede utilizar para acumular el número de subprocesos en cada urdimbre con un predicado verdadero.

De hecho, el prototipo de atomicOr es

int atomicOr(int* address, int val); 

y atomicOr lee el valor apuntado por address, realiza una función OR operación con val, y escribe el valor de nuevo a address y devuelve su valor antiguo como parámetro de retorno

Por otro lado, __popc devuelve la cantidad de bits configurados con un parámetro 32 -bit.

En consecuencia, las instrucciones

volatile __shared__ u32 warp_shared_ballot[MAX_WARPS_PER_BLOCK]; 

const u32 warp_sum = threadIdx.x >> 5; 

atomicOr(&warp_shared_ballot[warp_num],__ballot(data[tid]>threshold)); 

atomicAdd(&block_shared_accumulate,__popc(warp_shared_ballot[warp_num])); 

se pueden utilizar para contar el número de hilos para que el predicado es cierto.

Para más detalles, ver Shane Cook, CUDA programación, Morgan Kaufmann

Answer 3

Como ejemplo de algoritmo que utiliza la API __ballot yo mencionaría la corriente en el núcleo de compactación por D.M Hughes et al. Se usa en la parte de suma de prefijos de la compactación de flujo para contar (por warp) la cantidad de elementos que pasaron el predicado.

Here the paper. In-k Stream Compaction