Unkoalestiger globaler Speicherzugriff, der durch indirekten Zugriff in CUDA verursacht wird

Mein CUDA-Programm leidet unter unkoalesziertem globalen Speicherzugriff. Obwohl der idx-te Thread nur die [idx] -te Zelle in einem Array behandelt, gibt es viele indirekte Speicherzugriffe, wie unten gezeigt.Unkoalestiger globaler Speicherzugriff, der durch indirekten Zugriff in CUDA verursacht wird

int idx=blockDim.x*blockIdx.x+threadIdx.x; 

.... = FF[m_front[m_fside[idx]]];

Für m_fisde [idx], wir haben Zugriffe koalesziert, aber was wir wirklich brauchen, ist FF [m_front [m_fside [idx]]]. Es gibt einen indirekten Zugang auf zwei Ebenen.

Ich habe versucht, einige Muster der Daten in m_front oder m_fsied zu finden, um dies zu einem direkten sequentiellen Zugriff zu machen, aber herausgefunden, dass sie fast "zufällig" sind.

Gibt es eine Möglichkeit, dies anzugehen?

Quelle

2013-02-28 thierry

Dies ist tatsächlich das gleiche Problem wie die Adressierung von Sparse-Matrix, und es hat ziemlich viel Arbeit getan, um zu verstehen, wie man das verbessert. Sie könnten einige Ideen aus der Literatur über Sparse Matrix Operationen auf GPUS bekommen. – talonmies

Wenn es einen Ort in den Zugängen gibt, könnte [diese Frage] (http://stackoverflow.com/questions/12938333/coalesced-global-memory-writes-using-hash/12938726#12938726) von Interesse sein. –

@RobertCrovella ... Der Link ** Texture Mechanism ** in der obigen verknüpften Antwort ist abgelaufen. Können Sie den Link bitte aktualisieren? – sgarizvi

Accelerating globalen Speicherdirektzugriff: Invalidierung der L1-Cache-Zeile

Fermi und Kepler-Architekturen zwei Arten von Lasten aus dem globalen Speicher unterstützen. Vollständige Zwischenspeicherung ist der Standardmodus, es versucht, in L1, dann L2, dann GMEM zu schlagen, und die Ladegranularität ist 128-Byte-Zeile. Nur L2- versucht, in L2, dann GMEM und die Ladegranularität ist 32-Byte. Für bestimmte Direktzugriffsmuster kann die Speichereffizienz erhöht werden, indem L1 ungültig gemacht wird und die geringere Granularität von L2 ausgenutzt wird. Dies kann durch Kompilieren mit –Xptxas –dlcm=cg Option zu nvcc erfolgen.

Allgemeine Richtlinien für den globalen Speicherzugriff beschleunigt: Deaktivieren ECC-Unterstützung

Fermi und Kepler GPUs Unterstützung Error Correcting Code (ECC) und ECC ist standardmäßig aktiviert. ECC reduziert die maximale Speicherbandbreite und wird zur Verbesserung der Datenintegrität in Anwendungen wie der medizinischen Bildverarbeitung und Cluster-Computing im großen Maßstab eingesetzt. Wenn sie nicht benötigt wird, kann sie deaktiviert werden, um die Leistung mit dem Dienstprogramm nvidia-smi unter Linux (siehe link) oder über die Systemsteuerung auf Microsoft Windows-Systemen zu verbessern. Beachten Sie, dass das Aktivieren oder Deaktivieren des ECC einen Neustart erfordert, um wirksam zu werden.

Allgemeine Richtlinien für den globalen Speicherzugriff auf Kepler Beschleunigung: read-only unter Verwendung von Daten-Cache

Kepler eine 48KB-Cache für Daten verfügt, die bekannt ist schreibgeschützt werden für die Dauer der Funktion. Die Verwendung des Nur-Lese-Pfads ist vorteilhaft, da sie den Shared/L1-Cache-Pfad auslagert und uneingeschränkten Speicherzugriff bei voller Geschwindigkeit unterstützt. Die Verwendung des schreibgeschützten Pfads kann automatisch vom Compiler (verwenden Sie das Schlüsselwort const __restrict) oder explizit (verwenden Sie die __ldg() intrinsische) vom Programmierer.

Quelle

2013-03-01 22:16:43 JackOLantern

Vielen Dank. Das Invalidieren der L1-Cache-Zeile ist für mich sinnvoll. Ich werde dies versuchen, um zu sehen, ob die Mem-Effizienz erhöht werden kann. Und ein schreibgeschützter Datencache in Kepler ist auch ein guter Rat. Ich kenne die neue Eigenschaft von Kepler nicht ganz. Lass es mich versuchen. Ich brauche jedoch ECC-Unterstützung. – thierry

Unkoalestiger globaler Speicherzugriff, der durch indirekten Zugriff in CUDA verursacht wird

Antwort

Verwandte Themen