Rechnen auf der Grafikkarte unter Linux - Seite 4

Parallele Reduktion

Die Grundidee bei der parallelen Reduktion findet sich auch in Sportwettkämpfen wieder. Beim jährlichen Tennisturnier in Wimbledon tritt eine Vielzahl von Spielern gegeneinander an. Im K.o.-Verfahren kommt jener Spieler eine Runde weiter, der das Spiel gewinnt. Die Anzahl der Spieler reduziert sich in jeder Runde um die Hälfte bis zum Finale.

Der folgende Algorithmus macht sich dies zunutze und ist die Basis für eine ganze Klasse ähnlicher Algorithmen. Er zerlegt die Summe in eine Vielzahl von Teilsummen zweier Zahlen und berechnet diese parallel. Der nächste Schritt addiert immer zwei Teilsummen - so lange, bis die Gesamtsumme ermittelt ist.

Abbildung 7 zeigt dieses Berechnungsschema für die Addition von acht Zahlen. Der erste Schritt bildet vier Teilsummen, die vier Threads parallel berechnen. Der nächste Schritt fasst die Ergebnisse in zwei Teilsummen zusammen, die sich wieder parallel berechnen lassen. Im dritten und letzten Schritt gibt es nur noch eine Summe, die von einem der vier Threads berechnet wird.

Abbildung 7: Parallele Reduktion berechnet die Summe einer Zahlenfolge analog einem K.-o.-Ausscheidungsverfahren im Sport. Die Anzahl der rechnenden Threads halbiert sich in jedem Schritt.

Jeder Schritt verkleinert die Anzahl der parallelen Berechnungen. Im Beispiel sind statt sieben sequenzieller Additionen nur drei parallele Reduktionsschritte nötig. Im allgemeinen Fall sind nur dem Logarithmus entsprechend viele Schritte notwendig: Bei der Summe von 1024 Zahlen statt 1023 also nur 10. Das führt zu einer starken Reduzierung der Rechenzeit und ist nicht auf die Addition beschränkt. Weitere Operationen mit dieser günstigen Eigenschaft sind die logischen Verknüpfungen Und sowie Oder und die Minimum- und Maximum-Bestimmung.

Listing 5 zeigt die Kernelfunktion »parallel_reduce()«. Nach jedem parallelen Schritt synchronisiert Zeile 14 die Threads, da die Reduktion im Gleichtakt erfolgt. Die notwendige Barriere ist durch »__syncthreads()« gegeben. Da es nicht möglich ist, über Blockgrenzen hinweg die Threads im Gleichklang zu halten, reduziert jeder Block nur 1024 Zahlen. Anschließend summiert ein sequenzieller Schritt die Teilergebnisse aller Blöcke, aber das ist ein anderes Thema.

Listing 5: »reduce()«
01 // --Kernel-Funktion fuer die Parallele Reduktion 02 03 __global__ void parallel_reduce(int* arr, int radix2_size,int nsteps) { 04 // --Wer bin ich? 05 int index=(blockDim.xblockIdx.x+threadIdx.x)2,s=0; 06 07 for(int i=0;i<nsteps;i++) { 08 s=1<<i; // --s: Schrittweite 09 // --Ueberpruefe, ob Thread aktiv rechnen muss. 10 if (threadIdx.x % s == 0) 11 arr[index]+=arr[index+s]; 12 13 // --Synchronisiere mit allen anderen Threads im Block 14 __syncthreads(); 15 } 16 17 } 18 19 20 // --Hier werden die Gitter- und Blockdimensionen definiert und der Kernel gestartet. 21 22 const int radix2Threads = 9; // --2^.. Threads 23 24 extern "C" void reduce(int* arr, int radix2_size) { 25 26 // --Dimension eines eindimensionalen Blocks 27 dim3 dimBlock(1<<radix2Threads); 28 29 // --Berechnung der Anzahl der Bloecke im eindimensionalen Gitter 30 dim3 dimGrid(1<<(radix2_size-radix2Threads-1)); 31 32 // --Berechnung der Anzahl der Reduktionsschritte pro Block 33 int nreduction=radix2Threads+1; 34 35 // --Starte die Kernelfunktion auf dem Gitter 36 parallel_reduce<<<dimGrid,dimBlock>>>(arr,radix2_size,nreduction); 37 38 }

Listing 5:
»reduce()«

