Zahlenfluss

In der Regel lassen sich Recheneinheiten fast aller modernen Computer über die Open Computing Language (OpenCL) der Khronos Group programmieren. Das gilt aber nicht nur für aktuelle Grafikhardware, sondern auch für viele Prozessoren samt den in den meisten enthaltenen SIMD-Vektoreinheiten (Single Instruction Multiple Data). Und selbst die CPUs von Smartphones dürften bald via OpenCL programmierbar sein – jedenfalls wenn sie über eine „Mali T604 GPU“ von ARM verfügen. Andere Hersteller dürften ähnliche Ambitionen haben.

Die in GPUs verfügbare Rechenleistung kann Dimensionen erreichen, die noch vor 20 Jahren einem Supercomputer vorbehalten waren: Eine GeForce GTX 690 von NVIDIA etwa kann mit ihren 3072 CUDA-Cores rund 5,6 Teraflops in einfacher Genauigkeit erreichen – mit nur 300 Watt. Die Karte ist für unter 1000 Euro zu haben, und im Stromverbrauch dürfte sie einem Supercomputer deutlich überlegen sein.

Man sollte meinen, dass sich ob der breiten Verfügbarkeit OpenCL-fähiger Hardware die Softwareentwickler dem Trend zur Nutzung aller parallelen Recheneinheiten angeschlossen hätten. Schließlich ist OpenCL ein Standard, und darin verfasste Programme sind praktisch ohne Änderungen zwischen unterschiedlicher Hardware übertragbar.

Zusatzprozessoren sind im Kommen

Auf aktuelle, meist unter Linux betriebene Supercomputer mag das zutreffen – kaum ein System der Top500-Liste, das einen Einsatz von GPUs nicht zumindest plant. Vertreten sind in der aktuellen Liste 62 solcher Hybridrechner, vor einem Jahr waren es 39. Auf PCs, speziell unter Linux, sind OpenCL-Anwendungen weit weniger verbreitet, obwohl es vielfältige Aufgaben gibt, die speziell von GPUs profitieren könnten. Und wer etwa im Rahmen seiner Ausbildung oder seines Berufs größere Berechnungen plant, kann mit OpenCL oft einen schier unglaubliche Beschleunigungszuwachs erzielen.

Um einen Faktor 100 schnellere Berechnungen gegenüber der CPU sind nach Erfahrungen des Autors mit GPUs und etwas Mühe erreichbar. Und die aktuell mit OpenCL unter Linux (noch) nicht ansprechbare HD 4000 GPU in Intels Ivy-Bridge-Prozessoren kann bei manchen Aufgaben ein Mehrfaches der Geschwindigkeit ihrer Host-CPU erreichen.

Einige Gründe für die mangelnde Verbreitung von OpenCL lassen sich identifizieren: Das Installieren der Treiber und Bibliotheken kann kompliziert sein. OpenCL lehnt sich zwar eng an den C99-Standard an – wer in C programmiert hat, sollte mit OpenCL keinerlei Schwierigkeiten haben –, das Aufsetzen des hostseitigen Codes (meist in C oder C++) ist aber komplex und – obwohl logisch strukturiert – auf den ersten Blick schwer zu durchschauen.

Beim Ausführen ähnelt OpenCL-Code dem Verfahren eines Remote Procedure Call (RPC). Meist muss man vor der Ausführung Daten in den Speicher der Zielplattform transferieren, bevor das Programm dort starten kann. Und nach dem Ende der Berechnung müssen die Ergebnisse wieder auf den Host zurück. Das erhöht die Komplexität des hostseitigen Codes und erschwert das Debuggen sowie das Profiling.

Speziell auf Grafikhardware ist man den Grenzen der Architektur unterworfen, denn die Entwickler haben sie für Programme des Typs „Single Instruction Multiple Data“ (SIMD) entworfen. Code mit vielen Verzweigungen ist deshalb auf GPUs oft wenig performant – da sind herkömmliche CPUs die bessere Wahl. Nicht jeder Algorithmus ist für die Ausführung auf der Grafikhardware geeignet.

OpenCL-Code ist besonders performant, wenn er die Eigenheiten der zugrunde liegenden Hardware mitberücksichtigt. Da sie sich je nach Hersteller und Geräteklassen naturgemäß stark unterscheidet, kann ein für eine Hardware von Hand optimierter Code auf Geräten eines anderen Herstellers selten dessen Potenzial ausreizen.

Treibersuche oft erforderlich

Im ersten Schritt beim Einrichten einer OpenCL-Plattform geht es um das Installieren passender Treiber und Bibliotheken für Grafikhardware. Wer seine Plattform „headless“ betreiben, die GPUs ausschließlich zum Berechnen nutzen will, muss sich oft auf die Suche nach aktuellen Treibern begeben.

