OpenCL (AMD GCN) globales Speicherzugriffsmuster für vektorisierte Daten: fortlaufend vs. zusammenhängend

Question

Ich werde die OCL-Kernleistung verbessern und klären, wie Speichertransaktionen funktionieren und welches Speicherzugriffsmuster wirklich besser ist (und warum). Der Kernel wird mit Vektoren von 8 ganzen Zahlen gespeist, die als Array definiert sind: int v [8], das heißt, bevor irgendeine Berechnung durchgeführt wird, muss der gesamte Vektor in GPRs geladen werden. Also, ich glaube, der Flaschenhals dieses Codes ist die anfängliche Datenlast.

Zuerst betrachte ich einige theoretische Grundlagen.

Ziel-HW ist Radeon RX 480/580, das über 256-Bit-GDDR5-Speicherbus verfügt, auf dem die Burst-Lese-/Schreibtransaktion eine Granularität von 8 Wörtern aufweist. Daher liest eine Speichertransaktion 2048 Bit oder 256 Byte. Das, glaube ich, was CL_DEVICE_MEM_BASE_ADDR_ALIGN bezieht sich auf:

Alignment (bits) of base address: 2048. 

So meine erste Frage: Was ist der physikalische Sinn von 128-Byte-Cache-Line? Behält es den Teil der Daten bei, der von einem einzelnen Burst gelesen, aber nicht wirklich angefordert wurde? Was passiert mit dem Rest, wenn wir zum Beispiel 32 oder 64 Bytes angefordert haben - also übersteigt der Rest die Cache-Zeilengröße? (Ich nehme an, es wird nur verworfen - dann, welcher Teil: Kopf, Schwanz ...?)

Nun zurück zu meinem Kernel, ich denke, dass Cache in meinem Fall keine signifikante Rolle spielt, weil ein Burst liest 64 Ganzzahlen -> Eine Speichertransaktion kann theoretisch 8 Arbeitsaufgaben gleichzeitig liefern, es sind keine zusätzlichen Daten zu lesen und der Speicher wird immer zusammengeführt.

Aber dennoch kann ich meine Daten mit zwei unterschiedlichen Zugriffsmustern platzieren:

1) angrenzende

a[i] = v[get_global_id(0) * get_global_size(0) + i]; 

(Weiche als eigentlich perfomed)

*(int8*)a = *(int8*)v; 

2) verschachtelt

a[i] = v[get_global_id(0) + i * get_global_size(0)]; 

Ich erwarte in meinem Fall zusammenhängend wäre schneller, weil wie oben erwähnt eine Speichertransaktion vollständig 8 Arbeitsaufgaben mit Daten stopfen kann. Allerdings weiß ich nicht, wie der Scheduler in der Recheneinheit physikalisch funktioniert: Müssen alle Daten für alle SIMD-Lanes bereit sein oder reicht der erste Teil für 4 parallele SIMD-Elemente? Nichtsdestoweniger nehme ich an, dass es schlau genug ist, zuerst mindestens eine CU mit Daten zu versorgen, sobald die CUs die Befehlsströme unabhängig ausführen können. Während im zweiten Fall müssen wir 8 * global_size/64 Transaktionen durchführen, um einen vollständigen Vektor zu erhalten.

Also, meine zweite Frage: ist meine Annahme richtig?

Nun, die Praxis.

Eigentlich habe ich die gesamte Aufgabe in zwei Kerne aufgeteilt, weil ein Teil weniger Registerdruck hat als ein anderes und daher mehr Arbeitsaufgaben übernehmen kann. Also spielte ich zuerst mit pattern, wie die Daten im Übergang zwischen den Kerneln gespeichert (mit vload8/vstore8 oder Casting zu int8 das gleiche Ergebnis) und das Ergebnis war etwas seltsam: Kernel, der Daten in zusammenhängender Weise liest arbeitet etwa 10% schneller (beide in CodeXL und durch OS Zeitmessung), aber der Kernel, der Daten zusammenhängend speichert, führt überraschend langsamer. Die Gesamtzeit für zwei Kernel ist dann ungefähr gleich. In meinen Gedanken müssen sich beide mindestens genauso verhalten - entweder langsamer oder schneller, aber diese umgekehrten Ergebnisse schienen unerklärlich.