01 // --Kernel-Funktion fuer die Parallele Reduktion
02 
03 __global__ void parallel_reduce(int* arr, int radix2_size,int nsteps) {
04 // --Wer bin ich?
05 int index=(blockDim.x*blockIdx.x+threadIdx.x)*2,s=0;
06 
07 for(int i=0;i&lt;nsteps;i++) {
08     s=1&lt;&lt;i;           // --s: Schrittweite
09     // --Ueberpruefe, ob Thread aktiv rechnen muss.
10     if (threadIdx.x % s == 0)
11         arr[index]+=arr[index+s];
12 
13     // --Synchronisiere mit allen anderen Threads im Block
14     __syncthreads();
15     }
16 
17 }
18 
19 
20 // --Hier werden die Gitter- und Blockdimensionen definiert und der Kernel gestartet.
21 
22 const int radix2Threads = 9;            // --2^.. Threads
23 
24 extern "C" void reduce(int* arr, int radix2_size) {
25 
26 // --Dimension eines eindimensionalen Blocks
27 dim3 dimBlock(1&lt;&lt;radix2Threads);
28 
29 // --Berechnung der Anzahl der Bloecke im eindimensionalen Gitter
30 dim3 dimGrid(1&lt;&lt;(radix2_size-radix2Threads-1));
31 
32 // --Berechnung der Anzahl der Reduktionsschritte pro Block
33 int nreduction=radix2Threads+1;
34 
35 // --Starte die Kernelfunktion auf dem Gitter
36 parallel_reduce&lt;&lt;&lt;dimGrid,dimBlock&gt;&gt;&gt;(arr,radix2_size,nreduction);
37 
38 }

Die reine parallele Reduktion ist nicht besonders effizient. Bereits nach der ersten Reduktion ist die Hälfte der Kerne ohne Arbeit. Daher wird häufig ein effizienteres Verfahren eingesetzt, das eher dem jetzt wieder bei der Fußball-Weltmeisterschaft in Südafrika eingesetzten entspricht, es unterscheidet zwischen einer Vor- und einer Hauptrunde.

Cuda bietet daneben einige (nicht objektorientierte) Konzepte aus der Programmiersprache C++. Dazu gehören Polymorphismus von Funktionen, die Möglichkeit, Default-Parameter anzugeben, die Überladung von Operatoren sowie Namespaces und Funktionstemplates. Das Klassenkonzept unterstützt Cuda nicht.

Abbildung 8: Ist die Installation des Nvidia-Treibers für Linux erfolgreich abgeschlossen, zeigt der Eintrag »Nvidia X Server Settings« unter anderem Informationen zu Grafikkarte, Displaynamen und der Treiberversion an.

Nvidia stellt die Werkzeuge, das SDK und die Dokumentation für alle Interessierten kostenfrei zur Verfügung. Allerdings handelt sich nicht um freie Software im GPL-Sinne; einige Komponenten sind auch nur binär verfügbar. Neben Nvidia ist auch ATI [3] bei der Entwicklung von GPGPU-Systemen sehr aktiv. Der Notwendigkeit einer Standardisierung der Programmierung von Grafikkarten trägt die Khronos-Gruppe [4] Rechnung. Die Organisation, Hüterin mehrerer offener Standards, beispielsweise Open GL, hat inzwischen eine erste Version von Open CL vorgestellt, einer standardisierten Sprache für die Programmierung von GPUs wie auch CPUs.

Das Know-how entscheidet

Der Einstieg in das Rechnen auf der Grafikkarte ist schnell gemacht. Die umfangreiche Dokumentation und die ausführlichen Codebeispiele im SDK sind eine taugliche Starthilfe für Interessierte. Trotzdem erfordert die Programmierung einige Kenntnisse über die eingesetzte Hardware, da ansonsten der Superrechner unterm Schreibtisch seine Leistung nicht ausspielen kann. (mhu)