OpenCL lässt sich aus einer Reihe von Programmiersprachen ansprechen, darunter Python, C und C++. Als IDE bietet sich unter Linux Eclipse mit der CDT-Erweiterung an (C/C++ Development Tooling). Neuerdings bietet NVIDIA mit Nsight ein Eclipse-Plug-in an, mit dem das Debuggen und Profilen des Codes erheblich leichter von der Hand gehen soll.

Zur Steuerung der Übersetzung lässt sich (wie so oft) CMake einsetzen. Zwar existiert (anders als im Fall von CUDA) von der derzeitigen Version 2.8 kein „offizielles“ OpenCL-Modul, per Google ist aber „FindOpenCL.cmake“ schnell zu orten.

Grundsätzlich muss man bei OpenCL-Programmen zwischen dem Host- und dem Device-seitigen Code unterscheiden. Ersterer ist selbst bei Verwendung von C++ recht komplex, bedarf jedoch bei Wechseln zwischen Einsatzbereichen kaum einer Anpassung.

OpenCL möchte nicht nur das Programmieren von GPUs erleichtern, sondern ganz allgemein Hardware mit parallelen Recheneinheiten ansprechen. Es kommen ebenso gut Rechenkerne der CPU infrage. Intel und AMD bieten entsprechende Anpassungen für ihre CPUs. Außerdem gibt es Varianten etwa für Field Programmable Gate Arrays (FPGAs). Grundsätzlich sollten Implementierungen verschiedener Hersteller nebeneinander laufen können.

Deshalb setzt OpenCL auf das Konzept der „Installable Compiler Driver“. Vereinfacht ausgedrückt liefert jeder Hersteller die Funktionen, die zum Ansprechen seiner jeweiligen Hardware nötig sind, etwa zum Kopieren von Daten oder Starten von Programmen, sowie einen „Just-in-time“-Compiler für den geräteseitigen OpenCL-Code. Beim Code für den Host geht es erst einmal darum, die Zahl der Plattformen eines Systems zu ermitteln, wobei jede OpenCL-Implementierung der des jeweiligen Herstellers entspricht.

Man erzeugt für die gewünschten Plattformen cl::Context-Objekte – sie stellen eine Art Container aller erforderlichen Informationen dar. Im nächsten Schritt benötigt man für jedes Context- ein cl::CommandQueue-Objekt. Es dient der Kommunikation mit den Devices, etwa zum Kopieren von Daten. Das kann übrigens multi-threaded erfolgen, denn bis auf wenige Ausnahmen sind die OpenCL-Objekte und -Funktionen Thread-sicher.

Zu guter Letzt übersetzt man den bereitgestellten OpenCL-Code für die Geräteseite, wobei die Installable Compiler Driver zum Einsatz kommen. Einzelne, als „kernel“ identifizierte Funktionen des Codes kann man mit Argumenten versehen und direkt zur Ausführung an die Queue übergeben. Von dort wandert der Code zum Beispiel auf die GPU und startet parallel auf den verfügbaren und benötigten Recheneinheiten – bei einer GPU können das durchaus einige Tausend sein.

Rechnen auf allen Cores

Als das „Hello World“ der OpenCL-Welt darf die Addition zweier Vektoren gelten. Listing 1 zeigt den benötigten OpenCL-Kernel. Es handelt sich um eine vom Nutzer ansprechbare Einheit des OpenCL-Codes, der auf der benötigten Zahl an Recheneinheiten parallel läuft. Der Kernel erhält drei Argumente: zwei Pointer a und b, deren Inhalte nicht verändert werden dürfen, sowie einen „Ausgabepointer“ result. Der Kernel sollte immer gemeinsam mit dem hostseitigen Code entstehen. Er ist dafür verantwortlich, dass die drei Pointer innerhalb der GPU auf korrekt initialisierte Speicherbereiche verweisen.

Der Kernel fragt seine Position in einem globalen Index-Raum ab und muss sich darauf verlassen können, dass die Position einem gültigen Array-Index entspricht. Das sicherzustellen ist Aufgabe des Programmierers, die er ausschließlich über den hostseitigen Code bewältigen kann.

Funktion des Kernels ist es, „seine“ Werte in a und b zu addieren und an der richtigen Position in result zu speichern. Der Programmierer muss die OpenCL-Runtime passend instruieren, sodass der Kernel genau so oft (gegebenenfalls parallel) läuft, wie Einträge im Array existieren. Bei mehreren Tausend Recheneinheiten und dementsprechend vielen Einträgen im Vektor lässt sich die gesamte Operation parallel und damit schnell ausführen. Der Start der GPU-seitigen Threads erfolgt in Hardware und damit selbst für viele parallele Ausführungsstränge quasi instantan. Der entstehende Overhead ist mit dem einer CPU nicht vergleichbar.

