CUDA/openCL; Umschreiben Verzweigungen als nicht-verzweigende Ausdruck

Question

meiste Zeit ein Zweig in einer CUDA oder OpenCL Programm erforderlich ist, wie:

for (int i=0; i<width; i++) 
{ 
    if(i % threadIdx.x == 0) 
    quantity += i*i; 
} 

kann der Code immer (oder zumindest die meiste Zeit) umgeschrieben werden in nicht-Verzweigung Stil:

for (int i=0; i<width; i++) 
{ 
    quantity += i*i* (i % threadIdx.x != 0); 
} 

scheint der Kompromiss in einer einzigen Kette Schlitz wird entweder läuft im Vergleich zu mehr Berechnungen auf all Threads zu tun (im zweiten Fall wird die Summe ausgeführt wird, immer, nur, dass manchmal der Wert Null)

Unter der Annahme, dass Verzweigungsoperationen mehrere Warp-Slots für jeden möglichen Zweig benötigen, würde man erwarten, dass der zweite konsistent besser ist als der erste. Nun ist meine Frage; Kann ich mich auf den Compiler verlassen, um 1) in 2) zu optimieren, wann immer es Sinn macht, oder gibt es keine breit anwendbaren Kriterien, was bedeutet, dass man ohne Probieren nicht generell entscheiden kann, welches besser ist?

Answer 1

Ich habe nicht viele Erinnerungen an CUDA, aber warum parallelisieren Sie Ihre Schleife nicht? Sie sollten atomare Operationen [1] verwenden, um Ihre Berechnung hinzuzufügen. Ich hoffe, dies wird dir helfen! Tut mir leid, wenn das nicht der Fall ist.

1. Atomic-Operationen: http://supercomputingblog.com/cuda/cuda-tutorial-4-atomic-operations/

Answer 2

Nach meiner Erfahrung - es ist völlig bis zu den Compiler-Autoren diese Art von Grenzfällen zu optimieren.

Also kann ich an Fälle denken, in denen 1) nicht zu 2) gedreht werden kann? Hier ist einer: Ich habe Kernel geschrieben, wo es effizienter war, bestimmte Teile der Berechnungen alle 10 Threads oder etwas ähnliches auszuführen, in welchem ​​Fall eine solche Optimierung nicht abgeleitet werden kann, obwohl es eine mathematische Operation (Dividieren/Subtrahieren) gibt das gleiche Ergebnis unabhängig von bedingten versus "Lauf bei allen, aber null Ergebnisse".

Aber selbst wenn das Überprüfen auf threadId == 0 ein Szenario ist, das häufig genug ist, habe ich keine Ahnung, ob es tatsächlich optimiert ist. Ich wette, dass es auf die Implementierung und sogar das Gerät selbst (CPU vs GPU) abhängt.

Sie müssen es versuchen, um herauszufinden, was am besten funktioniert, nicht nur aus dem oben genannten Grund, sondern auch weil der Arbeitsplaner sich möglicherweise anders verhält, je nachdem, wie teuer es ist, eine Reihe von Threads zu planen/zu starten/zu stoppen im Gegensatz dazu, dass sie alle laufen (und die meisten ein Null-/Identitätsergebnis liefern).

Hoffe, das hilft!

Answer 3

Modulo-Operationen sind ziemlich teuer: Ich bin ziemlich sicher, dass das Hinzufügen des Modulo mehr Zeit in Anspruch nehmen würde als nur eine einzige Anweisung, die nur 1 Thread ausführt. Ihre einzige Branching-Anweisung, eine if ohne else, hängt nur die anderen Threads währenddessen, wenn die Statement ausgeführt wird. Da GPUs für sehr schnelle Kontextwechsel optimiert sind, sollten die Kosten dafür sehr gering sein.

Es wird jedoch davon abgeraten, lange Branching-Anweisungen zu verwenden: zu viel serielle Berechnung auf der GPU (dh ein Thread, der all die Arbeit erledigt) negiert den Vorteil der Parallelität.