Installation und Konfiguration
Die Installation der Komponenten erfordert teilweise Handarbeit. Die nachfolgenden Hinweise sind eine Zusammenfassung diverser Foreneinträge, Readme-Dateien und Erfahrungen des Autors. Der Test fand statt mit der Ubuntu-Version 9.10 (Karmic Koala) in 64 Bit, zum einen auf einem Core-i7-System mit einer Geforce-GTX260-Grafikkarte sowie zum anderen auf einem Lenovo-x200s-Notebook, dort aber nur mit der Emulation, da das Gerät keine Nvidia-Grafikkarte enthält. Die Kernelversion ist auf beiden getesteten Systemen 2.6.31-14. Grafiktreiber, Toolkit und SDK Ubuntu bietet im Software-Center einige ältere Versionen von Nvidias Cuda-Treiber zur Installation an. Um die zuletzt erschienene Cuda-Version 2.3 zu installieren, ist allerdings ein Download direkt von Nvidia [5] notwendig. Nach der Auswahl des Betriebssystems (Linux 64 Bit) ist die Distribution (Ubuntu 9.04) zu wählen. Die für diesen Artikel genutzte Version ist 190.18 Beta. Anschließend ist der Treiber zu installieren, allerdings muss dies im Runlevel 3 erfolgen, also ohne laufenden X-Server. Wenn alles glattgegangen ist, steht nach einem erneuten Start des X-Servers der Eintrag »System \| Einstellungen \| Nvidia X Server Settings« zur Verfügung (Abbildung 8). Das Toolkit enthält alle notwendigen Dateien, Bibliotheken und Werkzeuge, um Programme für die Grafikkarte zu entwickeln. Auch hier muss der Admin das Paket von Nvidia direkt laden und anschließend installieren (vorzugsweise unter »/opt«). Die Umgebungsvariable »PATH« ergänzt er dann um »/opt/cuda/bin« sowie »/etc/ld.so.conf« um »/opt/cuda/lib64«. Zusätzlich muss er einmalig »ldconfig« als Root ausführen, sodass der Linker die Cuda-Bibliotheken findet. Alternativ kann er auch die »LD_LIBRARY_PATH«-Variable erweitern. Nun lassen sich Cuda-Programme übersetzen und auszuführen. Hilf- und lehrreich ist es allerdings, auch das separate SDK zu installieren, das eine Vielzahl von Beispielen im Sourcecode bereithält. Nach dem Auspacken des Archivs nach »/opt« oder in ein Homeverzeichnis gilt es, das SDK zu übersetzen. Und hier fangen die Schwierigkeiten an: Ubuntu 9.10 hat standardmäßig GCC 4.4 installiert, damit kommt aber das SDK nicht klar. Zum anderen ist die SDK-Software etwas empfindlich, wenn man andere Pfade verwendet als der Nvidia-Mitarbeiter, der das Paket zusammengebaut hat (es also beispielsweise unter »/opt« installiert). Ubuntu bietet die Möglichkeit, parallel verschiedene Versionen des Compilers zu verwalten. Hierzu sind über die Paketverwaltung die beiden Pakete »gcc-4.3« und »g++-4.3« und Abhängigkeiten nachträglich zu laden und zu installieren. Hinzu kommen noch die Pakete »freeglut3« und »freeglut3-dev«. Dann ist alles beisammen zur Übersetzung der Beispiele und Bibliotheken des SDK. Links zu Bibliotheken Das SDK erwartet, dass die Bibliotheken zum Linken eine bestimmte Namenskonvention erfüllen. Das ist aber für die getestete Ubuntu-Version nicht komplett der Fall. Daher sind symbolische Links notwendig, die als Root zu erzeugen sind: ln -s /usr/lib/libglut.so.3 /usr/lib/Ulibglut.so ln -s /usr/lib/libGLU.so.1 /usr/lib/UlibGLU.so ln -s /usr/lib/libX11.so.6 /usr/lib/UlibX11.so ln -s /usr/lib/libXi.so.6 /usr/lib/UlibXi.so ln -s /usr/lib/libXmu.so.6 /usr/lib/UlibXmu.so Nach der Installation des GCC 4.3 ist das zentrale Makefile »common.mk« unter »/opt/NVIDIA_GPU_Computing_SDK/C/common« entsprechend zu modifizieren. Die folgenden Änderungen erzwingen, dass die erwünschte ältere Compilerversion beim Übersetzen des SDK zum Einsatz kommt: CUDA_INSTALL_PATH ?= /opt/cuda CXX := g++-4.3 CC := gcc-4.3 LINK := g++-4.3 -fPIC NVCCFLAGS := --compiler-bindir=U/usr/bin/gcc-4.3 Nach dem Aufruf von Make in »/opt/NVIDIA_GPU_Computing_SDK/C« und nach ein paar Minuten Übersetzungszeit (nicht von den vielen Compiler-Warnungen irritieren lassen) stehen in »/opt/NVIDIA_GPU_Computing_SDK/C/bin/linux/release« die Programme zur Ausführung bereit. Das Programm »deviceQuery« liefert beispielsweise einige Informationen zur verbauten Nvidia-Grafikkarte. Emulation Ist keine Cuda-fähige Grafikkarte im System vorhanden, dann sollte man die Emulation im Release- oder Debug-Modus erstellen. Das geschieht durch die Eingabe von »make emu=1« und »dbg=1«. Die ganze Cuda-Applikation läuft dann auf dem Host, also der CPU des Linux-PC. Diese Emulation ist übrigens die einzige Möglichkeit, den Cuda-Code im Debugger zu inspizieren. Auf der Grafikkarte selbst ist dies nicht möglich. Für die Cuda-Entwicklung bietet sich Eclipse Galileo mit dem CDT-Plugin an. Der folgende Wrapper hilft gegen einen lästigen GTK+-Bug, der sich unter Ubuntu 9.10 eingeschlichen hat: #!/bin/bash export GDK_NATIVE_WINDOWS=1 /opt/eclipse/eclipse Das im Cuda mitgelieferte Template-Projekt lässt sich als Vorlage für Cuda-Projekte in Eclipse verwenden.

