Computational Sciences Center

OpenCL

Viele Computer enthalten eine Grafikkarte, die sich auch für allgemeinere Rechenoperationen einsetzen lässt. Die Programmierung der Grafikkarte und den Transfer von Daten zwischen Haupt- und Grafikspeicher können wir über verschiedene Standards bewerkstelligen: CUDA (Compute Unified Device Architecture) ist maßgeschneidert für Grafikkarten des Herstellers NVIDIA, OpenCL (Open Compute Language) ist ein offener Standard, der neben NVIDIA auch von AMD und Intel unterstützt wird, und schließlich bieten aktuelle Versionen des Standards OpenMP (Open Multiprocessing) ebenfalls die Möglichkeit, Rechenaufgaben auszulagern.

Wir konzentrieren uns hier auf OpenCL, da es herstellerunabhängig und relativ einfach zu installieren ist. Allerdings sollte erwähnt werden, dass CUDA durch eine engere Integration in den Compiler erheblich kürzere und "schönere" Programme erlaubt, während bei OpenCL ein Preis für die wesentlich bessere Portabilität zu zahlen ist.

Im einfachsten Fall, also falls die OpenCL-Bibliotheken in einem normalen Bibliothekspfad liegen und passende Header-Dateien installiert sind, genügt ein relativ einfaches Makefile, um die folgenden Beispielprogramme zu übersetzen. Die Bibliothek libOpenCL lädt die herstellerspezifischen Treiber automatisch nach.

Ein knappe Einführung in OpenCL findet sich in Kapitel 8 meines HPC-Skripts.

Systeminformationen

Bevor Arbeit auf Grafikkarten verteilt werden kann, müssen wir zunächst heraus finden, welche Grafikkarten im System vorhanden sind und welche Eigenschaften sie haben. OpenCL kann mit Implementierungen verschiedener Hersteller umgehen, beispielsweise können auf einem System maßgeschneiderte Implementierungen für AMD- und NVIDIA-Grafikkarten parallel installiert sein. Jede solche Implementierung wird in OpenCL durch eine cl_platform_id dargestellt, und mit der Funktion clGetPlatformIDs können wir die verfügbaren Plattformen erfragen. Informationen über eine Plattform erhalten wir mit der Funktion clGetPlatformInfo.

Jede Plattform verwaltet eine Reihe von Geräten, die durch eine cl_device_id dargestellt werden. Mit clGetDeviceIDs können wir die vorhandenen Geräte erfragen und mit clGetDeviceInfo ihre Eigenschaften ermitteln.

cl_info.c

Arithmetische Operationen

Aufgrund seines sehr allgemeinen Programmiermodells erfordert OpenCL relativ viele vorbereitende Schritte, um ein Programm auf einer Grafikkarte auszuführen: Wir müssen eine Plattform auswählen, dann ein von dieser Plattform verwaltetes Gerät, dann einen Kontext, in dem die Kommunikation mit diesem Gerät erfolgen soll. Wenn der Kontext vorliegt, können wir ein OpenCL-Programm definieren und übersetzen. Wir können auch Grafikspeicher anfordern. Die Kommunikation zwischen dem Hauptprozessor und der Grafikkarte erfolgt über eine durch cl_command_queue modellierte Warteschlange, in die wir Arbeitsaufträge einreihen können, beispielsweise Anweisungen für den Transfer von Daten zwischen Haupt- und Grafikspeicher oder Aufrufe des für die Grafikkarte übersetzten Programms.

Für den Datentransfer sind dabei die Funktionen clEnqueueReadBuffer und clEnqueueWriteBuffer zuständig, für die Ausführung von Funktionen ist es clEnqueueNDRangeKernel. Um letztere Funktion aufrufen zu können, müssen wir ein separates Objekt des Typs cl_kernel anlegen, das festlegt, welche Funktion mit welchen Parametern aufgerufen werden soll. Mit der Funktion clFinish können wir schließlich darauf warten, dass alle in einer Warteschlange vorhandenen Aufträge ausgeführt wurden.

Programme für die Grafikkarte werden in der Programmiersprache OpenCL C formuliert, einer Variante der Sprache C, die für den Einsatz auf Grafikkarten angepasst wurde. Beispielsweise gibt es keine Standardbibliothek, aber dafür sind viele der dort vorhandenen Funktionen Bestandteil der Sprache. Darüber hinaus bietet OpenCL C eine Reihe von Erweiterungen, beispielsweise sehr vielfältige Funktionen für Vektoren und Matrizen.

Die Funktion clEnqueueNDRangeKernel erzeugt in der Regel eine Vielzahl von Threads, die alle dieselbe OpenCL-C-Funktion ausführen. Um Arbeit zu verteilen, müssen wir Funktionen wie get_global_id verwenden, um heraus zu finden, welcher Thread gerade ausgeführt wird und welchen Teil der größeren zu bearbeitenden Aufgabe ihm zufällt.

cl_axpy.c

Exponentialfunktion

Das vorige Beispiel bietet einer modernen Grafikkarte nicht die Möglichkeit, die vorhandene Rechenleistung voll auszuspielen, weil im Wesentlichen die Speicherbandbreite die Geschwindigkeit festlegt, da nur sehr wenige Operationen pro aus dem Speicher gelesenen Element ausgeführt werden. Das Rechnen auf Grafikkarten sieht erheblich attraktiver aus, wenn mehr Rechenoperationen pro Speicherzugriff ausgeführt werden, beispielsweise wenn wir die Exponentialfunktion auswerten. Das kann entweder mit einer Taylor-Reihe oder mit der in OpenCL C bereits enthaltenen Exponentialfunktion geschehen.

cl_exp.c

Verschachtelung von Transfers und Rechenoperationen

Falls wir Datenmengen verarbeiten müssen, die nicht in den, häufig auf wenige Gigabyte beschränkten, Grafikspeicher passen, müssen wir sie stückweise dorthin kopieren, verarbeiten, und wieder zurückholen. Damit die Rechenwerke der Grafikkarte während der Kopiervorgänge nicht ungenutzt bleiben, empfiehlt es sich, Transfers und Rechenoperationen zu verschachteln, beispielsweise indem wir mit zwei Warteschlangen arbeiten. Das OpenCL-System kann dann nämlich aus einer Warteschlange einen Speichertransfer bearbeiten und gleichzeitig aus der zweiten einen Rechenvorgang.

cl_exp_readwrite.c

Statt explizit mit clEnqueueWriteBuffer und clEnqueueReadBuffer Daten zu transferieren, können wir auch den Grafikspeicher mit clEnqueueMapBuffer in den Adressraum eines Prozesses einblenden, dort wie ganz gewöhnlichen Speicher manipulieren, und anschließend mit clEnqueuUnmapMemObject signalisieren, dass wir fertig sind. Damit ergibt sich die folgende Variante des obigen Beispiels:

cl_exp_map.c