CUDA: __synctreads() innerhalb if Anweisungen

Question

Ich habe eine Frage über CUDA-Synchronisierung. Insbesondere brauche ich eine Klarstellung bezüglich der Synchronisierung von if-Anweisungen. Ich meine, wenn ich eine __synctreads() unter den Rahmen einer if-Anweisung setze, die von einem Bruchteil der Threads innerhalb des Blocks getroffen wird, was passiert dann? Ich dachte, dass einige Threads "für immer" auf die anderen Threads warten werden, die den Synchronisationspunkt nicht treffen. Also, ich schrieb und ausgeführtem Code einige Beispiel zu inspizieren:

__global__ void kernel(float* vett, int n) 
{ 
    int index = blockIdx.x*blockDim.x + threadIdx.x; 
    int gridSize = blockDim.x*gridDim.x; 

    while(index < n) 
    { 
     vett[index] = 2; 
     if(threadIdx.x < 10) 
     { 
      vett[index] = 100; 
      __syncthreads(); 
     } 
     __syncthreads(); 

     index += gridSize; 
    } 
} 

Erstaunlicherweise habe ich festgestellt, dass der Ausgang ein ziemlich "normal" war (64 Elemente, Blockgröße 32):

100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 

Also ich leicht mein Code in folgenden Weise geändert:

__global__ void kernel(float* vett, int n) 
{ 
    int index = blockIdx.x*blockDim.x + threadIdx.x; 
    int gridSize = blockDim.x*gridDim.x; 

    while(index < n) 
    { 
     vett[index] = 2; 
     if(threadIdx.x < 10) 
     { 
      vett[index] = 100; 
      __syncthreads(); 
     } 
     __syncthreads(); 
      vett[index] = 3; 
     __syncthreads(); 

     index += gridSize; 
    } 
} 

Und der Ausgang war:

3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 

Wieder lag ich falsch: Ich dachte, dass die Threads in der if-Anweisung nach dem Ändern des Elements des Vektors in einem Wartezustand bleiben und nie aus dem if-Bereich herauskommen würden. Also ... könnten Sie bitte klären, was passiert ist? Wird ein Thread, der nach einem Synchronisationspunkt kommt, die an der Barriere wartenden Threads entsperren? Wenn Sie meine Situation reproduzieren müssen, habe ich CUDA Toolkit 5.0 RC mit SDK 4.2 verwendet. Vielen Dank im Voraus.

Answer 1

Sie dürfen __syncthreads() nicht verwenden, solange die Anweisung nicht immer in allen Threads innerhalb eines Threadblocks erreicht wird. Vom programming guide (B.6):

__syncthreads() wird in bedingtem Code erlaubt, aber nur dann, wenn die Bedingung auswertet identisch über den gesamten Thread-Block, da sonst die Codeausführung ist wahrscheinlich unbeabsichtigte Nebenwirkungen hängen oder zu produzieren.

Grundsätzlich ist Ihr Code kein wohlgeformtes CUDA-Programm.

Answer 2

CUDA-Modell ist MIMD, aber aktuelle NVIDIA-GPUs implementieren __syncthreads() mit Warp-Granularität anstelle von Thread. Es bedeutet, das sind warps inside a thread-block, die nicht unbedingt synchronisiert sind threads inside a thread-block. __syncthreds() wartet auf alle 'Warps' des Thread-Blockes, um die Barriere zu treffen oder das Programm zu verlassen. Weitere Informationen finden Sie unter Henry Wong's Demistifying paper.

Answer 3

Kurz gesagt, ist das Verhalten undefined. So kann es manchmal tun, was Sie wollen, oder es wird nicht, oder (sehr wahrscheinlich) wird nur Ihren Kernel hängen oder zum Absturz bringen.

Wenn Sie wirklich neugierig sind, wie die Dinge intern arbeiten, müssen Sie daran denken, dass Threads nicht unabhängig ausgeführt werden, sondern eine Warp (Gruppe von 32 Threads) gleichzeitig.

Dies erzeugt natürlich ein Problem mit bedingten Verzweigungen, bei denen die Bedingung im gesamten Warp nicht einheitlich ausgewertet wird. Das Problem wird gelöst, indem nacheinander beide Pfade ausgeführt werden, wobei jeweils die Threads deaktiviert sind, die diesen Pfad nicht ausführen sollen. IIRC bei bestehender Hardware wird zuerst die Verzweigung genommen, dann wird der Pfad ausgeführt, wo die Verzweigung nicht genommen wird, aber dieses Verhalten ist undefiniert und somit nicht garantiert.

