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

Question

Soweit ich weiß, ist geteilter Speicher in Banken aufgeteilt und Zugriffe von mehreren Threads auf ein einzelnes Datenelement innerhalb der gleichen Bank wird einen Konflikt (oder Broadcast) verursachen).

Im Moment zuzuteilen I einen ziemlich großen Array, das mehrere Paare von zwei Matrizen konzeptionell darstellt:

__shared__ float A[34*N] 

Wo N die Anzahl von Paaren sind, und die ersten 16 Hin- und Herbewegungen eines Paares sind eine Matrix und die folgenden 18 Schwimmer sind die zweite.

Die Sache ist, der Zugriff auf die erste Matrix ist konfliktfrei, aber der Zugriff auf die zweite hat Konflikte. Diese Konflikte sind unvermeidlich, aber mein Denken ist, dass aufgrund der zweiten Matrix 18 alle zukünftigen Matrizen zu den Banken falsch ausgerichtet sein werden und deshalb mehr Konflikte als notwendig auftreten werden.

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

Jedes Mal, wenn ich Shared Memory zuteile, beginnt es bei einer neuen Bank? So könnte ich möglicherweise tun

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

Irgendwelche Ideen?

Dank

Answer 1

Wenn Paare von Matrizen zusammenhängend gespeichert werden, und wenn Sie die Elemente linear von Thread-Index zugreifen, dann haben Sie keine Speicherbank Konflikte geteilt.

Mit anderen Worten, wenn Sie:

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 greifen Sie dies mit:

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

Dann benachbarte Fäden benachbarte Elemente in der Matrix und Sie haben keine Konflikte zugreifen.

Als Antwort auf Ihre Kommentare (unten) scheint es, dass Sie in Ihrem Verständnis falsch sind. Es stimmt, dass es 16 Bänke (in aktuellen Generationen, 32 in der nächsten Generation, Fermi) gibt, aber aufeinanderfolgende 32-Bit-Wörter in aufeinanderfolgenden Bänken liegen, d. H. Der Adressraum ist über die Bänke verschachtelt. Das bedeutet, dass Sie keine Bankkonflikte haben werden, vorausgesetzt, Sie haben immer einen Array-Index, der in x + threadIdx.x zerlegt werden kann (x ist nicht abhängig von threadIdx.x oder zumindest über Gruppen von 16 Threads hinweg konstant).

Wenn Sie auf die Matrizen weiter entlang des Arrays zugreifen, greifen Sie immer noch auf sie in einem zusammenhängenden Chunk zu, und Sie haben daher keine Bankkonflikte. Nur wenn Sie auf nicht benachbarte Elemente zugreifen, werden Sie Bankkonflikte haben.

Die Reduktion im Beispiel im SDK veranschaulicht Bankkonflikte sehr gut, indem sie von einer naiven Implementierung zu einer optimierten Implementierung führt, die möglicherweise einen Blick wert ist.

Answer 2

Die Bänke sind so eingerichtet, dass die nächsten 32 Bits in der nächsten Bank liegen. Wenn Sie also ein Array von 4 Byte Floats deklarieren, befindet sich jeder nachfolgende Float im Array in der nächsten Bank (Modulo 16 oder 32, abhängig von Ihrer Architektur). Ich gehe davon aus, dass Sie die Rechenkapazität 1.x haben, also haben Sie eine Bank der Breite 16.

Wenn Sie Arrays von 18 und 16 haben, können die Dinge lustig sein. Sie können, indem er erklärt es wie

__shared__ float sixteen[16][16+1] 

die Bank vermeidet Konflikte Bank Konflikte im 16x16-Array vermeiden, wenn transponieren Elemente Zugriff auf threadIdx.x mit (wie ich Sie tun, wenn Sie davon ausgehen, Konflikte sind immer). Wenn Sie auf Elemente in der ersten Zeile einer 16x16-Matrix zugreifen, befinden sie sich alle in der 1. Bank. Was Sie tun möchten, ist, dass jeder von diesen in einer aufeinander folgenden Bank ist. Padding macht das für dich. Sie behandeln das Array genau wie vorher, als sechzehn [Zeile] [Spalte] oder ähnlich für eine abgeflachte Matrix, als sechzehn [Zeile * (16 + 1) + Spalte], wenn Sie möchten.

Beim 18x18-Fall bewegen Sie sich beim Transponieren gleichmäßig. Die Antwort ist wieder aufzufüllen durch 1.

__shared__ float eighteens[18][18+1] 

So, jetzt, wenn Sie in der transponieren zugreifen (sagen Elemente in der ersten Spalte Zugriff auf), wird sie zugreifen, als (18 + 1)% 16 = 3 und Sie greifen auf die Banken 3, 6, 9, 12, 15, 2, 5, 8 usw. zu, so dass Sie keine Konflikte bekommen sollten.

Die besondere Ausrichtung Verschiebung aufgrund einer Matrix der Größe 18 ist nicht das Problem, weil der Ausgangspunkt des Arrays keinen Unterschied macht, ist es nur die Reihenfolge, in der Sie darauf zugreifen. Wenn Sie die oben vorgeschlagenen Arrays abflachen und sie zu 1 zusammenführen möchten, ist das in Ordnung, solange Sie auf ähnliche Weise darauf zugreifen.