Embedded Grafikprozessoren als Rechenmaschinen GPGPU-Programmierung mit OpenCL

Heutige Embedded Grafikprozessoren sind massiv-parallele Rechenmaschinen, die frei programmierbar sind. In dieser Eigenschaft werden sie General Purpose Graphics Processing Units (GPGPUs) genannt. Programmieren kann man sie mit dem offenen und standardbasierten OpenCL Framework, mit dem man Rechenaufgaben auf CPUs, GPUs und DSPs verteilen kann, um so die Gesamtleistung eines Systems zu optimieren.
Neben dem guten „GFLOPS-pro-Dollar“-Verhältnis bieten GPGPUs eine hohe Performance pro Watt. Hochleistungsplattformen wie die AMD Radeon HD 6970 Grafikkarte liefern beispielsweise 10 GFLOPS pro Watt. Sie eignen sich jedoch aufgrund der Thermal Design Power (TDP) von 250 Watt vorrangig für serverbasierte Applikationen, bei denen man mit dieser Leistungsaufnahme problemlos umgehen kann. Bei vielen Embedded Applikationen muss der Grenzwert für die TDP deutlich geringer liegen. Gründe hierfür sind oft umgebungsbedingte Beschränkungen hinsichtlich Größe, Gewicht und Leistungsaufnahme (Size, Weight and Power = SWaP). Portable Ultraschallgeräte beispielsweise, sollen möglichst klein sein, benötigen aber eine hohe Rechenleistung für die Bildverarbeitung in Echtzeit. Auch Telekommunikations-Infrastrukturen, bei denen die Computing-Anforderung hoch, die Grenzen für die Leistungsaufnahme aber limitiert sind, können von GPGPU-Technologie profitieren. Viele Wehrtechnik- und Luftfahrt-Applikationen, wie beispielsweise Sonar, Radar und Videoüberwachung benötigen hohe Rechenleistung auf Embedded-Formfaktoren. Für diesen Bedarf liefert die AMD Radeon E6760 Embedded GPU mit 480 Stream-Prozessoren 16,5 GFLOPS pro Watt bei einer TDP von nur 35 Watt und damit eine Leistungsaufnahme, die für alle gängigen Embedded Systeme geeignet ist, die auf slotbasierten Rackmoutsystemen basieren wie beispielsweise PICMG 1.x, CompactPCI, VME, VPX, MicroTCA sowie AdvancedTCA.

Effiziente Algorithmen

Um von dem vollen Leistungspotential eingebetteter GPGPU-Anwendungen zu profitieren, ist zudem die Entwicklung grundlegender Algorithmen für spezifische GPU-Architekturen erforderlich. Zu den wichtigsten Algorithmen zählen lineare Gleichungen, Matrizenmultiplikationen, schnelle Fourier-Transformationen, Zufallszahlenerzeugung, elementare Funktionen wie z.B. Subtraktion, Summe, Teilung, Sortieren, etc. sowie domänenspezifische Algorithmen in der Bildverarbeitung . Eine weitere Aufgabe ist die Entwicklung paralleler Algorithmen. Einige Workloads sind beispielsweise nativ parallel. Entsprechende parallele Algorithmen liefern deshalb gegenüber Implementierungen auf Multi-Core CPUs eine viel höhere Leistung. Andere Algorithmen benötigen etwas mehr Aufwand, um sie auf eine massiv-parallele Umgebung zu portieren, aber die Ergebnisse können den Aufwand rechtfertigen. Untersuchungen haben ergeben, dass die Ausführung einer Single Precision General Matrix Multiply Routine (SGEMM) auf einer ATI Radeon HD5870 GPU bis zu 73 Prozent der theoretischen Single-Precision Floating-Point Leistung der GPU erreicht kann. Mit ihrer hohen Rechendichte und parallelen Natur ist SGEMM deshalb gut für die Implementierung auf GPGPU-Plattformen geeignet – insbesondere, wenn es um die Verarbeitung großer Matrizen geht.

Offener Standard

