Seite 4: Beschleuniger und Coprozessoren

Inhaltsverzeichnis

Vollgas

Beschleuniger und Coprozessoren bieten das verlockende Angebot von mehr Rechenleistung bei geringerem Energieverbrauch im Verhältnis zu herkömmlichen Prozessoren. Dieser Vorteil wird durch eine signifikant veränderte Mikroarchitektur erkauft, die eine spezielle Programmierung einfordert; außerdem ist der direkt angeschlossene Speicher deutlich kleiner als heute übliche Hauptspeicher. Typische Vertreter dieser Geräteklassen sind Grafikkarten, die in den letzten Jahren immer leistungsfähiger wurden und mittlerweile nahezu frei programmierbar sind, und der Xeon-Phi-Coprocessor, der aus vielen x86-kompatiblen Rechenkernen mit spezieller Vektoreinheit besteht. Beide Gerätetypen haben ihre individuellen Vor- und Nachteile gemeinsam haben sie, dass sie nicht für alle Anwendungsklassen geeignet sind. Wenn aber doch und Anwendungen entsprechend parallelisiert wurden, dann lässt ihre Rechenleistung herkömmliche Prozessoren oftmals "alt" aussehen.

OpenMP 4.0 bringt als die wesentlichste Neuerung die Unterstützung solcher Techniken mit sich. Dabei wurde das Gerätemodell vom Host mit den traditionellen OpenMP-Threads um das Vorhandensein von einem oder mehreren Beschleunigern erweitert, die derzeit alle vom selben Typ sein müssen. Das target-Konstrukt überträgt die Ausführung vom Host zum entsprechenden Gerät (siehe Abb.). Dabei ist es möglich über die map-Klausel Variablen mitsamt ihren Werten mitzunehmen beziehungsweise auf dem Gerät Speicher zu reservieren.

Um hier möglichst viele Anwendungsfälle zu unterstützen, kann man für Arrays in C und C++ Teilbereiche angeben. Beispielhaft überträgt map(to:input[:N])) die Elemente 0 bis N - 1 vom Array input vom Host zum Gerät. Anstelle von to ließe sich auch tofrom verwenden, um den Array-Ausschnitt am Ende des Konstrukts wieder zurück zum Host zu transferieren; nur from würde lediglich eine Rückkopie anfordern.

Da heutige Beschleuniger und Coprozessoren über den PCI Express Bus mit dem Host verbunden sind, ist der Datentransfer eben darüber der Flaschenhals. Eine grundlegende und wichtige Optimierung ist somit, nur die Daten zu transferieren, die wirklich über diese "dünne" Leitung zu schicken sind. Beispielsweise können Variablen, die in einer nachfolgenden target-Region unverändert wieder benötigt werden, auf dem Gerät verbleiben. Zu diesem Zweck existiert das target data-Konstrukt, das Daten zwischen Host und Gerät transferieren kann, aber keinen Übergang des Kontrollflusses zur Folge hat. Das Konstrukt kann somit etwa die Aufrufe mehrerer target-Konstrukte umfassen, die auf das Gerät ausgelagert werden.

Verdeutlicht sei das an einem Beispiel: Die SAXPY-Operation (Single-Precision Real Alpha X Plus Y) y = a * x + y mit den Vektoren y und x und Skalar a ist hervorragend für eine breite Menge von Beschleunigern und Coprozessoren geeignet und wurde auch zur Illustration von Programmierparadigmen für GPUs verwendet. Soll die Operation zweimal hintereinander ausgeführt werden, lassen sich dank des target data-Konstrukts Datentransfers einsparen. Das folgende Codebeispiel verdeutlicht die OpenMP Parallelisierung für Geräte, der zweite SAXPY-Aufruf zeigt dabei die komprimierte Syntax.

int n = 10240; float a = 2.0f; float b = 3.0f;
float *x = (float*) malloc(n * sizeof(float));   // init x
float *y = (float*) malloc(n * sizeof(float));   // init y

