Ich muss Shared Memory verwenden, der 64 Elemente groß ist, doppelt so viele wie die 32-Bänke. So ist die Anzahl der Speicherzugriffe doppelt so groß wie die Anzahl der Threads in einem Warp. Wie soll ich sie ansprechen, um einen bankkonfliktfreien Zugang zu erhalten?Bankkonfliktfreier Zugriff im gemeinsamen Speicher
Bei 32-Bit-Speicherzugriff können Sie das Standardspeicherzugriffsmuster verwenden.
__shared__ int shared[32];
int data = shared[base + stride * tid];
dort stride ist ungerade.
Wenn Sie 64-Bit-Zugang haben, können Sie einige Trick wie folgt verwenden:
struct type
{
int x, y, z;
};
__shared__ struct type shared[32];
struct type data = shared[base + tid];
Vielen Dank. Als ein Beispiel habe ich ein Array mit 32 Elementen, das mit 16 Threads geteilt wird, so dass jeder Thread auf die zwei Elemente dieses Arrays zugreifen muss. Wie sollte dann die richtige Adressierung gemäß dem oben genannten Problem aussehen? – BehzadX
erstes Muster funktioniert gut in Ihrem Fall. Könnte ich eine kleine Frage stellen: Warum 16? ein beliebiges Array mit dem Qualifikationsmerkmal '__shared__', das für alle Threads im Thread-Block freigegeben ist. – geek
Es ist nur ein Beispiel. Ich beschäftige mich tatsächlich mit FEM-Simulation, die einige Arrays wie Massenmatrizen unter den Berechnungen von Threads teilt. – BehzadX
Nehmen wir an, Sie Compute Capability 1.x verwenden, so dass Ihre Shared-Memory-16 Banken hat, und jeder Thread hat Zugriff auf 2 Elemente im Shared Memory.
Sie möchten, dass ein Thread auf die gleiche Speicherbank für beide Elemente zugreift. Wenn Sie also so organisieren, dass die erforderlichen Elemente 16 voneinander entfernt sind, sollten Sie Bankkonflikte vermeiden.
__shared__ int shared[32];
int data = shared[base + stride * tid];
int data = shared[base + stride * tid + 16];
benutzte ich dieses Muster komplexe Schwimmern zum Speichern, aber ich hatte eine Reihe von komplexen schwebt, so sah es aus wie
#define TILE_WIDTH 16
__shared__ float shared[TILE_WIDTH][2*TILE_WIDTH + 1];
float real = shared[base + stride * tid];
float imag = shared[base + stride * tid + TILE_WIDTH];
Wo die +1 Serialisierung in transponiert Zugriffsmuster zu vermeiden ist.
Wie greifen Sie auf Ihren gemeinsamen Speicher zu? Benötigen Sie einen 64-Bit-Zugriff pro Thread oder zwei 32-Bit-Zugriffe? – geek