Das Potenzial zur parallelen Programmierung ist schon länger bekannt. Den Pionieren der GPGPU-Technologie standen jedoch noch keine Programmiersprachen zur Verfügung, um die massiv-parallele Rechenleistung der GPUs einfach abzurufen. Stattdessen griffen sie auf Grafikoperationen unter OpenGL zurück, welche vergleichbare mathematische Funktionen nutzten, wie sie zur Berechnungen einer bestimmten Aufgabe benötigt wurden. Anschließend kopierten sie die Ergebnisse dann aus dem Frame Buffer. Proprietäre GPGPU-Programmiersprachen wie CUDA und Brook+ vereinfachten zwar diese Nutzung, was jedoch nicht gelöst ist, ist die Portabilität der Algorithmen. Diese Situation führte zur Entwicklung von OpenCL (Open Computing Language) mit der Entwickler ihre Applikationen auf unterschiedliche Hardware (CPUs, GPUs und DSPs) verteilen können. OpenCL wurde 2008 von einem Industrie-Konsortium geschaffen. Heute sind viele Chip-Hersteller, Software-Firmen und Forschungseinrichtungen Mitglieder des Konsortiums, das sich um die Weiterentwicklung dieses offenen Standards für die parallele Programmierung in heterrogenen Umgebungen kümmert.

OpenCL

OpenCL ermöglicht vor allem eine Effizienzsteigerung bei der Entwicklung von heterogenen Computing-Anwendungen: Mit Hilfe von OpenCL können Applikationsentwickler effizient serielle Aufgaben auf die CPU und parallele Aufgaben auf die GPU verteilen. GPUs sind besonders gut für die parallele Datenverarbeitung geeignet, besonders wenn es um ähnliche Berechnungen großer Datenmengen geht (auch datenparallele Berechnungen genannt). Zur Illustration nutzen wir ein einfaches Beispiel: Eine einfache elementweise Addition von zwei Vektoren a und b, deren Ergebnis in den Vektor c geschrieben wird. Anstatt jeweils ein Paar von Elementen zu addieren, wie es der CPU Code handhabt, kann man OpenCL dazu verwenden, viele Additionen parallel auf der GPU durchzuführen. Das folgende Beispiel zeigt einen typischen Code-Schnipsel, um die Addition auf einer Single-Core CPU durchzuführen, sowie den Code eines sogenannten OpenCL-Kernels, welcher dieselbe Addition auf der GPU durchführt. Programmcode in C: void vec_add (const float *a, const float *b, float *c) { for (int i = 0; i++) c[i] = a[i] + b[i]; } Dieselbe Addition als OpenCL Kernel: __kernel void vec_add (__global const float *a, __global const float *b, __global float *c) { int i = get_global_id(0); c[i] = a[i] + b[i]; } Die Operation für jedes Element i wird Work-Item genannt. Konzeptionell werden alle Work-Items in einem Kernel parallel ausgeführt. Man kann nicht sagen, ob das Work-Item i = x vor, gleichzeitig oder nach dem Work-Item i = y ausgeführt wird. Wir wissen lediglich, dass auf GPUs hunderte und sogar tausende von Work-Items parallel in der Ausführung sein können. OpenCL bietet allerdings einen Weg, Sätze von Work-Items in Arbeitsgruppen zusammenzufassen. In der Regel können Work-Items sich nicht untereinander synchronisieren oder Daten austauschen, aber Work-Items innerhalb der gleichen Arbeitsgruppe können es. Das ermöglicht es, OpenCL-Kernel zu schreiben, die intelligenter sind, als wir in diesem Beispiel darstellen können. Neben dem Code, der auf der GPU ausgeführt wird, muss noch ein Host Programm geschrieben werden, um die GPU zu steuern und zu nutzen. Dieses Host Programm findet und initialisiert die GPU(s), sendet die Daten und den Kernel Code an die GPU, weist die GPU an, die Ausführung zu starten und liest die Ergebnisse aus der GPU aus.

Fazit

Da viele bestehende Algorithmen bereits hervorragend mit der GPGPU Architektur korrelieren und im Vergleich zu traditionellen Multi-Core CPU-Umsetzungen deutliche Leistungszuwächse zeigen, verspricht der Einsatz von OpenCL viele Vorteile bei der Applikationsauslegung. Performancesprünge alleine sind jedoch keine Rechtfertigung für entsprechende Entwicklungsaufwendungen. Auch der Preis muss passen. Da die GPGPU Technologie ansprechende GFLOPS-pro-Watt und -Dollar bietet und zudem neue Möglichkeiten für größen-, gewichts- und leistungssensitive Embedded Applikationen eröffnet, kann der Einsatz von GPGPU wesentlich zur Steigerung Performance und Systemeffizienz und somit letztlich zur Steigerung der Wettbewerbsfähigkeit beitragen. Mit OpenCL steht eine standard-basierte, parallele Entwicklungsumgebung zur Verfügung, mit deren Hilfe Entwickler diese Vorteile effizient in leistungsstarke Applikationen überführen können.