GPU Shared Memory Jordanland Konflikt

Question

Ich versuche zu verstehen, wie Bank Konflikte stattfinden.
Wenn ich ein Array der Größe 256 in der globalen Speicher habe, und ich habe 256 Threads in einem einzigen Block, und ich möchte das Array gemeinsam genutzten Speicher kopieren. daher jeder Faden Kopien eines Elements.

shared_a[threadIdx.x]=global_a[threadIdx.x]

hat diese einfache Aktion Ergebnis in einem Bankkonflikt?

nehmen nun an, dass die Größe des Arrays ist größer als die Anzahl der Threads, so bin ich jetzt diese den globalen Speicher auf den gemeinsamen Speicher kopieren Sie mit:

tid = threadIdx.x;
for(int i=0;tid+i<N;i+=blockDim.x)
     shared_a[tid+i]=global_a[tid+i];

hat den obigen Code Ergebnis in einem Bankkonflikt?

Answer 1

Der beste Weg, dies zu überprüfen, wäre der Code zum Profil des „Compute von Visual Profiler“ verwenden; dies kommt mit dem CUDA Toolkit. Auch gibt es einen großen Abschnitt in GPU Gems 3 auf diesem - „39.2.3 Vermeiden Bank Konflikte“.

„ Wenn mehrere Threads in der gleichen Kette Zugriff auf die gleiche Bank, eine Bankkonflikt, es sei denn alle Fäden der Kette Zugriff erfolgt die gleiche Adresse innerhalb der gleichen 32-Bit-Wort “ - Das erste, was es gibt 16 Speicherbänke jeweils 4 Bytes breit. So im Wesentlichen, wenn Sie jeder Faden in einem halben Kette Speicher aus dem gleichen 4bytes in einer gemeinsam genutzten Speicherbank zu lesen, Sie gehen Bankkonflikte haben und Serialisierung usw.

OK, damit Ihr erstes Beispiel :

Zuerst läßt vermuten, Ihre Arrays zum Beispiel von der Art sagen werden int ( ein 32-Bit-Wort ). Ihr Code speichert diese ints in der gemeinsamen Speicher, über jede Hälfte WARP der K-ten Thread zu der K-ten Speicherbank zu speichern. So zum Beispiel Thread 0 der ersten Hälfte warp wird shared_a[0] speichern, die in der ersten Speicherbank ist, 1 Thread zu shared_a[1] speichern, die jeweils die Hälfte warp hat 16 Fäden diese Karte auf die 16 4Byte Banken. In der nächsten halben Kette wird der erste Thread jetzt speichert seinen Wert in shared_a [16], die in der ersten Speicherbank wieder. Wenn Sie also ein 4-Byte Wort so int verwenden, schwimmen etc dann Ihr erstes Beispiel wird in einem Bankkonflikt führen. Wenn Sie ein 1-Byte-Wort wie char, in der ersten Hälfte Kettfäden verwenden 0, 1, 2 und 3 werden alle speichern ihre Werte an die erste Bank des gemeinsam genutzten Speichers, der einen Bankkonflikt verursacht.

Zweites Beispiel :

Auch hier wird das alles hängt von der Größe des Wortes Sie verwenden, aber für das Beispiel werde ich ein 4-Byte Wort verwenden. So suchen Sie in der ersten Hälfte Kette:

Anzahl der Threads = 32

N = 64

Thread 0: Wird auf 0 schreiben, 31, 63 Thema 1: Schreiben wird auf 1, 32

Alle Fäden über die Hälfte Warp gleichzeitig ausgeführt, so dass die Schreibvorgänge auf Shared Memory sollte nicht Ursache Bank Konflikte. Ich werde diese verdoppeln müssen überprüfen, though.

Hope, das hilft, tut mir leid für die große Antwort!

Answer 2

In beiden Fällen Threads Zugriffsspeicher mit konsekutiven Adressen geteilt. Es hängt von der Elementgröße des gemeinsam genutzten Speichers, aber aufeinanderfolgenden Zugriff auf gemeinsam genutzten Speicher durch eine Kette von Fäden führt nicht zu einem Bankkonflikt für „kleine“ Elementgrößen.

dieser Code mit NVIDIA von Visual Profiler zeigt, dass für Elementgröße von weniger als 32 und ein Vielfaches von 4 (4, 8, 12, ..., 28) aufeinanderfolgenden Zugriff auf den gemeinsamen Speicher führt nicht zu einem Bankkonflikt. Elementgröße von 32, führt jedoch zu einer Bankkonflikt.

---

Antwort von Ljdawson enthält einige veraltete Informationen:

... Wenn Sie ein 1-Byte-Wort wie char verwenden, in der ersten Hälfte Kettfäden 0, 1, 2 und 3 werden alle speichern ihre Werte an die erste Bank des gemeinsam genutzten Speichers, die einen Bankkonflikt verursachen.

Dies kann für alten GPUs wahr sein, aber für die letzten GPUs mit cc> = 2.x, sie verursacht keine Bank Konflikte, effektiv aufgrund des Broadcast-Mechanismus ( link ). Zitat aus dem CUDA C Programmieranleitung (v8.0.61) G3.3. Shared Memory .

Eine gemeinsame Speicheranforderung für eine Kette keinen Bankkonflikt zwischen zwei Threads, das jede Adresse innerhalb der gleichen 32-Bit-Wort-Zugriff generiert (auch wenn die beiden Adressen in der gleichen Bank fallen): In diesem Fall für Lesezugriffe wird das Wort an den anfordernden Threads Broadcast (mehrere Worte in einer einzelnen Transaktion übertragen werden können) und für Schreibzugriffe wird jede Adresse, die von nur einem des Gewindes (welcher Thread führt das schreiben wird nicht definiert).

geschrieben

Das bedeutet insbesondere, dass es keine Bank Konflikte, wenn ein Array von char zugegriffen wird, wie folgt, zum Beispiel:

   extern __shared__ char shared[];
   char data = shared[BaseIndex + tid];