Diese separate Ausführung von Pfaden setzt sich bis zu einem gewissen Punkt fort, für den der Compiler bestimmen kann, dass er von allen Threads der zwei separaten Ausführungspfade ("Konvergenzpunkt" oder "Synchronisationspunkt") garantiert erreicht wird. Wenn die Ausführung des ersten Code-Pfades diesen Punkt erreicht, wird sie gestoppt und stattdessen der zweite Code-Pfad ausgeführt. Wenn der zweite Pfad den Synchronisationspunkt erreicht, werden alle Threads erneut aktiviert und die Ausführung wird von dort aus gleichmäßig fortgesetzt.

Die Situation wird komplizierter, wenn vor der Synchronisation eine andere bedingte Verzweigung auftritt. Dieses Problem wird mit einem Stapel von Pfaden gelöst, die noch ausgeführt werden müssen (zum Glück ist das Wachstum des Stapels begrenzt, da wir höchstens 32 verschiedene Codepfade für einen Warp haben können).

Wo die Synchronisationspunkte eingefügt werden ist undefined und variiert sogar leicht zwischen den Architekturen, so wieder gibt es keine Garantien. Der einzige (inoffizielle) Kommentar, den Sie von Nvidia bekommen, ist, dass der Compiler ziemlich gut darin ist, optimale Synchronisationspunkte zu finden. Es gibt jedoch häufig subtile Probleme, die den optimalen Punkt weiter nach unten verschieben können, als Sie erwarten würden, insbesondere wenn die Threads frühzeitig beendet werden.

nun das Verhalten der __syncthreads zu verstehen() Richtlinie (das in eine bar.sync Anweisung in PTX übersetzt) ​​ist es wichtig, diese Anweisung zu erkennen, dass nicht pro Thread ausgeführt wird, sondern für die gesamte Kette auf einmal (unabhängig davon, ob irgendwelche Threads deaktiviert sind oder nicht), da nur die Warps eines Blocks synchronisiert werden müssen. Die Threads eines Warps werden bereits synchron ausgeführt, und die weitere Synchronisierung hat entweder keine Auswirkungen (wenn alle Threads aktiviert sind) oder führt zu einem Deadlock, wenn versucht wird, die Threads von verschiedenen bedingten Codepfaden zu synchronisieren.

Sie können sich von dieser Beschreibung zu dem Verhalten Ihres bestimmten Codeabschnitts hinarbeiten. Aber denken Sie daran, dass dies alles undefined ist, gibt es keine Garantien, und sich auf ein bestimmtes Verhalten verlassen kann Ihren Code jederzeit brechen.

Sie können sich die PTX manual für weitere Details ansehen, insbesondere für die bar.sync Anweisung, die __syncthreads() kompiliert. Henry Wong's "Demystifying GPU Microarchitecture through Microbenchmarking" paper, unten von Ahmad referenziert, ist auch lesenswert. Auch wenn die Architektur und die CUDA-Version veraltet sind, scheinen die Abschnitte über bedingte Verzweigungen und __syncthreads() immer noch allgemein gültig zu sein.

Answer 4

__synctreads() wird verwendet, um Threads innerhalb eines Blocks zu synchronisieren. Das bedeutet, dass alle Threads im Block darauf warten, dass alle abgeschlossen sind, bevor Sie fortfahren.

Betrachten Sie den Fall, in dem sich einige Threads in einem Block befinden, der die if-Anweisung eingibt, und einige nicht. Die wartenden Threads werden blockiert. für immer warten.

Im Allgemeinen ist es kein guter Stil, in eine if-bedingte Anweisung zu synchronisieren. Am besten, es zu vermeiden und Ihren Code neu zu gestalten, wenn Sie ihn haben. Der Zweck der Synchronisierung besteht darin, sicherzustellen, dass alle Threads gemeinsam ablaufen, warum filtern Sie sie dann mithilfe der if-Anweisung aus?

Hinzufügen, wenn Synchronisation über die Blöcke hinweg erforderlich ist. Relaunch des Kernels ist erforderlich.