CUDA: Synchronisieren von Threads

Question

Fast überall las ich über die Programmierung mit CUDA gibt es einen Hinweis auf die Bedeutung, die alle Fäden in einer Kette das gleiche tun.
In meinem Code habe ich eine Situation, wo ich nicht eine bestimmte Bedingung vermeiden. Es sieht wie folgt aus:

// some math code, calculating d1, d2
if (d1 < 0.5)
{
    buffer[x1] += 1;  // buffer is in the global memory
}
if (d2 < 0.5)
{
    buffer[x2] += 1;
}
// some more math code.

Einige der Themen eingehen könnten in einen für die Bedingungen, könnten einige in beiden eingeben und andere möglicherweise nicht in entweder eingeben.

Um nun alle die Fäden wieder zu machen „die gleiche Sache“ wieder nach den Bedingungen, soll ich sie nach den Bedingungen mit __syncthreads() synchronisieren? Oder hat irgendwie das geschieht automatisch?
Kann zwei Threads sein nicht das gleiche tun auf Grund einer von ihnen eine Operation ist hinter, damit es für alle zu ruinieren? Oder gibt es einige hinter den Kulissen Anstrengung, um sie die gleiche Sache wieder nach einem Zweig zu tun?

Answer 1

Innerhalb einer Kette wird kein Thema „get ahead“ von irgendwelchen anderen. Wenn es eine bedingte Verzweigung und es durch einige Fäden in der Kette genommen wird, andere aber nicht (aka warp „Divergenz“), wird die anderen Threads nur im Leerlauf, bis der Zweig abgeschlossen ist und sie alle „zusammenlaufen“ wieder zusammen auf einem gemeinsamen Unterricht . Also, wenn Sie nur innerhalb verzugs Synchronisation von Threads müssen, was geschieht „automagically.“

Aber verschiedene Verwerfungen sind nicht auf diese Weise synchronisiert. Also, wenn Ihr Algorithmus erfordert, dass bestimmte Operationen in vielen Verwerfungen abgeschlossen sein, dann müssen Sie explizite Synchronisation Anrufe verwenden (siehe das CUDA-Programmierhandbuch, Abschnitt 5.4).

---

EDIT:. in den nächsten Absätzen neu organisiert, einige Dinge klarstellen

Es gibt wirklich zwei verschiedene Dinge hier:. Instruction Synchronisation und Speicher Sichtbarkeit

• __syncthreads() erzwingt Anweisung Synchronisation und sorgt für Speicher Sichtbarkeit, aber nur innerhalb eines Blocks, nicht über Blöcke (CUDA-Programmierhandbuch, Anhang B.6). Es ist nützlich für Schreib dann ablesbare auf gemeinsam genutzten Speicher, ist aber nicht geeignet für die globalen Speicherzugriff zu synchronisieren.

• __threadfence() sorgt für globale Speicher Sichtbarkeit aber führt keine Anweisung Synchronisation, so in meiner Erfahrung es nur von begrenztem Nutzen ist (aber Beispielcode in Anhang B.5).

• Globale Anweisung Synchronisation innerhalb eines Kernels nicht möglich. Wenn müssen Sie auf alle Threads getan f() vor g() auf jedem Thread, Split f() und g() in zwei verschiedene Kernel aufrufen und nennen sie seriell von dem Host.

• Wenn Sie nur geteilt oder globale Zähler müssen erhöhen, sollten Sie mit der Atom Inkrementfunktion atomicInc() (Anlage B.10). Im Fall des Codes oben, wenn x1 und x2 nicht global eindeutig sind (über alle Fäden in Ihrem Netz), nicht-atomare Schritte werden in einem Rennen-Zustand, ähnlich den letzten Absatz der Anlage B.2.4 führen.

Schließlich beachten Sie, dass alle Operationen auf den globalen Speicher und Synchronisationsfunktionen insbesondere (einschließlich atomics) für die Leistung schlecht sind.

Ohne zu wissen, das Problem, das Sie lösen es schwer ist, zu spekulieren, aber vielleicht können Sie Ihren Algorithmus neu zu gestalten gemeinsam genutzten Speicher zu verwenden, anstatt die globalen Speicher an einigen Stellen. Dadurch wird die Notwendigkeit zur Synchronisation reduzieren und geben Sie eine Leistungssteigerung.

Answer 2

Aus dem Bereich 6.1 der CUDA Best Practices Guide:

  

Jeder Ablaufsteuerbefehl (wenn, Schalter, tun, denn während) erheblich beeinflussen   der Befehlsdurchsatz von Fäden der gleichen Kette zu divergieren zu verursachen; das ist,   unterschiedliche Ausführungspfaden zu folgen. Wenn dies geschieht, um die verschiedenen Ausführungspfade   müssen serialisiert werden, denn dies ist die Gesamtzahl der ausgeführten Befehle zu erhöhen   Kette. Wenn alle verschiedenen Ausführungspfade abgeschlossen haben, konvergieren die Fäden   zurück auf den gleichen Ausführungspfad.

Also, Sie brauchen nichts Besonderes zu tun.

Answer 3

In Gabriels Antwort:

"Globale Anweisung Synchronisation ist nicht möglich, innerhalb eines Kernel. Wenn Sie f () auf allen Threads erfolgt vor dem Aufruf von g () auf jedem Thread, Split f () und g () in zwei verschiedene Kernel und nennen sie seriell von der Wirt. "

Was passiert, wenn der Grund, Sie f () und g () in demselben Thread benötigen, weil Sie Registerspeicher verwenden, und Sie möchten sich registrieren oder gemeinsam genutzte Daten von f zu g bekommen? Das heißt, für mein Problem, der ganze Grund für Blöcke für die Synchronisierung ist, da die Daten von f in g benötigt wird - und zu einem Kernel Ausbrechen würde eine große Menge an zusätzlichen globalen Speicher erfordern Registerdaten von f zu g zu übertragen, was ich d‘vermeiden

Answer 4

Die Antwort auf Ihre Frage lautet: Nein. Sie brauchen nichts Besonderes zu tun. Wie auch immer, Sie können dieses Problem beheben, anstelle des Codes Sie etwas tun können:

buffer[x1] += (d1 < 0.5);
buffer[x2] += (d2 < 0.5);

Sie sollten überprüfen, ob Sie gemeinsam genutzten Speicher und Zugriff globale Speicher in einem verschmolzenen Muster verwenden können. Auch sicher sein, dass Sie nicht mehr als 1 Thread zu dem gleichen Index schreiben tun wollen.