Rationalisierung, was in meinem einfachen OpenCL -Kernel in Bezug auf das globale Gedächtnis vor sich geht

Question

const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

Der obige Kernel ist eine Vektor -Addition, die zehnmal pro Schleife durchgeführt wird. Ich habe den Programmierleitfaden und den Stack -Überlauf verwendet, um herauszufinden, wie globales Speicher funktioniert, aber ich kann immer noch nicht herausfinden, indem ich meinen Code betrachte, wenn ich auf eine gute Weise auf globales Speicher zugriff. Ich greife auf zusammenhängende Weise darauf zu und rate auf eine ausgerichtete Weise. Laden die Karten 128 KB -Stücke globaler Speicher für Arrays A, B und C? Lädt es dann die 128 KB -Stücke für jedes Array einmal für jede 32 verarbeitete GID -Indizes? (4*32 = 128) Es scheint, als würde ich dann keine globale Speicherbandbreite verschwenden, oder?

Übrigens zeigt der Rechenprofiler eine GLD- und GST -Effizienz von 1,00003, was seltsam erscheint. Wie ist es über 1.0?

Answer 1

Ja, Ihr Speicherzugriffsmuster ist ziemlich optimal. Jeder Halbwarp greift auf 16 aufeinanderfolgende 32 -Bit -Wörter auf. Darüber hinaus ist der Zugang 64Byte ausgerichtet, da die Puffer selbst ausgerichtet sind und der Startindex für jeden Halbwarp ein Vielfaches von 16 ist. So erzeugt jeder Halbwarp eine 64Byte -Transaktion. Sie sollten also keine Speicherbandbreite durch unbekotete Zugriffe verschwenden.

Da Sie in Ihrer letzten Frage nach Beispielen gefragt haben, können Sie diesen Code für andere ändern (weniger optimales Zugriffsmuster (da die Schleife nicht wirklich etwas tut, werde ich das ignorieren):

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

Lassen Sie sich zunächst se SE.

Für die Schreibvorgänge erzeugt dies ein etwas unoptimales Muster (nach dem Halbware, die durch ihren ID -Bereich und das entsprechende Zugriffsmuster identifiziert wurden):

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

Grundsätzlich verschwenden wir ungefähr die Hälfte unserer Bandbreite (je weniger als verdoppelte Zugangsbreite für den ungeraden Halbwarps nicht viel hilft, da sie mehr Zugriffe erzeugt, was nicht schneller ist als sozusagen mehr Bytes).

Für die Lesevorgänge von B zugreifen nur die Threads nur Elemente des Arrays. Für jeden Halbwarp liegen alle Zugriffe in einem 128 -byte -ausgerichteten Block (das erste Element befindet Der Index ist ein Vielfaches von 32 für 4 Byte -Elemente, dh der Adressversatz ist ein Vielfaches von 128b). Der AccessPattern erstreckt sich über den gesamten 128B -Block, sodass dies für jeden Halbwarp eine 128B -Übertragung durchführt, die erneut die Hälfte der Bandbreite tailliert.

Die Lesevorgänge von C erzeugen eines der schlimmsten Szenarien, in dem jeder Thread in seinem eigenen 128B -Block indiziert ist. So benötigt jeder Thread seine eigene Übertragung, die eine Hand ein bisschen wie ein Serialisierungsszenario ist (obwohl nicht ganz so schlimm wie normal. da die Hardware in der Lage sein sollte, die Überweisungen zu überlappen). Was noch schlimmer ist, ist die Tatsache, dass dies einen 32B -Block für jeden Thread überträgt und 7/8 der Bandbreite verschwendet (wir greifen auf 4B/Thread, 32b/4b = 8 zu, sodass nur 1/8 der Bandbreite verwendet wird). Da dies die AccessPattern von naiven Matrixtransposes ist, ist es sehr ratsam, diejenigen zu tun, die lokaler Speicher verwenden (sprechen aus Erfahrung).

Berechnen Sie 1.0 (G80)

Hier ist das einzige Muster, das einen guten Zugriff erzeugt, das Original. Alle Muster im Beispiel erzeugen einen vollständig unbekoteten Zugriff und verschwenden 7/8 der Bandbreite (32B Transfer/Thread, siehe oben). Für G80 -Hardware jeden Zugriff, bei dem der n -te Thread in einem Halbwarp nicht auf das n -te Element zugreift

Berechnen Sie 2.0 (Fermi)

Hier erstellt jeder Zugriff auf Speicher 128B -Transaktionen (so viele, wie dies erforderlich ist, um alle Daten zu sammeln, also 16x128b im schlimmsten Fall). Nehmen wir im Moment an, der Cache ist groß genug, um alle Daten zu halten, und es gibt keine Konflikte. Daher wird jede 128B -Cacheline höchstens einmal übertragen. Nehmen wir weiterhin eine serialisierte Ausführung des Halbwarps an, sodass wir einen deterministischen Cache -Beruf haben.

Zugriff auf B überträgt immer noch immer 128B -Blöcke (keine anderen Threadindizes in der korrespondierenden MemoryArea). Der Zugriff auf C erzeugt 128B Transfers pro Thread (schlechtestes AccessPattern möglich).

Für den Zugriff auf einen folgenden (behandeln Sie sie wie Lesevorgänge für den Moment):

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

Für große Arrays verschwendet der Zugriff auf einen theoretisch fast keine Bandbreite. In diesem Beispiel ist die Realität natürlich nicht ganz so gut, da die Zugriffe auf C den Cache ziemlich gut abwerfen werden

Für den Profiler würde ich davon ausgehen, dass die Effizienz von über 1,0 einfach die Ergebnisse von Schwimmpunkten sind.

Ich hoffe, das hilft