Installation und
Konfiguration

Die Installation der Komponenten erfordert teilweise Handarbeit. Die nachfolgenden Hinweise sind eine Zusammenfassung diverser Foreneinträge, Readme-Dateien und Erfahrungen des Autors. Der Test fand statt mit der Ubuntu-Version 9.10 (Karmic Koala) in 64 Bit, zum einen auf einem Core-i7-System mit einer Geforce-GTX260-Grafikkarte sowie zum anderen auf einem Lenovo-x200s-Notebook, dort aber nur mit der Emulation, da das Gerät keine Nvidia-Grafikkarte enthält. Die Kernelversion ist auf beiden getesteten Systemen 2.6.31-14.

Grafiktreiber, Toolkit und SDK

Ubuntu bietet im Software-Center einige ältere Versionen von Nvidias Cuda-Treiber zur Installation an. Um die zuletzt erschienene Cuda-Version 2.3 zu installieren, ist allerdings ein Download direkt von Nvidia [5] notwendig. Nach der Auswahl des Betriebssystems (Linux 64 Bit) ist die Distribution (Ubuntu 9.04) zu wählen. Die für diesen Artikel genutzte Version ist 190.18 Beta. Anschließend ist der Treiber zu installieren, allerdings muss dies im Runlevel 3 erfolgen, also ohne laufenden X-Server. Wenn alles glattgegangen ist, steht nach einem erneuten Start des X-Servers der Eintrag »System | Einstellungen | Nvidia X Server Settings« zur Verfügung (Abbildung 8).

Das Toolkit enthält alle notwendigen Dateien, Bibliotheken und Werkzeuge, um Programme für die Grafikkarte zu entwickeln. Auch hier muss der Admin das Paket von Nvidia direkt laden und anschließend installieren (vorzugsweise unter »/opt«). Die Umgebungsvariable »PATH« ergänzt er dann um »/opt/cuda/bin« sowie »/etc/ld.so.conf« um »/opt/cuda/lib64«. Zusätzlich muss er einmalig »ldconfig« als Root ausführen, sodass der Linker die Cuda-Bibliotheken findet. Alternativ kann er auch die »LD_LIBRARY_PATH«-Variable erweitern. Nun lassen sich Cuda-Programme übersetzen und auszuführen.

Hilf- und lehrreich ist es allerdings, auch das separate SDK zu installieren, das eine Vielzahl von Beispielen im Sourcecode bereithält. Nach dem Auspacken des Archivs nach »/opt« oder in ein Homeverzeichnis gilt es, das SDK zu übersetzen. Und hier fangen die Schwierigkeiten an: Ubuntu 9.10 hat standardmäßig GCC 4.4 installiert, damit kommt aber das SDK nicht klar. Zum anderen ist die SDK-Software etwas empfindlich, wenn man andere Pfade verwendet als der Nvidia-Mitarbeiter, der das Paket zusammengebaut hat (es also beispielsweise unter »/opt« installiert).

