Über Warp Voting-Funktion

Question

Die CUDA Programmieranleitung eingeführt das Konzept der Warp-Vote-Funktion, "_ alle", " _any" und "__ballot".

Meine Frage ist: Welche Anwendungen werden diese 3 Funktionen verwenden?

Answer 1

__ballot in CUDA-histogram und in CUDA NPP-Bibliothek für die schnelle Erzeugung von Bitmasken verwendet und mit __popc eigen macht eine sehr effiziente Implementierung von boolean Reduktion kombiniert.

__all und __any wurde in der Reduktion vor der Einführung von __ballot verwendet, obwohl ich an keine andere Verwendung von ihnen denken kann.

Answer 2

Der Prototyp ist die folgende __ballot

unsigned int __ballot(int predicate); 

Wenn predicate ungleich Null ist, __ballot Wert mit dem eingestellten N te Bit zurückkehrt, wo der Faden N Index ist.

Kombiniert mit atomicOr und __popc kann es verwendet werden, um die Anzahl der Threads in jedem Warp mit einem echten Prädikat zu akkumulieren.

Tatsächlich ist der Prototyp von atomicOr ist

int atomicOr(int* address, int val); 

und atomicOr der Wert liest von address zu spitz, führt eine bitweise OR Betrieb mit val und schreibt den Wert wieder auf address und gibt seinen alten Wert als Rückgabeparameter. Auf der anderen Seite gibt __popc die Anzahl der Bits zurück, die mit einem 32 -bit-Parameter gesetzt wurden.

Dementsprechend werden die Anweisungen

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])); 

kann die Anzahl von Threads, für die das Prädikat wahr ist zu zählen, verwendet werden.

Weitere Einzelheiten finden Sie Shane Cook CUDA-Programmierung, Morgan Kaufmann

Answer 3

Als Beispiel eines Algorithmus, der __ballot API verwendet i die In-Kernel-Stream Kompaktierung von D. M. Hughes et al erwähnen würde. Er wird im Präfix-Summen-Teil der Stream-Komprimierung verwendet, um (pro Warp) die Anzahl der Elemente zu zählen, die das Prädikat passiert haben.

Here the paper. In-k Stream Compaction