La guida alla programmazione CUDA ha introdotto il concetto di funzione di voto a curvatura, "_ all", " _any" e "__ballot".Informazioni sulla funzione di voto ordito
La mia domanda è: quali applicazioni utilizzeranno queste 3 funzioni?
__ballot viene utilizzato in CUDA-histogram e in libreria CUDA NPP per una rapida generazione di bitmasks, e combinandola con __popc intrinseca per fare un molto efficiente attuazione della riduzione booleana.
__all e __any è stato utilizzato nella riduzione prima dell'introduzione di __ballot, anche se non riesco a pensare a nessun altro uso.
Il prototipo del __ballot è il seguente
unsigned int __ballot(int predicate);
Se predicate è diverso da zero, __ballot restituisce un valore con il set N esimo bit, dove N è l'indice di thread.
Combinato con atomicOr e __popc, può essere utilizzato per accumulare il numero di thread in ogni ordito con un predicato vero.
In effetti, il prototipo di atomicOr è
int atomicOr(int* address, int val);
e atomicOr legge il valore puntato da address, esegue un OR un'operazione bit per bit con val, e scrive il valore a address e restituisce il vecchio valore come parametro di ritorno.
Dall'altro lato, __popc restituisce il numero di bit impostato con un parametro 32 -bit.
Pertanto, le istruzioni
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]));
può essere utilizzato per contare il numero di fili per il quale il predicato è vero.
Per maggiori dettagli, si veda Shane Cook, CUDA programmazione, Morgan Kaufmann
Come esempio di algoritmo che utilizza API __ballot vorrei menzionare l'nel kernel stream compattazione da D.M Hughes et Al. Viene utilizzato nella parte somma del prefisso della compattazione del flusso per contare (per distorsione) il numero di elementi che hanno superato il predicato.
questo suona super-interessante. C'è qualche implementazione che posso guardare? – aatish
sì, ho scritto una versione migliorata di tale algoritmo. https://github.com/knotman90/cuStreamComp. Per favore chiedimi se hai bisogno di chiarimenti o benchmark. –
In realtà sarebbe bello avere alcuni benchmark contro la libreria thrust. Inoltre, penso che nella riga 78 di cuCompactor.cuh, dovrebbe essere possibile avere un altro array globale chiamato d_output_index che conterrebbe il valore di idx corrispondente alla provenienza dei dati originali. Ho ragione? – aatish