Seite 3: Beschleuniger

Inhaltsverzeichnis

Mehr Komfort mit Beschleunigern

In OpenMP 4.0 wurde das Gerätemodell vom Host mit den traditionellen OpenMP-Threads um das Vorhandensein eines oder mehrerer Beschleuniger 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.), im Folgenden häufig Beschleuniger genannt. Der Code, der auf dem Beschleuniger ausgeführt wird, wird auch als Kernel bezeichnet und umfasst in der Regel einen rechenintensiven Bereich der Anwendung. Dabei ist es möglich, über die map-Klausel Variablen mitsamt ihren Werten mitzunehmen beziehungsweise auf dem Gerät Speicher für den Kernel zu reservieren.

OpenMP 4.0 brachte sogenannte Datenregionen, wie sie OpenACC besitzt, in denen Daten auf dem Beschleuniger auch über mehrere Kernel-Aufrufe hinweg erhalten bleiben, und zwar für die gesamte Lebensdauer der Region. Ein Beispiel ist im folgenden Codebeispiel gezeigt, in dem sich die Datenregion von Zeile 3 bis zum Ende erstreckt.

double var1[N];

#pragma omp target data map(tofrom:var1[:N])
{

class C *c = new C();

#pragma omp target
    c.kernel1();   // verwendet var1 und Member von C
    
#pragma omp target
    c.kernel2();   // verwendet var1 und Member von C

delete c;
}

Eine wesentliche Beschränkung dieses Ansatzes ist allerdings, dass sich die Dauer einer Datenregion durch den strukturierten Block ergibt, der dem Pragma nachfolgt. Das erlaubt zwar das Umfassen mehrerer Kernelaufrufe wie hier in einfach strukturiertem Code, ermöglicht es aber nicht, in einer Datenregion weitere Variablen auf den Beschleuniger zu mappen oder aber an weiteren Stellen im Code Vorkehrungen zum Mappen von Daten zu treffen. Ein Beispiel hierfür wäre das Mappen von Daten im Konstruktor der Klasse C.

Dieses oftmals benötigte Feature bringt OpenMP 4.5 in Form der beiden neuen target data-Konstrukte, wie im Folgenden gezeigt. Beide stehen im Code für sich selbst. Das target enter data-Konstrukt führt dabei ein Mapping vom Host zum Gerät durch, target exit data entfernt dieses Mapping entsprechend wieder mit der Option, die Daten zum Host zu kopieren.

/* Variablen in var-list werden zum Beschleuniger gemappt */
#pragma omp target enter data map(map-type: var-list) [clauses]

/* Variablen in var-list werden vom Beschleuniger gemappt */
#pragma omp target exit data map(map-type: var-list) [clauses]

Um nun die Motivation für die neue Funktion noch einmal aufzugreifen, zeigt das nächste Beispiel, wie sich im Konstruktor der Klasse C Daten auf ein Gerät im Rahmen einer bestehenden Datenregion mappen lassen.

class C {

 public:
    C() {
#pragma omp target enter data map(alloc:values[M])
    }

    ~C() {
#pragma omp target exit data map(delete:values[M])
    }
 private:
    double *values;
};

Hierbei wird das Array values auf dem Beschleuniger angelegt. Das target enter data-Konstrukt erzeugt dabei eine sogenannte unstrukturierte Datenregion. Im Destruktor der Klasse C lässt sich entsprechend mit einem target exit data-Konstrukt und dem map-Typ delete das Feld abräumen.

Eine weitere Neuerung im Kontext der Datenverwaltung zwischen Host und Geräten ist die Unterstützung vom Transfer strukturierter Daten. Bisher war es lediglich möglich, einfache skalare Variablen und Felder oder eine Instanz eines strukturierten Datentyps als Ganzes via bitweisem Kopieren zu mappen. OpenMP 4.5 unterstützt nun auch das Mappen einzelner Datentypteile, wobei sich diese Funktion wiederum mit den Mitteln zum Ausdruck von Teilen von Arrays kombinieren lässt. Der folgende Code zeigt ein paar Beispiele der nun vorhandenen Möglichkeiten, womit zu hoffen ist, dass diese den Transfer komplizierter Datenstrukturen bei der Programmierung von Beschleunigern deutlich vereinfachen.

struct A {
   int field;
   double array [N];
} a;

#pragma omp target map(a.field)
#pragma omp target map(a.array[23:42])