Ubuntu bietet die Möglichkeit, parallel verschiedene Versionen des Compilers zu verwalten. Hierzu sind über die Paketverwaltung die beiden Pakete »gcc-4.3« und »g++-4.3« und Abhängigkeiten nachträglich zu laden und zu installieren. Hinzu kommen noch die Pakete »freeglut3« und »freeglut3-dev«. Dann ist alles beisammen zur Übersetzung der Beispiele und Bibliotheken des SDK.

Links zu Bibliotheken

Das SDK erwartet, dass die Bibliotheken zum Linken eine bestimmte Namenskonvention erfüllen. Das ist aber für die getestete Ubuntu-Version nicht komplett der Fall. Daher sind symbolische Links notwendig, die als Root zu erzeugen sind:

ln -s /usr/lib/libglut.so.3 /usr/lib/Ulibglut.so
ln -s /usr/lib/libGLU.so.1 /usr/lib/UlibGLU.so
ln -s /usr/lib/libX11.so.6 /usr/lib/UlibX11.so
ln -s /usr/lib/libXi.so.6  /usr/lib/UlibXi.so
ln -s /usr/lib/libXmu.so.6  /usr/lib/UlibXmu.so

Nach der Installation des GCC 4.3 ist das zentrale Makefile »common.mk« unter »/opt/NVIDIA_GPU_Computing_SDK/C/common« entsprechend zu modifizieren. Die folgenden Änderungen erzwingen, dass die erwünschte ältere Compilerversion beim Übersetzen des SDK zum Einsatz kommt:

CUDA_INSTALL_PATH ?= /opt/cuda
CXX        := g++-4.3
CC         := gcc-4.3
LINK       := g++-4.3 -fPIC
NVCCFLAGS  := --compiler-bindir=U/usr/bin/gcc-4.3

Nach dem Aufruf von Make in »/opt/NVIDIA_GPU_Computing_SDK/C« und nach ein paar Minuten Übersetzungszeit (nicht von den vielen Compiler-Warnungen irritieren lassen) stehen in »/opt/NVIDIA_GPU_Computing_SDK/C/bin/linux/release« die Programme zur Ausführung bereit. Das Programm »deviceQuery« liefert beispielsweise einige Informationen zur verbauten Nvidia-Grafikkarte.

Emulation

Ist keine Cuda-fähige Grafikkarte im System vorhanden, dann sollte man die Emulation im Release- oder Debug-Modus erstellen. Das geschieht durch die Eingabe von »make emu=1« und »dbg=1«. Die ganze Cuda-Applikation läuft dann auf dem Host, also der CPU des Linux-PC. Diese Emulation ist übrigens die einzige Möglichkeit, den Cuda-Code im Debugger zu inspizieren. Auf der Grafikkarte selbst ist dies nicht möglich.

Für die Cuda-Entwicklung bietet sich Eclipse Galileo mit dem CDT-Plugin an. Der folgende Wrapper hilft gegen einen lästigen GTK+-Bug, der sich unter Ubuntu 9.10 eingeschlichen hat:

#!/bin/bash
export GDK_NATIVE_WINDOWS=1
/opt/eclipse/eclipse

Das im Cuda mitgelieferte Template-Projekt lässt sich als Vorlage für Cuda-Projekte in Eclipse verwenden.

Infos
[1] Cuda-Seite bei Nvidia: [http://www.nvidia.de/object/cuda_home_new_de.html] [2] Sobel-Operator: [http://de.wikipedia.org/wiki/Sobel-Operator] [3] GPGPU bei ATI: [http://developer.amd.com/gpu/ATIStreamSDK/Pages/default.aspx] [4] Khronos-Gruppe: [http://www.khronos.org] [5] Cuda-Download: [http://www.nvidia.com/object/cuda_get.html]

Infos

[1] Cuda-Seite bei Nvidia: [http://www.nvidia.de/object/cuda_home_new_de.html]

[2] Sobel-Operator: [http://de.wikipedia.org/wiki/Sobel-Operator]

[3] GPGPU bei ATI: [http://developer.amd.com/gpu/ATIStreamSDK/Pages/default.aspx]

[4] Khronos-Gruppe: [http://www.khronos.org]

[5] Cuda-Download: [http://www.nvidia.com/object/cuda_get.html]

Der Autor
Michael Uelschen ist Professor für Software-Engineering für technische Systeme an der Fachhochschule Osnabrück. Er befasst sich unter anderem mit Multicore- und GPGPU-Systemen.