Und meine dritte Frage ist: kann jemand solch ein Ergebnis erklären? Oder kann ich etwas falsch machen? (Oder ganz falsch?)

Answer 1

Schauen Sie sich das Kapitel 2.1 im AMD OpenCL Optimization Guide an. Es konzentriert sich hauptsächlich auf Karten der älteren Generation, aber die GCN-Architektur hat sich nicht vollständig geändert, sollte also immer noch für Ihr Gerät (Polaris) gelten.

Im Allgemeinen verfügen AMD-Karten über mehrere Speichercontroller, auf die in jedem Taktzyklus Speicheranforderungen verteilt werden. Wenn Sie zum Beispiel auf Ihre Werte in der Spaltenhaupt- statt in der Zeilenhauptlogik zugreifen, wird Ihre Leistung schlechter, da die Anforderungen an denselben Speichercontroller gesendet werden. (nach Spalte major) Ich meine, auf eine Spalte Ihrer Matrix wird gemeinsam mit allen Arbeitselementen zugegriffen, die im aktuellen Taktzyklus ausgeführt werden. Dies ist das, was Sie als koalesziertes vs. interleaved bezeichnen. Wenn Sie in einem einzigen Taktzyklus auf eine Reihe von Elementen (dh zusammengewachsene Elemente) zugreifen (dh alle Arbeitselementzugriffswerte in derselben Zeile), sollten diese Anforderungen an andere Speichercontroller als an dieselben verteilt werden.

In Bezug auf Ausrichtung und Cache-Zeilengrößen, frage ich mich, ob dies wirklich hilft, die Leistung zu verbessern. Wenn ich in Ihrer Situation wäre, würde ich versuchen, einen Blick darauf zu werfen, ob ich den Algorithmus selbst optimieren kann oder ob ich oft auf die Werte zugreife und es sinnvoll wäre, sie in den lokalen Speicher zu kopieren. Aber wiederum ist es schwer zu sagen, ohne Wissen darüber, was Ihre Kernel ausführen.

Mit besten Grüßen,

Michael

Answer 2

Nun, nicht beantworten wirklich all meine Frage, aber einige Informationen in Weiten des Internet zu finden Dinge zusammen freiere Art und Weise, zumindest für mich (im Gegensatz zu oben genannten AMD-Optimierung Leitfaden, der unklar und manchmal verwirrend erscheint):

«die Hardware führt einige Koaleszenz, aber es ist kompliziert ... Speicherzugriffe in einem Warp müssen nicht unbedingt zusammenhängend sein, aber Es spielt eine Rolle, wie viele 32-Byte-Speichersegmente (und 128 Byte-1-Cache-Segmente) in sie fallen. der Speichercontroller kann 1, 2 oder 4 dieser 32-Byte-Segmente in einer einzigen Transaktion laden, aber das wird in Cache-Zeilen mit 128 Byte gelesen.
Wenn also jede Spur in einem Warp ein zufälliges Wort in einem 128-Byte-Bereich lädt, gibt es keine Strafe; es ist 1 Transaktion und das Lesen ist bei voller Effizienz. Aber, wenn jede Spur in einem Warp 4 Bytes mit einem Schritt von 128 Bytes lädt, dann ist das sehr schlecht: 4096 Bytes werden geladen, aber nur 128 verwendet, was ~ 3% Effizienz ergibt. »

Also, für meine In diesem Fall spielt es keine Rolle, wie die Daten gelesen/gespeichert werden, während sie immer zusammenhängend sind, aber die Reihenfolge, in der die Teile von Vektoren geladen werden, kann die nachfolgende Befehlsablauf- (Neu-) Planung durch den Compiler beeinflussen.
Ich kann mir auch vorstellen, dass neuere GCN-Architektur zwischengespeicherte/zusammengewachsene Schreibvorgänge ausführen kann, deshalb unterscheiden sich meine Ergebnisse von denen, die von diesem Optimierungshandbuch angefordert werden.