Kernel können zwar erheblich komplexer sein als das gezeigte Beispiel. Im Vergleich mit dem hostseitigen Teil dürfte man aber häufig erheblich einfachere Programmteile antreffen. Es ist oft günstiger, Code mit vielen Verzweigungen in mehrere Kernel zu splitten und der GPU wenig komplizierte Fragmente zu präsentieren. Schleifen und Verzweigungen sind eher eine Sache für die CPU denn für die GPU.

Alle drei Pointer sind mit dem Bezeichner __global versehen. Er gibt den Ort an, an dem die Daten in der GPU gespeichert sind. Sie verfügt über eine Speicherhierarchie, angefangen beim „Global Memory“ (bei modernen Grafikkarten üblicherweise 2 bis 3 GByte GDDR5-Speicher) über mehrere Cache-Level bis zum „Local Memory“ und den Registerdateien. Die teils unterschiedliche Speichertechniken beeinflussen die Zugriffsgeschwindigkeit. Der Transfer von Daten aus dem Host-Speicher in das Global Memory einer diskreten GPU braucht grundsätzlich die meiste Zeit.

Wer häufig mit bestimmten „global“ Daten arbeitet, sollte versuchen, sie in schnellere Speicherbereiche zu kopieren. Das kann durchaus asynchron erfolgen. Latenzen des Speicherzugriffs lassen sich damit teilweise maskieren. Das ist allerdings ein zweischneidiges Schwert: Denn der lokale Speicher ist knapp und ein Überlauf führt zum Rückgriff auf den Global Memory, ohne dass dies für den Programmierer ohne längere Profiling Sessions ersichtlich ist.

Da lokal verfügbarer Speicher je nach GPU-Typen unterschiedlich groß ist, kann das dazu führen, dass derselbe OpenCL-Code auf nominell gleich schnellen GPUs unterschiedlich performant abläuft. Ein weiterer Grund liegt im Vorhandensein spezialisierter Funktionen, etwa kryptografischer Natur. Letztlich bedeutet das, dass man für eine optimale Performance wieder OpenCL-Code für unterschiedliche Devices schreiben muss. Das kann wegen unterschiedlicher Kernel-Argumente oder gar der Aufsplittung eines Kernels in kleinere Fragmente selbst den Code für den Host betreffen.

Fazit

OpenCL-Code ist damit zwar portabel, liefert ohne Nacharbeit aber selten die optimale Performance. Dies sollte man sich vor Augen führen, wenn nach dem Release einer neuen GPU der Zahlenfluss durch erste Benchmarks anschwillt. Sonst kommt zwangsläufig ein Äpfeln-Birnen-Vergleich heraus.

Als wohl größte Hürde vor OpenCL unter Linux darf die mangelnde Integration der Installer in das Betriebssystem gelten. Wenn das Einrichten von OpenCL samt Treibern für die jeweilige Hardware mit den üblichen Hilfsmitteln wie apt-get install einfach von der Hand gehen würde, gäbe es erheblich mehr Programme, die direkt auf die installierten Recheneinheiten zurückgriffen. Windows-Nutzer haben es derzeit einfacher.

Das Aufsetzen des hostseitigen Codes ist im Vergleich einfacher. Es erscheint anfänglich schwierig, folgt aber logischen Regeln und ändert sich zwischen verschiedenen Aufgaben nur wenig. Anspruchsvoller ist das Erstellen performanten Codes – man kommt nicht umhin, sich mit der Architektur der Zielplattform genau auseinanderzusetzen.

Dr. Rüdiger Berlich

ist der Hauptautor der Geneva Bibliothek und forscht am Steinbuch Centre for Computing des Karlsruhe Institute of Technology in den Bereichen Simulationen, Grid und Cloud Computing sowie verteilte parametrische Optimierung.

Alle Links: www.ix.de/ix1301124

 /**
  * @file main.cpp
  */
 /*
  * Copyright (C) Gemfony scientific UG (haftungsbeschraenkt)
  *
  * Contact: contact [at] gemfony (dot) com
  *
  * This file is part of the vectorAdd program.
  *
  * vectorAdd is free software: you can redistribute and/or modify it 
  * under the terms of version 3 of the GNU Affero General Public 
  * License as published by the Free Software Foundation.
  *
  * vectorAdd is distributed in the hope that it will be useful,
  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
  * GNU Affero General Public License for more details.
  *
  * You should have received a copy of the GNU Affero General Public 
  * License along with the Geneva library. If not, see 
  * <http://www.gnu.org/licenses/>.
  *
  * For further information on Gemfony scientific, visit 
  * http://www.gemfony.com
  */

 #include <iostream>
 #include <algorithm>
 

// This will force OpenCL C++ classes to raise exceptions
 // rather than to use an error code
 #define __CL_ENABLE_EXCEPTIONS
 

