Shared-Memory-Bank-Konflikte in CUDA: Wie Speicher ist ausgerichtet auf die Banken

Soweit mein Verständnis geht, shared memory ist unterteilt in Banken und Zugriffe von mehreren threads auf eine einzelne Daten-element innerhalb der gleichen bank wird zu einem Konflikt (oder broadcast).

Ich im moment weisen Sie eine relativ große Auswahl, die konzeptionell stellt mehrere Paare von zwei Matrizen:

__shared__ float A[34*N]

Wo N ist die Anzahl der Paare und die ersten 16 schwimmt ein paar sind eine matrix und die folgenden 18 Schwimmer sind das zweite.

Die Sache ist, Zugang zu dem ersten matrix-Konflikt ist kostenlos, aber der Zugang zu der zweiten hat man Konflikte. Diese Konflikte sind unvermeidlich, aber mein denken ist, dass, da die zweite matrix ist 18 alle zukünftigen Matrizen wird ausgerichtet sein auf die Banken und damit auch mehr Konflikte als nötig auftreten.

Ist das wahr, wenn ja, wie kann ich es vermeiden?

Jedesmal, wenn ich allocate shared memory, fängt er bei einer neuen bank? Also könnte ich tun,

__shared__ Apair1[34]
__shared__ Apair2[34]
...

Irgendwelche Ideen?

Dank

  • Können Sie näher erläutern, wie Sie den Zugriff auf die Elemente?
InformationsquelleAutor zenna | 2010-02-17
  1. 5

    Wenn Ihr Paaren von Matrizen sind zusammenhängend gespeichert, und wenn Sie den Zugriff auf die Elemente Linear mit der thread-index, dann werden Sie nicht haben shared-memory-bank-Konflikte.

    In anderen Worten, wenn Sie haben:

    A[0]  <- mat1 element1
A[1]  <- mat1 element2
A[2]  <- mat1 element3
A[15] <- mat1 element16
A[16] <- mat2 element1
A[17] <- mat2 element2
A[33] <- mat2 element18

    Und Sie dies mit:

    float element;
element = A[pairindex * 34 + matindex * 16 + threadIdx.x];

    Dann angrenzenden Gewinde, der Zugriff auf benachbarte Elemente in der matrix und Sie haben keine Konflikte.

    In der Antwort auf Ihre Kommentare (unten) scheint es, dass Sie Irre in Ihrem Verständnis. Es ist wahr, es sind 16 Banken (in der aktuellen Generation, 32 in der nächsten generation, Fermi), sondern aufeinander folgende 32-bit-Worte befinden sich in aufeinander folgenden Banken, d.h. der Adressraum ist interleaved über die Banken. Dies bedeutet, dass vorausgesetzt, Sie haben immer ein array-index, die zerlegt werden können, um x + threadIdx.x (wo x ist nicht abhängig von threadIdx.x ist oder zumindest Konstanten in Gruppen zu je 16 threads), werden Sie nicht haben, die bank zu Konflikten.

    Beim Zugriff auf die Matrizen weiter entlang der array, das Sie immer noch Zugriff auf Sie in einem zusammenhängenden Stück und daher Sie nicht haben, die bank zu Konflikten. Es ist nur, wenn Sie starten Sie den Zugriff auf nicht-benachbarte Elemente, dass Sie ein bank-Konflikte.

    Den Reduktion Beispiel in der SDK veranschaulicht bank-Konflikte sehr gut durch den Bau von einer naiven Umsetzung, um eine optimale Durchführung, eventuell lohnt sich ein Blick.

    • Danke. Hätte ich nur ein einziges paar von Matrizen (eigentlich ist diese matrix Zeilen wie mache ich eine QR-ZERLEGUNG mit givens Rotationen), dann gäbe es keine oder nur wenige Konflikte. Das problem ist denke ich, dass die folgenden Paare von Matrizen wird nun disaligned auf die shared-memory-Banken. In anderen Worten, Daten, gehört das zweite paar nicht beginnen am Anfang der bank, und damit Konflikte auftreten.
    • Having said that, ich denke, mein Verständnis von Banken war verwirrt. Ich dachte, mehrere 32-bit-Elemente gehörten zu einer einzigen bank, jetzt scheint es jede einzelne 32-bit-element gehört zu seiner eigenen bank. Aber dann verstehe ich nicht, was in der Dokumentation bedeutet, durch "gibt es 16 Banken 16', wie das entspräche einer Gesamtzahl von 64 Byte shared memory an.
    • Aktualisiert meine Antwort Antwort...
    InformationsquelleAutor Tom
  2. 2

    Banken sind so angelegt, dass jeder der aufeinander folgenden 32 bits werden in den nächsten bank. Also, wenn Sie deklarieren ein array von 4-byte schwebt jede weitere schweben in der array wird in der nächsten bank (modulo 16 oder 32, abhängig von Ihrer Architektur). Ich nehme an, du bist auf der compute capability 1.x, so haben Sie eine bank in Breite 16.

    Wenn Sie arrays von 18 und 16, die Dinge können lustig sein. Sie können vermeiden, bank-Konflikte in der 16x16 array deklarieren, wie

    __shared__ float sixteen[16][16+1]

    die bank vermeidet Konflikte beim Zugriff auf die transpose-Elemente mit threadIdx.x (wie ich nehme an, Sie tun, wenn Sie mit Konflikten). Beim Zugriff auf Elemente in, sagen wir, die erste Zeile eines 16x16-matrix, werden Sie alle befinden sich in der 1. bank. Was Sie tun möchten, ist jedes von diesen in einer aufeinanderfolgenden bank. Polsterung tut dies für Sie. Sie behandeln das array genau wie vorher, als sechzehn[Zeile][Spalte], oder auch für eine abgeflachte matrix, als sechzehn[row*(16+1)+Spalte], wenn Sie wollen.

    Für die 18x18 Fall, wenn der Zugriff in das transponieren, verschieben Sie auf eine noch Schrittlänge. Die Antwort ist wieder auf pad 1.

    __shared__ float eighteens[18][18+1]

    So, jetzt, wenn Sie Zugang in die Umsetzung (sagen Sie den Zugriff auf Elemente in der ersten Spalte), wird es von access als (18+1)%16 = 3, und Sie erhalten Zugang auf die Banken 3, 6, 9, 12, 15, 2, 5, 8 etc, man sollte also keine Konflikte.

    Die Besondere Ausrichtung der Verschiebung zu, da Sie eine matrix der Größe 18 ist nicht das problem, denn der Ausgangspunkt des array macht keinen Unterschied, es ist nur die Reihenfolge, in der Sie darauf zugreifen. Wenn Sie möchten, glätten Sie die arrays habe ich oben vorgeschlagen, und führen Sie Sie in 1, das ist in Ordnung, solange Sie auf Sie in ähnlicher Weise.

    InformationsquelleAutor P O'Conbhui
Schreibe einen Kommentar