#pragma omp target data map(to:x)
{
  int num_blcks = 61;
  int num_thrds = 4;
#pragma omp target map(tofrom:y)
#pragma omp teams num_teams(num_blcks) num_threads(num_thrds)
#pragma omp distribute
  for (int b = 0; b < n; b += num_blcks) {
#pragma omp parallel for
    for (int i = b; i < b + num_blcks; ++i){
      y[i] = a*x[i] + y[i];
    }
  }

// do something with y

#pragma omp target map(tofrom:y)
#pragma omp teams distribute parallel for \
      num_teams(num_blcks) num_threads(num_thrds)
   for (int i = 0; i < n; ++i){
      y[i] = b*x[i] + y[i];
   }

free(x); free(y);

Der gesamte Codebereich für einen Xeon Phi Coprozessor 7120P (60 Rechenkerne mit je vier Hardware-Threads) wird von einem target data-Konstrukt eingefasst, das das Array x auf das Gerät transferiert und dort für beide SAXPY-Aufrufe vorhält. Das target-Konstrukt überträgt den Kontrollfluss auf den Coprozessor und aktualisiert dabei das Array y. Die Parallelisierung der eigentlichen Operation erfolgt dann hierarchisch mit folgenden Konstrukten:

• teams zum Erzeugen einer Liga unabhängiger OpenMP-Teams, deren Master-Thread den Rumpf des Konstrukts ausführt.
• distribute zum Verteilen der Iterationen einer Schleife auf die Liga von Teams, wobei deren Master-Thread die verteilten Iterationen ausführt.
• parallel for zur Verteilung einer Schleife auf alle Arbeits-Threads der jeweiligen Thread-Teams.

Wie das Beispiel zeigt, ist es möglich, einen kombinierten Ausdruck zu verwenden.

Für eine Nvidia-Grafikkarte sind die Werte von num_teams und num_threads entsprechend der Anzahl von Blöcken pro Streaming Multiprocessor und Threads pro Block einzustellen, um eine gute Leistung zu erreichen. Für einen Coprozessor von Intel ist das teams- und distribute-Konstrukt derzeit nicht unbedingt nötig. Allerdings schadet es auch nicht und kann die Leistung des Coprozessors verbessern sowie den Overhead von OpenMP-Konstrukten wie Barrieren verringern.

Der Transfer des Kontrollflusses auf ein Gerät findet immer blockierend statt. Der Host-Thread, der das target-Konstrukt ausführt, bleibt so lange stehen, bis der Kontrollfluss vom Gerät zurückkommt und alle Datentransfers abgeschlossen sind. Das ist hinderlich, wenn Host-Thread und Gerät gleichzeitig etwas Sinnvolles tun sollen. Die Lösung liegt hier in der Verwendung von OpenMP-Tasks. Wird ein target-oder target update-Konstrukt in einem OpenMP-Task eingepackt, lässt sich der OpenMP-Task zur Kommunikation mit dem Gerät von einem anderen OpenMP-Thread ausführen. Der ursprüngliche OpenMP-Thread kann dann mit einem anderen OpenMP-Task eine Berechnung ausführen oder aber auch ein weiteres Gerät mit Arbeit versorgen.

Fazit

Mit 4.0 macht OpenMP einen wahren Sprung nach vorne. Die neu eingeführten Konstrukte vor allem für die Nutzung von Coprozessoren weisen den Weg in die neu – heterogene – Welt. Mit den zusätzlichen Neuerungen und Erweiterungen der bestehenden Features tritt OpenMP 4.0 aus seinem bisherigen Schatten und bietet sich auch für Anwendungen weit über das wissenschaftliche Umfeld hinaus an. Schluss ist noch lange nicht: mit Veröffentlichung des Standards haben die Arbeiten an der Nachfolgeversion bereits begonnen.

Dr. Michael Klemm
ist Mitglied der Intel Software and Services Group und arbeitet in der Developer Relations Division mit Fokus auf Höchstleistungsrechnen. Er arbeitet im OpenMP Language Committee an unterschiedlichsten Fragestellungen und leitet die Projektgruppe zur Entwicklung von Fehlerbehandlungsmechanismen für OpenMP.

Christian Terboven
ist stellvertretender Leiter der Abteilung Hochleistungsrechnen am Rechen- und Kommunikationszentrum der RWTH Aachen. Seit 2006 arbeitet er im OpenMP Language Committee und ist dort Leiter der Gruppe Affinity (Nähe von Threads, Tasks und Daten zu Rechenkernen) und in der Gruppe Tasking (Weiterentwicklung des Tasking-Programmiermodells) aktiv.