// OpenCL includes
 #include "LOCAL_CL/cl.hpp"
 

/*********************************************************************/
 /**
  * The OpenCL source
  */
 std::string OCLSource =
 "__kernel void vector_add ("
 "  const __global float *a"
 "  , const __global float *b"
 "  , __global float * result"
 "){"
 "  size_t pos = get_global_id(0);"
 "  result[pos] = a[pos] + b[pos];"
 "}";
 

/*********************************************************************/
 /**
  * Returns the platforms available on this computer
  */
 std::vector<cl::Platform> getPlatforms() {
    std::vector<cl::Platform> platforms;
    // Extract the platforms
    cl::Platform::get(&platforms);
    // Let the audience know
    return platforms;
 }
 

/*********************************************************************/
 /**
  * Emits the devices of a specified device type for a given platform
  */
 std::vector<cl::Device> getDevices(const cl::Platform& platform,  const cl_device_type& t) {
    std::vector<cl::Device> devices;
    // Extract a vector of devices
    platform.getDevices(t, &devices);
    // Let the audience know about the platforms
    for(std::size_t nd=0; nd<devices.size(); nd++) {
      std::cout << "Found device with name " << devices[nd].getInfo<CL_DEVICE_NAME>() << std::endl;
    }
    // Let the audience know
    return devices;
 }
 

/*********************************************************************/
 /**
  * The main function
  */
 const std::size_t NENTRIES = 1000;
 int main(int argc, char** argv)
 {
   // Some OpenCL declarations
   std::vector<cl::Platform> platforms;
   std::vector<cl::Device> devices;
   cl::Context context;
   cl::CommandQueue queue;
   cl::Program::Sources sources;
   cl::Program program;
   cl::Kernel kernel;
   cl::Buffer a_buffer, b_buffer, result_buffer;
 

  // Initialize and fill some input and output buffers
   cl_float *a      = new cl_float[NENTRIES];
   cl_float *b      = new cl_float[NENTRIES];
   cl_float *result = new cl_float[NENTRIES];
   for(std::size_t i=0; i<NENTRIES; i++) {
     a[i]      = (cl_float)i;
     b[i]      = (cl_float)i;
     result[i] = 0.f; // result[i] should be 2*i
   }
   try {
 

    // Retrieve devices
     platforms = getPlatforms();
     std::cout << "Found " << platforms.size() << " platforms" <<  std::endl;
 

    // Use the first platform only
     devices = getDevices(platforms[0], CL_DEVICE_TYPE_GPU);
 

    // Build a context for the devices
     context = cl::Context(devices);
 

    // Build a queue  for the first device -- we will ignore the others
     cl::CommandQueue queue(context, devices[0]);
 

    // Create a program object
     sources = cl::Program::Sources(1, std::make_pair(OCLSource.c_str(),  OCLSource.length()));
     program = cl::Program(context, sources);
 

    // Build the program for our devices. Note that we can pass  compiler options here
    // We could just as well pass a DEFINE into the kernel here
     program.build(devices,"-cl-fast-relaxed-math");
 

    // Create a kernel object
     kernel = cl::Kernel(program, "vector_add");
 

    // Create buffers for the kernel arguments
     a_buffer = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_ PTR, NENTRIES*sizeof(cl_float),  (void *)a);
      b_buffer = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_ PTR, NENTRIES*sizeof(cl_float), (void *)b);
     result_buffer = cl::Buffer(context, CL_MEM_WRITE_ONLY,  NENTRIES*sizeof(cl_float)); // No data is copied to the device
 

    // Set the kernel arguments
     kernel.setArg(0, a_buffer);
     kernel.setArg(1, b_buffer);
     kernel.setArg(2, result_buffer);
 

    // Submit the kernel and wait for its termination
     cl::Event event;
     queue.enqueueNDRangeKernel(kernel, cl::NullRange,  cl::NDRange(NENTRIES), cl::NullRange, NULL, &event); 
     event.wait();
 

    // Extract the result data
     queue.enqueueReadBuffer(result_buffer, CL_TRUE, 0,  NENTRIES*sizeof(cl_float), (void *)result);
 

    // Output a sub-set of the retrieved data
     for(std::size_t i=0; i < std::min<std::size_t>(10,NENTRIES); i++) {
       std::cout << result[i] << " // Should be " << 2*i << std::endl;
     }
   } catch(cl::Error& err) {
     std::cout 
       << err.what() << std::endl
       << err.err() << std::endl;
 

    // If we ran into an error, it was likely due to a build failure
     std::cout << "Build log:\t"
            << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(devices[0]) 
            << std::endl;
   }
 

  // Get rid of the buffers
   delete [] a;
   delete [] b;
   delete [] result;

   return 0;
 } 

(rh)