CUDA: Wie atomar diese Anweisungen ausführen (4 Adressen)

Question

ich einen Algorithmus in Cuda implementieren, die folgenden Schritte durchführen muss:

ein Array x (in gemeinsam genutzten mem) und zeigen einige Gerätefunktion f,

• ein Paar von Indizes wählen (i,j)x (zufällig).
• Berechnen y = f(x[i], x[i - 1], x[j], x[j + 1]).
• Basierend auf y entscheiden, ob die Positionen von x [i] und x [j] ausgetauscht werden sollen.

Das Problem ist, dass die Funktion f auf 4 Werten im gemeinsam genutzten Speicher abhängt, von denen alle garantiert werden müssen, bis nach dem Tausch unverändert bleiben.

Für eine Minute dachte ich, dies könnte der Körper einer critical section sein, aber ich sehe nicht, wie ich eine einzige Lock-Adresse verwenden könnte, um 4 Variablen zu sperren. Das Hauptproblem, denke ich, ist, dass, wenn ein Thread an (i,j) arbeitet, andere Threads nicht auf einem Paar (k,l) arbeiten dürfen, wobei k oder l{i, i-1, j, j+1} sind.

EDIT

Gleich nach der Veröffentlichung, eine Idee in den Sinn kam ... Wäre es möglich, eine Kaskade von Schlössern zu haben? Erstes Schloss x[i], wenn das gelingt, sperren Sie x[i-1] usw. für alle 4 Werte. Nur wenn die endgültige Sperre erfolgreich ist, führen Sie die oben genannten Schritte aus. Ich werde experimentieren und diese Frage für andere Vorschläge offen halten.

Answer 1

CUDA ist sehr lock-unfreundlich und kritisch-Abschnitt-unfreundlich :) Einer von vielen Gründen ist, dass es in einem 32-weiten SIMD-Modus arbeitet. Dies kann zu unerwarteten Deadlocks führen. Betrachten Sie zum Beispiel:

__shared__ int crit; 
crit = 0; 
__syncthreads(); 
int old; 
do { 
    old = atomicCas(&crit, 0, 1); 
} while (old==1); 
//critical section 
crit = 0; 

Die intetion ist, dass Threads aktiv an der do-while-Schleife warten. Nur ein Thread gibt die Schleife zur gleichen Zeit aus, führt eine Aktion in der kritischen Sektion durch und setzt dann crit auf 0. In CUDA gibt ein Warp-Scheduler 31 Threads in der Schleife immer Vorrang vor dem 1 Thread, der beendet wird. Da Warps in SIMD ausgeführt werden, wird der Thread im kritischen Abschnitt nie ausgeführt und Sie erhalten einen unerwarteten Deadlock.

---

Aus diesem Grund würde ich dringend empfehlen, kritische Abschnitte vollständig zu vermeiden.

Jetzt kenne ich die Details Ihres Algorithmus nicht. Ich nehme an, dass Sie einige "Master" haben for/while Schleife und in jeder Iteration wählen Sie ein zufälliges Paar für einen möglichen Austausch.

Sie sagen, die Kollisionen passieren nicht oft. Wenn dies der Fall ist, könnten Sie einfach eines der Konfliktpaare vollständig löschen, anstatt darauf zu warten, dass es erfolgreich ist?

Wenn das etwas ist, was Sie akzeptieren würden, dann wäre nur das Erkennen der Kollision das Problem, nicht die Aktion, die Sie danach ausführen. Zur Erkennung von Kollisionen Sie könnten zum Beispiel:

• Nach jedem Thread mit einem Paar Kandidaten kommt, sortieren Sie den Paar-Indizes und dann prüfen, die von den Nachbarn gehaltenen Werte.

• Haben Sie eine Flagge Array f von der gleichen Größe wie x und atomicCas auf sie 4 mal, ähnlich zu dem, was Sie vorgeschlagen. Wenn f im Shared Memory ist, sollte es nicht teuer sein.

Jetzt, wenn ein Thread sieht, dass es in Konflikt ist, tut es nichts. Wartet nur auf alle anderen Threads, um ihre Arbeit abzuschließen, __syncthreads, und geht dann für die nächste Iteration des Masters for/while Schleife.

Der Unterschied zu Ihrer vorgeschlagenen Lösung besteht darin, dass Ihr Thread seine Arbeit abbricht, anstatt zu warten, wenn Sie die Sperre nicht bestehen.

Answer 2

Es scheint mir, dass Sie dies stark überdenken. Wenn alle Speichertransaktionen serialisiert werden müssen, damit die Operation threadsicher ist, dann ist die einfachste Lösung, einen Thread pro Block die Operation ausführen zu lassen. So etwas wie

if (threadIdx.x == 0) // assume 1D block for simplicity 
{ 
    y = f(x[i], x[i - 1], x[j], x[j + 1]); 
    compare_and_swap(y, x[i], x[j]; 
} 
__syncthreads(); 

wird gut funktionieren, da das Array auf ist im gemeinsam genutzten Speicher betrieben wird, so ein garantierter einzelner Thread pro Block, der die Operation ausführt, und es gibt keine Lese-nach-Schreib-Gefahren. In der Praxis sollte dieser Ansatz nicht langsamer sein als ein gesamter Block von Threads, der in Konkurrenz um eine Sperre steht, oder eine große Anzahl von serialisierten atomaren Speichertransaktionen.