Koaleszierte vs. gestridete Lesezugriffe

Die Koaleszierung hängt davon ab, welche Adresse jeder Thread anspricht. Die Zuordnung von Thread zu Adresse entscheidet darüber, ob eine Transaktion den gesamten Warp bedient.

Wenn benachbarte Threads benachbarte Adressen lesen, deckt der Warp einen zusammenhängenden Bereich ab. Dieses übersichtliche Layout ist ein koaleszierter Zugriff.

Thread 0 greift auf Index 0 zu, Thread 1 auf Index 1 und so weiter. Alle 32 Adressen liegen in einer ausgerichteten Cache-Zeile, daher benötigt die GPU nur eine einzige Transaktion.

Ein Stride ist der feste Abstand zwischen den Adressen, auf die aufeinanderfolgende Threads zugreifen. Sobald dieser Abstand größer als ein Element wird, beginnt die Koaleszierung zu scheitern.

Bei einem Stride von 2 liest Thread 0 den Wert 0, Thread 1 den Wert 2 und Thread 2 den Wert 4. Der Warp erstreckt sich nun über doppelt so viel Speicher und benötigt daher mehr Transaktionen.

Bei doppeltem Stride werden ungefähr doppelt so viele Cache-Zeilen berührt. Große Strides können dazu führen, dass ein Warp viele Transaktionen benötigt, obwohl er aus jeder Cache-Zeile nur einen kleinen Teil verwendet.

Wenn Sie jedem Thread eine ganze Spalte einer im Zeilenformat gespeicherten Matrix zuweisen, entsteht ein riesiger Stride. Im Code sieht das ordentlich aus, verteilt aber stillschweigend jeden Warp über viele Speicherstellen.

Oft besteht die Lösung einfach darin, den Index zu ändern, der sich mit dem Thread am schnellsten ändert. Lassen Sie aufeinanderfolgende Threads auf aufeinanderfolgende Speicherstellen zugreifen, und die Lesezugriffe werden wieder koalesziert.

Selbst ein koalesziertes Muster leidet, wenn es in der Mitte einer Cache-Zeile beginnt. Ein nicht ausgerichteter Offset verschiebt den Warp über eine Grenze und teilt einen Lesezugriff in zwei auf.

Um ein Zugriffsmuster zu beurteilen, betrachten Sie nicht nur einen Thread. Stellen Sie sich alle 32 Lanes gleichzeitig vor und fragen Sie sich, wie viele Cache-Zeilen ihre Adressen gemeinsam abdecken. 🔍

Lassen Sie dem sich am schnellsten ändernden Index threadIdx.x folgen. Diese eine Gewohnheit hält die meisten globalen Lesezugriffe kostenlos koalesziert.

Sie haben gesehen, dass koaleszierte Lesezugriffe benachbarte Elemente zusammenhalten, während ein Stride sie über mehrere Cache-Zeilen verteilt. Lassen Sie threadIdx.x den sich am schnellsten ändernden Index bestimmen. 🎉