Um das Thema Mapping von Daten abzuschließen, sei noch erwähnt, dass nun skalare Variablen in Bezug auf Datenregionen wie "firstprivate" behandelt werden: Sind sie vor einer Datenregion im lokalen Scope deklariert, werden sie genau dann auf das Gerät gemappt, wenn sie in der entsprechende Region referenziert werden. Werden sie nicht referenziert, werden sie auch nicht gemappt – in OpenMP 4.0 bestand der unnötige Zwang für eine Implementierung, alle Variablen zu mappen. Wer das neue Verhalten aber nicht mag, der kann die neue Klausel defaultmap(tofrom:scalar) verwenden, sodass die skalaren Variablen wieder alle wie gehabt gemappt werden.

Eine weitere wesentliche Neuerung im Bereich der Beschleunigerprogrammierung ist das asynchrone Offloading. Bisher war es so, dass ein target-Konstrukt den Offload initiierte, also den Wechsel des Kontrollflusses vom Host hin zum Gerät. Damit war aber der ausführende Thread auf dem Host so lange blockiert, bis der Kernel auf dem Gerät abgeschlossen war.

Im Sinne der Integration in den Rest von OpenMP wurde kein neues Konstrukt eingeführt, um Asynchronität zu ermöglichen, sondern das target-Konstrukt hat ein Upgrade erhalten und ist nun per definitionem ein Task-Konstrukt.

Damit das Verhalten bestehender Programme nicht verändert wird, ist ein target-Konstrukt standardmäßig allerdings ein Task, dessen Ausführung nicht aufgeschoben werden darf, sondern direkt erfolgen muss. Um die Asynchronität einzuschalten, ist die Klausel nowait hinzuzufügen. Das hat zur Folge, dass der Kernel – also der Inhalt des target-Konstrukts – auf dem Beschleuniger ausgeführt wird, während auf dem Host die Ausführung direkt nach dem Ende des target-Konstrukt fortgesetzt wird. Das gilt ebenso für die oben erwähnten enter/exit target data-Konstrukte, was ein asynchrones Mappen zwischen Host und Gerät ermöglicht.

Um das Ganze sinnvoll anwenden zu können, fehlt noch die Möglichkeit, auf den Abschluss einer asynchronen Operation zu warten. Hierzu wurde auf die mit OpenMP 4.0 eingeführten Task-Abhängigkeiten zurückgegriffen. Ein einfaches Beispiel für sowohl asynchrone Datentransfers und Kernelausführung als auch den Ausdruck von Abhängigkeiten ist im Folgenden dargestellt.

double data[N];

#pragma omp target enter data map(to:data[N]) \
            depend(inout:data[0]) nowait

do_something_on_the_host();

#pragma omp target depend(inout:data[0]) nowait
perform_kernel_on_device();

do_something_on_the_host();

#pragma omp target exit data map(from:data[N]) \
            depend(inout:data[0])

Das target enter data-Konstrukt in Zeile 3 initiiert das Mapping von data auf das Gerät, wegen der Angabe der Klausel nowait kann das aber asynchron geschehen, sodass die Ausführung auf dem Host direkt in Zeile 5 mit der Funktion do_something_on_the_host() fortgesetzt wird. Der eigentliche Kernelaufruf erfolgt in den Zeilen 7 und 8: Auch hier wird wieder Asynchronität ermöglicht, aber die Angabe der Abhängigkeit sorgt dafür, dass der Kernel erst gestartet wird, wenn der Datentransfer abgeschlossen ist. Das Mapping des Ergebnisses enthält ebenfalls wieder eine Abhängigkeit, sodass es sich erst ausführen lässt, wenn der Kernel beendet wurde.

Um dieses Feature zu ergänzen, bringt OpenMP 4.5 eine Reihe von API-Funktionen mit, welche die Speicherverwaltung auf dem Gerät vom Host aus erlauben. Beispielsweise alloziert omp_target_alloc() einen Speicherbereich auf dem Gerät. Es gibt zusätzlich einen Zeiger dazu zurück, der sich dann als Argument entsprechende Funktionen mitgeben lässt, die schließlich einen Kernel aufrufen. Mit der omp_target_memcpy()-Funktion können beliebige Datentransfers ausgedrückt werden. Diese API-Funktionen sind insbesondere im Zusammenspiel mit Bibliotheken notwendig geworden.

Die folgende abschließende persönliche Anmerkung sei erlaubt: Die Beschleunigerprogrammierung ist bestimmt nicht leicht, nicht unbedingt für jedermann gedacht, und auch nicht in jeder Anwendung ist Asynchronität oder der Ausdruck von Abhängigkeit notwendig. OpenMP 4.5 bringt diese Werkzeuge aber mit und versucht dadurch besonders attraktiv zu sein, dass – anstelle der Einführung neuer Konstrukte – soweit wie möglich auf vorhandene zurückgegriffen wurde.