18. Mär 2019
Lesedauer 17 Min.
Rechnen mit CUDA
Vollgas mit der GPU, Teil 1
Schnellere Ergebnisse dank einer modernen Grafikkarte.
dotnetpro
Das Rechnen auf einer Grafikkarte (GPU) oder einer speziellen Rechenkarte ist immer noch die „etwas andere Art, zu rechnen“. Es gibt verschiedene Bibliotheken, mit denen sehr schnelle Berechnungen auf einer Grafikkarte durchgeführt werden können.Angefangen hat alles im Jahr 2007 mit den CUDA-Bibliotheken des Grafikkartenherstellers Nvidia. CUDA, kurz für Compute Unified Device Architecture, bietet eine Möglichkeit, mathematische Algorithmen in einen speziellen Code zu übersetzen, der dann mit sehr vielen Threads auf einer Grafikkarte ausgeführt wird.In diesem Artikel soll die Benutzung der CUDA-Bibliotheken ausgehend von normalem C/C++-Code vorgestellt werden. Außerdem werden die Grundlagen der GPU-Programmierung vermittelt und die erforderlichen Werkzeuge beschrieben.In einem Desktop-Computer kann sowohl die CPU (Central Processing Unit) als auch die GPU zum Rechnen genutzt werden. Da in den Rechnern heute Mehrkernprozessoren mit in der Regel vier Prozessoren (Quad Core) eingebaut werden, lassen sich auf der CPU vier Threads starten und vier voneinander unabhängige Aufgaben ausführen. Auf einer GPU können jedoch mehrere Hundert oder mehrere Tausend Threads gestartet und parallel bearbeitet werden, wodurch man – bei geeigneter Aufgabenstellung – eine wesentlich höhere Ausführungsgeschwindigkeit erzielen kann.Vergleicht man die CPU mit einem Sportwagen und die GPU mit einem Lieferwagen, so bietet sich folgendes Beispiel an: Für einen Umzug sind 50 Kisten über 1 000 Kilometer zu transportieren. In den Sportwagen passen jeweils zwei Kisten, die er in fünf Stunden ans Ziel bringt. Der Sportwagen kommt dabei auf eine Leistung 0,4 Kisten/Stunde. Der Lieferwagen schafft dagegen alle 50 Kisten auf einmal, braucht dafür zwar 20 Stunden, kommt mit 2,5 Kisten/Stunde aber auf einen vielfach höheren Durchsatz. Wären dagegen nur vier Kisten zu transportieren, sähe die Rechnung anders aus.Man kann eine Grafikkarte also nicht für alle Berechnungen sinnvoll einsetzen, sondern muss zunächst prüfen, ob viele Datenwerte gleichzeitig und gleichartig verarbeitet werden können. Wenn dies der Fall ist, dann sind Sie bei der GPU sehr gut aufgehoben und können je nach Problem, Datenmenge und Grafikkarte durchaus enorme Performance-Steigerungen bis zum 1 000-Fachen erzielen.
Installation
In diesem Artikel geht es um die CUDA-Bibliotheken von Nvidia. Das bedeutet, dass Sie, um die Codebeispiele nachzuvollziehen, einerseits eine einigermaßen moderne Grafikkarte von Nvidia in Ihrem Rechner benötigen und andererseits das kostenlose CUDA-SDK aus dem Internet herunterladen und installieren müssen.Das Software Development Kit finden Sie unter [1] für die Betriebssysteme Microsoft Windows, Linux und macOS. Mit dem Toolkit werden alle benötigten Werkzeuge, Compiler, Editor-Erweiterungen und Bibliotheken installiert. Denken Sie auch daran, den neuesten Treiber für Ihre Grafikkarte zu installieren, um ein einwandfreies Arbeiten der CUDA-Komponenten sicherzustellen.Sie benötigen außerdem eine Entwicklungsumgebung. Unter Windows (Version 7, 8 oder 10) ist das meist Visual Studio ab Version 2015. Mit Visual Studio (VS) ist es möglich, CUDA-Anwendungen zu erstellen. Eine entsprechende Projektvorlage wird für die vorhandenen VS-Versionen installiert. Unter Linux oder macOS können die für diese Plattformen vorhandenen Entwicklungswerkzeuge benutzt werden.In diesem Artikel wird CUDA 9.1 (64 Bit) mit Microsoft Visual Studio 2015 unter dem Betriebssystem Microsoft Windows 10 (64 Bit) benutzt.Vorgehensweise mit CUDA
Wenn Sie auf der Grafikkarte etwas berechnen möchten, müssen Sie sowohl den Code für den Algorithmus als auch die benötigten Daten auf die GPU transferieren. Und hier ergibt sich das erste Problem in diesem Szenario: Die Übertragung großer Datenmengen kostet nämlich etwas Zeit. Bei den Beispielprogrammen, die mit dem CUDA-SDK installiert werden, ist das kleine Testprogramm bandwidthTest enthalten, welches die Übertragungsgeschwindigkeit für Ihren Rechner ermitteln kann. Diese liegt in der Regel zwischen 3 und 6 Gigabyte pro Sekunde.Die Menge der zu übertragenden Daten kann bei manchen Berechnungen sehr schnell anwachsen. Bei einer Matrixmultiplikation müssen zwei Arrays auf die Grafikkarte kopiert und anschließend das Ergebnis-Array zurück in den CPU-Speicher kopiert werden. Bei einer Matrix-Größe von 500 mal 500 Elementen in doppelter Genauigkeit (8 Byte) müssen insgesamt 6 Megabyte an Daten übertragen werden. Wenn mit diesen Daten nun nur wenige Rechenoperationen ausgeführt werden, benötigt die Datenübertragung mehr Zeit, als durch das parallele Ausführung des Algorithmus auf der GPU gewonnen wird.Das Übertragen des Programmcodes geht normalerweise sehr schnell, da es sich oft nur um einige Kilobyte handelt. Aber auch hier gibt es einen Haken: Bevor das Programm das erste Mal auf der GPU ausgeführt werden kann, muss der Code speziell für die Grafikkarte vorbereitet werden. Auch das kostet wieder Zeit. Dies bedeutet, dass der erste Aufruf eines GPU-Codes etwas länger dauert als der darauf folgende zweite Aufruf. Der allgemeine Ablauf einer Berechnung auf der GPU ist folgender:- Die GPU mit cudaSetDevice auswählen,
- den Speicher auf der GPU mit cudaMalloc bereitstellen,
- Daten mit cudaMemcpy auf die GPU kopieren,
- den Algorithmus auf der GPU ausführen,
- mit cudaDeviceSynchronize auf die GPU warten,
- die Ergebnisse der GPU mit cudaMemcpy zurückkopieren,
- den Speicher der GPU mit cudaFree wieder freigeben,
- die GPU mit cudaDeviceReset zurücksetzen.
Listing 1: Ein einfaches CUDA-Beispiel
#<span class="hljs-keyword">include</span> <span class="hljs-string">"cuda_runtime.h"</span> <br/>#<span class="hljs-keyword">include</span> <span class="hljs-string">"device_launch_parameters.h"</span> <br/><br/>#<span class="hljs-keyword">include</span> &lt;stdio.h&gt; <br/><br/>// Code, der auf der GPU laufen soll <br/>__global__ void addKernel(<br/> <span class="hljs-built_in">int</span> *c, const <span class="hljs-built_in">int</span> *a, const <span class="hljs-built_in">int</span> *b) <br/>{ <br/> <span class="hljs-built_in">int</span> i = threadIdx.x; <br/> c[i] = a[i] + b[i]; <br/>} <br/><br/><span class="hljs-built_in">int</span> main() <br/>{ <br/> // Arrays auf dem Host <br/> const <span class="hljs-built_in">int</span> arraySize = <span class="hljs-number">5</span>; <br/> const <span class="hljs-built_in">int</span> a[arraySize] = { <span class="hljs-number">1</span>, <span class="hljs-number">2</span>, <span class="hljs-number">3</span>, <span class="hljs-number">4</span>, <span class="hljs-number">5</span> }; <br/> const <span class="hljs-built_in">int</span> b[arraySize] = { <span class="hljs-number">10</span>, <span class="hljs-number">20</span>, <span class="hljs-number">30</span>, <span class="hljs-number">40</span>, <span class="hljs-number">50</span> }; <br/> <span class="hljs-built_in">int</span> c[arraySize] = { <span class="hljs-number">0</span> }; <br/><br/> // Zeiger auf Arrays <span class="hljs-keyword">in</span> der GPU <br/> <span class="hljs-built_in">int</span> *dev_a = <span class="hljs-number">0</span>; <br/> <span class="hljs-built_in">int</span> *dev_b = <span class="hljs-number">0</span>; <br/> <span class="hljs-built_in">int</span> *dev_c = <span class="hljs-number">0</span>; <br/><br/> // GPU auswählen <br/> cudaSetDevice(<span class="hljs-number">0</span>); <br/><br/> // Speicher <span class="hljs-keyword">in</span> der GPU anlegen <br/> cudaMalloc((void**)&amp;dev_c, arraySize * sizeof(<span class="hljs-built_in">int</span>)); <br/> cudaMalloc((void**)&amp;dev_a, arraySize * sizeof(<span class="hljs-built_in">int</span>)); <br/> cudaMalloc((void**)&amp;dev_b, arraySize * sizeof(<span class="hljs-built_in">int</span>)); <br/> // Arrays a und b <span class="hljs-keyword">in</span> die GPU kopieren <br/> cudaMemcpy(dev_a, a, arraySize * sizeof(<span class="hljs-built_in">int</span>), <br/> cudaMemcpyHostToDevice); <br/> cudaMemcpy(dev_b, b, arraySize * sizeof(<span class="hljs-built_in">int</span>), <br/> cudaMemcpyHostToDevice); <br/><br/> // Code auf der GPU ausführen <br/> addKernel &lt;&lt;&lt;<span class="hljs-number">1</span>, arraySize &gt;&gt;&gt;(dev_c, dev_a, dev_b); <br/><br/> // Warten auf die GPU <br/> cudaDeviceSynchronize(); <br/><br/> // Ergebnis-Array auf den Host zurückkopieren <br/> cudaMemcpy(c, dev_c, arraySize * sizeof(<span class="hljs-built_in">int</span>), <br/> cudaMemcpyDeviceToHost); <br/><br/> // Speicher auf der GPU freigeben <br/> cudaFree(dev_c); <br/> cudaFree(dev_a); <br/> cudaFree(dev_b); <br/> <br/> // Ausgabe der Ergebnisse <br/> for (<span class="hljs-built_in">int</span> i = <span class="hljs-number">0</span>; i &lt; <span class="hljs-number">5</span>; i++) <br/> { <br/> printf(<span class="hljs-string">"%d + %d = %d\n"</span>, a[i], b[i], c[i]); <br/> } <br/><br/> // GPU zurücksetzen <br/> cudaDeviceReset(); <br/><br/> <span class="hljs-keyword">return</span> <span class="hljs-number">0</span>; <br/>}
Es sei noch einmal darauf hingewiesen, dass in Listing 1 wegen der Übersichtlichkeit der gesamte Fehlerprüfcode weggelassen wurde. Wenn Sie das CUDA-Toolkit frisch auf Ihrem Rechner installiert haben, ist es sicherlich ratsam, zunächst einmal zwei oder drei Beispielprogramme aus dem Toolkit auszuführen, um die Korrektheit und Vollständigkeit der CUDA-Installation zu überprüfen.Im Beispielcode werden zunächst die wichtigen CUDA-Header-Dateien eingefügt. Danach wird die Funktion addKernel definiert, die den Code enthält, der auf der GPU ausgeführt werden soll. Diese Funktion muss mit dem Prefix __global__ eingeleitet werden, damit der Compiler diesen Code für die Ausführung auf der GPU vorbereitet. Der Funktionscode ist eigentlich sehr einfach, trotzdem gibt er uns ein Rätsel auf: Woher kommt die Variable threadIdx.x in der ersten Zeile der Funktion?Es handelt sich dabei um ein spezielle CUDA-Variable, die in der Header-Datei device_launch_parameter.h deklariert wird. Die Variable threadIdx.x stellt dem Entwickler eine eindeutige Nummer eines Threads zur Verfügung, in welchem der Funktionscode ausgeführt wird.Auf das Beispiel bezogen entsteht dann folgende Situation: Die Arrays a, b und c enthalten jeweils fünf Elemente. Die CUDA-Bibliothek erzeugt nun fünf Threads. Für den ersten Thread wird die Variable threadIdx.x auf den Wert 0 gesetzt, im zweiten Thread auf den Wert 1, und so weiter. Der fünfte Thread enthält in der Thread-spezifischen Variablen threadIdx.x den Wert 4. Da eine Grafikkarte fünf Threads problemlos gleichzeitig erzeugen und auch ausführen kann, erfolgt die Berechnung der fünf Ergebniswerte parallel – also gleichzeitig – in fünf GPU-Threads. In der zweiten Zeile der Funktion addKernel wird dieser Thread-Index benutzt, um die Elemente in den Arrays a, b und c anzusprechen.Die Datenübergabe an die Funktion addKernel erfolgt über ganz normale C-Zeiger. Es muss nun aber noch geklärt werden, wo diese Zeiger herkommen und wie die CUDA-Bibliothek erkennt, dass für das gezeigte Programm eben exakt fünf Threads benötigt werden, um das Ergebnis zu ermitteln. Dazu soll die main-Funktion des Beispiels etwas genauer betrachtet werden.Zuerst werden die drei erforderlichen C-Arrays deklariert und initialisiert. Dabei ist es wichtig zu verstehen, dass diese Arrays im normalen Arbeitsspeicher der CPU angelegt werden. Die Arrays a und b werden nicht verändert und können darum mit dem Attribut const versehen werden.Danach werden drei Integer-Zeiger angelegt und zunächst mit 0 initialisiert. Diese Zeiger werden in Kürze auf die Speicherbereiche zeigen, die für die Daten im Speicher der GPU angelegt werden.Als Nächstes wird eine CUDA-fähige Grafikkarte mit dem Befehl cudaSetDevice ausgewählt. In den meisten Fällen erübrigt sich diese Zeile, da es im Rechner oft nur eine solche Grafikkarte gibt. Diese wird von den CUDA-Bibliotheken dann automatisch benutzt.Jetzt werden die erforderlichen Speicherbereiche in der GPU mit cudaMalloc angelegt. Als Parameter wird die Adressvariable der jeweiligen Zeiger angegeben, die eben deklariert wurden. Zudem muss die gewünschte Größe des Speicherblocks übergeben werden. Die Funktion cudaMalloc legt den Speicherbereich im GPU-Speicher an und gibt die Position im Zeiger zurück. Allerdings ist hier etwas Vorsicht angesagt, denn in Listing 1 wird noch nicht überprüft, ob das Anlegen der Speicherblöcke erfolgreich war oder nicht. Der dazu erforderliche Code wird gleich im nächsten Beispiel vorgestellt.Im folgenden Schritt werden die Daten, die sich ja zunächst im CPU-Speicher in den Arrays a und b befinden, mithilfe der Funktion cudaMemcpy in den GPU-Speicher kopiert. Als Parameter werden Ziel- und Quellzeiger des Arrays sowie die Größe des Speicherblocks und die Kopierrichtung angegeben. Die Kopierrichtung kann mithilfe der Konstanten cudaMemcpyHostToDevice und cudaMemcpyDeviceToHost angegeben werden.Das Array c liegt zwar ebenfalls im Arbeitsspeicher der CPU und wurde mit 0 initialisiert, es wird aber nicht in die GPU kopiert, denn dieser Kopiervorgang wäre überflüssig und würde nur zusätzliche Zeit beanspruchen.Nun kommt im C-Quellcode wieder eine sehr sonderbare Zeile. Hier wird nun die Funktion addKernel aufgerufen und auf der GPU ausgeführt. Diese Kernel-Funktionen haben für CUDA die folgende Aufrufkonvention:
funcName <<<a, b<span class="hljs-meta">>>></span>(params,...);
Der Aufruf beginnt mit dem Namen der Funktion (hier heißt sie addKernel), die bei der Implementierung mit dem Prefix __global__ gekennzeichnet wurde. Der Compiler-Treiber
nvcc von Nvidia kann diesen Aufruf auflösen. In den dreifachen spitzen Klammern können Compiler-Konstanten oder Variablen angegeben werden, die im Prinzip die Anzahl der Thread-Blöcke (im Beispiel: 1) und die Anzahl der Threads in einem Block (im Beispiel: arraySize) definieren. Diese beiden Angaben werden im Folgenden noch genauer erläutert. Am Ende des Statements übergeben Sie die Parameter für den Funktionsaufruf. Im gezeigten Beispiel sind das die Zeiger auf die drei Arrays, deren Speicher in der GPU liegt.Da der Aufruf der Kernel-Funktion addKernel asynchron ausgeführt wird, muss nun auf die Beendigung durch den blockierenden Aufruf von cudaDeviceSynchronize gewartet werden. Dann können Sie die Ergebnisdaten mit cudaMemcpy aus dem GPU-Speicher zurück in den Arbeitsspeicher der CPU kopieren. Danach können Sie auf den Inhalt des Arrays c zugreifen, das nun die Ergebnisse enthält.Abschließend ist noch eine sehr wichtige Aufgabe zu erledigen: Die auf der GPU allokierten Speicherblöcke müssen mit cudaFree wieder frei gegeben werden.Im letzten Teil des Beispielprogramms werden die Ergebnisse ausgegeben und schließlich mit der Funktion cudaDeviceReset alle Ressourcen aufgelöst, die in der GPU für diesen Prozess angelegt wurden.Die Datei mit dem besprochenen C-Code bekommt (von Visual Studio) den Namen kernel.cu. Das Projekt kann nun mit den Standardbefehlen von Visual Studio übersetzt und ausgeführt werden. Alternativ dazu können Sie das Programm mit dem Befehl nvcc auch aus der Kommandozeile heraus übersetzen.Wie schon mehrfach erwähnt, wurde im Listing jeglicher Code für die Fehlerprüfung weggelassen. Jede CUDA-Funktion liefert beim Aufruf einen Fehlerstatus in Form einer Variablen vom Typ cudaError_t zurück. Hat alles geklappt, enthält diese Variable den Wert cudaSuccess, ansonsten liefert sie einen Fehlercode.Ein CUDA-Aufruf wird dann im Prinzip immer mit dem Fehlerbehandlungscode aus Listing 2 erweitert. Beim Starten der Kernel-Funktion auf der GPU kann der Fehlerstatus mit der Funktion cudaGetLastError abgefragt werden. Die zu einem Fehlerstatus gehörende Meldung ermitteln Sie durch einen Aufruf von cudaGetErrorString. Tritt ein Fehler auf, müssen Sie den auf der GPU allokierten Speicher unbedingt wieder freigeben. Dies wird im Beispiel durch einen Sprung auf das Label Error gewährleistet. Hier sind natürlich auch andere Vorgehensweisen denkbar. Den grundsätzlichen Aufbau des Codes zeigt Listing 2.
nvcc von Nvidia kann diesen Aufruf auflösen. In den dreifachen spitzen Klammern können Compiler-Konstanten oder Variablen angegeben werden, die im Prinzip die Anzahl der Thread-Blöcke (im Beispiel: 1) und die Anzahl der Threads in einem Block (im Beispiel: arraySize) definieren. Diese beiden Angaben werden im Folgenden noch genauer erläutert. Am Ende des Statements übergeben Sie die Parameter für den Funktionsaufruf. Im gezeigten Beispiel sind das die Zeiger auf die drei Arrays, deren Speicher in der GPU liegt.Da der Aufruf der Kernel-Funktion addKernel asynchron ausgeführt wird, muss nun auf die Beendigung durch den blockierenden Aufruf von cudaDeviceSynchronize gewartet werden. Dann können Sie die Ergebnisdaten mit cudaMemcpy aus dem GPU-Speicher zurück in den Arbeitsspeicher der CPU kopieren. Danach können Sie auf den Inhalt des Arrays c zugreifen, das nun die Ergebnisse enthält.Abschließend ist noch eine sehr wichtige Aufgabe zu erledigen: Die auf der GPU allokierten Speicherblöcke müssen mit cudaFree wieder frei gegeben werden.Im letzten Teil des Beispielprogramms werden die Ergebnisse ausgegeben und schließlich mit der Funktion cudaDeviceReset alle Ressourcen aufgelöst, die in der GPU für diesen Prozess angelegt wurden.Die Datei mit dem besprochenen C-Code bekommt (von Visual Studio) den Namen kernel.cu. Das Projekt kann nun mit den Standardbefehlen von Visual Studio übersetzt und ausgeführt werden. Alternativ dazu können Sie das Programm mit dem Befehl nvcc auch aus der Kommandozeile heraus übersetzen.Wie schon mehrfach erwähnt, wurde im Listing jeglicher Code für die Fehlerprüfung weggelassen. Jede CUDA-Funktion liefert beim Aufruf einen Fehlerstatus in Form einer Variablen vom Typ cudaError_t zurück. Hat alles geklappt, enthält diese Variable den Wert cudaSuccess, ansonsten liefert sie einen Fehlercode.Ein CUDA-Aufruf wird dann im Prinzip immer mit dem Fehlerbehandlungscode aus Listing 2 erweitert. Beim Starten der Kernel-Funktion auf der GPU kann der Fehlerstatus mit der Funktion cudaGetLastError abgefragt werden. Die zu einem Fehlerstatus gehörende Meldung ermitteln Sie durch einen Aufruf von cudaGetErrorString. Tritt ein Fehler auf, müssen Sie den auf der GPU allokierten Speicher unbedingt wieder freigeben. Dies wird im Beispiel durch einen Sprung auf das Label Error gewährleistet. Hier sind natürlich auch andere Vorgehensweisen denkbar. Den grundsätzlichen Aufbau des Codes zeigt Listing 2.
Listing 2: Fehlerbehandlung
<span class="hljs-regexp">//</span> ... <br/><br/>cudaError_t cudaStatus; <br/><br/><span class="hljs-regexp">//</span> ... <br/><br/>cudaStatus = <br/> cudaMalloc((void**)&amp;dev_c, size * sizeof(int)); <br/><br/><span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) <br/>{ <br/> fprintf(stderr, <span class="hljs-string">"cudaMalloc failed!"</span>); <br/> goto Error; <br/>} <br/><br/><span class="hljs-regexp">//</span> ... <br/><span class="hljs-regexp">//</span> Code auf GPU ausführen <br/><br/>addKernel&lt;&lt;&lt;<span class="hljs-number">1</span>, arraySize&gt;&gt;&gt;(dev_c, dev_a, dev_b); <br/>cudaStatus = cudaGetLastError(); <br/><br/><span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) <br/>{ <br/> fprintf(stderr, <span class="hljs-string">"addKernel launch failed: %s\n"</span>,<br/> cudaGetErrorString(cudaStatus)); <br/><br/> goto Error; <br/>} <br/><br/><span class="hljs-regexp">//</span> ... <br/><br/>Error: <br/> cudaFree(dev_c); <br/> cudaFree(dev_a); <br/> cudaFree(dev_b); <br/><br/><span class="hljs-regexp">//</span> ...
Den kompletten Code des Beispiels inklusive Fehlerbehandlung können Sie sich ansehen, wenn Sie mit Visual Studio ein CUDA-Projekt erstellen und sich die Datei kernel.cu ansehen. Der CUDA-spezifische Code ist dort in der separaten Kernel-Funktion addWithCuda implementiert. Dort werden der Speicher auf der GPU angelegt, die Daten kopiert, der Kernel aufgerufen und die Ergebnisse zurück in den Arbeitsspeicher transferiert. Nach jedem CUDA-Aufruf folgt eine Prüfung, ob die Aktion erfolgreich war. Wird ein Fehler festgestellt, werden alle angelegten Speicherblöcke freigegeben und das Programm beendet.
Was kann meine GPU?
Bevor es um die Organisation der Threads auf der GPU geht, soll noch gezeigt werden, wie Sie die Möglichkeiten einer bestimmten Grafikkarte ermitteln. Die CUDA-Bibliothek bietet dazu die Funktion cudaGetDeviceProperties an. Deren Aufruf erfolgt mit einem Zeiger auf die Datenstruktur cudaDeviceProp und die Nummer der jeweiligen GPU. Ist im Rechner nur eine Grafikkarte für CUDA-Anwendungen und Grafikausgabe vorhanden, so wird 0 als Gerätenummer verwendet. Listing 3 demonstriert den Aufruf der Funktion cudaGetDeviceProperties.Listing 3: Ausgabe einiger GPU-Daten
#include &lt;iostream&gt; <br/>#include &lt;cuda_runtime.h&gt; <br/><br/>int main() <br/>{ <br/> int deviceCount = <span class="hljs-number">0</span>; <br/> cudaError_t error_id = <br/> cudaGetDeviceCount(&amp;deviceCount); <br/><br/> <span class="hljs-keyword">if</span> (error_id != cudaSuccess) <br/> { <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">"Fehler Nr. %d\n-&gt; %s\n"</span>, <br/> (int)error_id, cudaGetErrorString(error_id)); <br/> exit(EXIT_FAILURE); <br/> } <br/><br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">"\nEs wurde(n) %d CUDA-Grafikkarte(n) </span><br/><span class="hljs-string"> gefunden.\n"</span>, deviceCount); <br/><br/> <span class="hljs-keyword">if</span>(deviceCount == <span class="hljs-number">0</span>) <br/> { <br/> <span class="hljs-built_in">return</span> <span class="hljs-number">0</span>; <br/> } <br/><br/> int driverVersion = <span class="hljs-number">0</span>; <br/> int runtimeVersion = <span class="hljs-number">0</span>; <br/><br/> // Schleife über alle gefundenen CUDA-Devices <br/> <span class="hljs-keyword">for</span> (int dev = <span class="hljs-number">0</span>; dev &lt; deviceCount; ++dev) <br/> { <br/> cudaSetDevice(dev); <br/> cudaDeviceProp deviceProp; <br/> cudaGetDeviceProperties(&amp;deviceProp, dev); <br/><br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">"\nDevice %d: \"%s\"\n\n"</span>, dev, <br/> deviceProp.name); <br/><br/> cudaDriverGetVersion(&amp;driverVersion); <br/> cudaRuntimeGetVersion(&amp;runtimeVersion); <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" CUDA Driver Version / Runtime Version: </span><br/><span class="hljs-string"> %d.%d / %d.%d\n"</span>, driverVersion / <span class="hljs-number">1000</span>, (driver<br/> Version <span class="hljs-symbol">%</span> <span class="hljs-number">100</span>) / <span class="hljs-number">10</span>, runtimeVersion / <span class="hljs-number">1000</span>, <br/> (runtimeVersion <span class="hljs-symbol">%</span> <span class="hljs-number">100</span>) / <span class="hljs-number">10</span>); <br/><br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" CUDA Capabilities: </span><br/><span class="hljs-string"> %d.%d\n"</span>, deviceProp.major, deviceProp.<span class="hljs-built_in">minor</span>); <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" Groesse des globalen Speichers: </span><br/><span class="hljs-string"> %llu Bytes\n"</span>, deviceProp.totalGlobalMem); <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" Anzahl der Multiprozessoren:</span><br/><span class="hljs-string"> %3d\n"</span>, deviceProp.multiProcessorCount); <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" GPU Taktfrequenz: </span><br/><span class="hljs-string"> %.0f MHz\n"</span>, deviceProp.clockRate * 1e-3f); <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" Max. Anzahl der Threads / MP:</span><br/><span class="hljs-string"> %d\n"</span>, deviceProp.maxThreadsPerMultiProcessor); <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" Max. Anzahl der Threads / Block:</span><br/><span class="hljs-string"> %d\n"</span>, deviceProp.maxThreadsPerBlock); <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" Groesse des Konstanten-Speichers:</span><br/><span class="hljs-string"> %lu Bytes\n"</span>, deviceProp.totalConstMem); <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">" Groesse des Shared Speichers / Block: </span><br/><span class="hljs-string"> %lu Bytes\n\n"</span>, deviceProp.sharedMemPerBlock); <br/> } <br/><br/> <span class="hljs-built_in">return</span> <span class="hljs-number">0</span>; <br/>}
Außerdem werden die Version des aktuell installierten Treibers und die Version der vorhandenen CUDA-Runtime ermittelt und ausgegeben. Im Anschluss daran listet die Funktion verschiedene Datenwerte aus der Struktur cudaDeviceProp auf dem Bildschirm. Im Beispiel handelt es sich nur um die allerwichtigsten Werte.Nach den Versionsnummern werden die sogenannten Compute Capabilities [2] ausgegeben. Dieser Wert gibt an, welche Möglichkeiten die GPU anbietet. Moderne Grafikkarten haben hier eine Versionsnummer größer als 5.x. Ältere Grafikkarten können natürlich auch noch mit CUDA benutzt werden. In diesen Fällen stehen aber die jüngeren Erweiterungen der CUDA-Bibliotheken nicht zur Verfügung.Die für Ihre Anwendung erforderlichen Compute Capabilities können Sie in den Eigenschaften des Visual-Studio-Projekts angeben und bei der Übersetzung entsprechend berücksichtigen.Danach werden die Größe des GPU-Speichers und die Anzahl der auf der Grafikkarte vorhandenen Multiprozessoren ausgegeben. Jeder der Prozessoren kann eine bestimmte Anzahl von Threads parallel ausführen. Wie groß die Anzahl der Threads ist, hängt von der Architektur der Grafikkarte ab. In aktuellen Nvidia-Grafikkarten findet man 64, 128 oder 192 Threads pro Multiprozessor. Schließlich folgt die Ausgabe der GPU-Taktfrequenz. Alle weiteren Angaben werden später noch genauer betrachtet.
Organisation der GPU-Threads
Die auf der GPU laufenden Threads werden in Grids und Blöcken organisiert (Bild 1). Alle Threads, die von einem Prozessorkern gleichzeitig gestartet werden, sind in einem Grid angeordnet. Diese Threads benutzen alle gemeinsam den gleichen globalen Speicher. Ein Grid wird aus Thread-Blöcken aufgebaut. Die Threads eines Blocks können interagieren, das heißt, die Threads lassen sich synchronisieren und alle Threads benutzen den Speicherbereich des jeweiligen Blocks gemeinsam (Shared Memory). Die Nummer eines Blocks im Grid wird durch die Variable blockIdx angegeben. Die Nummer eines Threads im Block wird mit der Variablen threadIdx angegeben, beide haben Sie bereits im ersten Beispielprogramm kennengelernt.
Die Architekturvon Grids, Blöcken und Threads(Bild 1)
Autor
Die Variablen blockIdx und threadIdx sind Vektoren, die aus drei vorzeichenlosen Integer-Variablen bestehen:
- blockIdx.x, blockIdx.y, blockIdx.z sowie
- threadIdx.x, threadIdx.y, threadIdx.z.
- blockDim.x, blockDim.y, blockDim.z.
threadNummer = blockIdx.<span class="hljs-keyword">x</span> * blockDim.<span class="hljs-keyword">x</span> + threadIdx.<span class="hljs-keyword">x</span>
Die Variable blockDim.x enthält dann die Zahl 10 und die Variable threadIdx.x läuft von 0 bis 9 für jeden Block.Nun folgt ein genauerer Blick auf den Aufruf der Kernel-Funktion mit den drei spitzen Klammern. Die dort angegebenen Zahlen kontrollieren nämlich die Anzahl der Blöcke und die Anzahl der Threads pro Block. In Listing 1 wurden hier einfache Integer-Zahlen angegeben, da das Thread-Arrangement nur eindimensional war. Da die Blöcke und Threads aber auch einen zwei- oder dreidimensionalen Aufbau haben können, gibt es in den CUDA-Bibliotheken eine spezielle Struktur für diese Parameter namens dim3. Diese Struktur enthält drei vorzeichenlose Integer-Variablen, um jeweils die Größen in x-, y- und z-Richtung für die Blöcke im Grid und für die Threads im Block aufzunehmen. Der Aufruf eines Kernels kann dann folgendermaßen aussehen:
/<span class="hljs-regexp">/ ... </span>
<span class="hljs-regexp">/</span><span class="hljs-regexp">/ Block-Größe (hier eindimensional) </span>
<span class="hljs-regexp">dim3 block(5); </span>
<span class="hljs-regexp">/</span><span class="hljs-regexp">/ Grid-Größe (hier eindimensional) </span>
<span class="hljs-regexp">dim3 grid(arraySize /</span> <span class="hljs-number">5</span>);
<span class="hljs-regexp">//</span> Kernel aufrufen
callKernel <<<grid, block<span class="hljs-meta">>>> </span>();
<span class="hljs-regexp">//</span> ...
Die Aufteilung der Threads und Blöcke ist in manchen Fällen mit etwas Ausprobieren verbunden, um die größte Performance bei der Kernel-Ausführung zu erzielen.
Die Matrixmultiplikation
Die Matrixmultiplikation ist ein relativ einfacher und überschaubarer Algorithmus, der aber dennoch in Wissenschaft und Technik von enormer Wichtigkeit ist. Dieser Algorithmus soll nun für die normale CPU und zum Vergleich mit CUDA für die GPU implementiert werden. Bei diesem Beispiel wird auch der Performancegewinn beobachtet. Standardalgorithmus für das Multiplizieren zweier quadratischer Matrizen (oder Arrays) sind drei ineinandergeschachtelte Schleifen. Das Ergebnis ist ein Array. Angenommen es gibt die drei gleich großen Arrays a, b und c, wobei c das Ergebnis-Array sein soll. Die Arrays haben jeweils die Größe n * n. Dann sieht der Code so aus:
// Arrays deklarieren und initialisieren
for(<span class="hljs-built_in">int</span> i = <span class="hljs-number">0</span>; i < <span class="hljs-built_in">n</span>; i++)
{
for(<span class="hljs-built_in">int</span> j = <span class="hljs-number">0</span>; j < <span class="hljs-built_in">n</span>; j++)
{
for(<span class="hljs-built_in">int</span> k = <span class="hljs-number">0</span>; k < <span class="hljs-built_in">n</span>; k++)
{
c[i,j] += a[i,k] * b[k,j];
}
}
}
Diese sehr einfache Implementierung funktioniert recht gut für kleine Arrays, die in den 1st-Level-Cache der CPU passen. Bei großen Arrays wird der Algorithmus allerdings überproportional langsamer. Man kann allgemein sagen: Verdoppelt man die Kantenlänge n der drei Arrays, so muss theoretisch achtmal soviel gerechnet werden. Bei sehr großen Arrays wird man aber wesentlich längere Rechenzeiten finden, da die benötigten Daten nicht mehr im Cache gehalten werden können.In Listing 4 finden Sie eine einfache Matrixmultiplikation sowohl für die Ausführung auf einer CPU als auch auf der GPU mit CUDA.
Listing 4: Einfache Matrixmultiplikation (Teil 1)
#include <span class="hljs-string">"cuda_runtime.h"</span> <br/>#include <span class="hljs-string">"device_launch_parameters.h"</span> <br/>#include &lt;stdio.h&gt; <br/>#include &lt;<span class="hljs-built_in">time</span>.h&gt; <br/>#<span class="hljs-built_in">define</span> BLOCK_WIDTH <span class="hljs-number">32</span> <br/><br/>cudaError_t cudaMethod(<span class="hljs-built_in">float</span>* a, <span class="hljs-built_in">float</span>* b, <br/> <span class="hljs-built_in">float</span>* c, int n); <br/>int cpuTest(<span class="hljs-built_in">float</span>* a, <span class="hljs-built_in">float</span>* b, <span class="hljs-built_in">float</span>* c, int n); <br/><br/>__global__ void matMultKernel(<span class="hljs-built_in">float</span>* d_a, <span class="hljs-built_in">float</span>* d_b, <br/> <span class="hljs-built_in">float</span>* d_c, int n) <br/>{ <br/> int iCol = blockIdx.x * blockDim.x + threadIdx.x; <br/> int iRow = blockIdx.y * blockDim.y + threadIdx.y; <br/><br/> // Was zu groß ist, wird nicht gerechnet! <br/> <span class="hljs-keyword">if</span> ((iRow &lt; n) &amp;&amp; (iCol &lt; n)) <br/> { <br/> <span class="hljs-built_in">float</span> <span class="hljs-built_in">sum</span> = <span class="hljs-number">0.</span><span class="hljs-number">0f</span>; <br/><br/> <span class="hljs-keyword">for</span> (int k = <span class="hljs-number">0</span>; k &lt; n; k++) { <br/> <span class="hljs-built_in">sum</span> += d_a[iRow * n + k] * d_b[k * n + iCol]; <br/> } <br/> d_c[iRow * n + iCol] = <span class="hljs-built_in">sum</span>; <br/> } <br/>} <br/><br/>int main() <br/>{ <br/> double start, ende; // für Zeiten <br/> const int n = <span class="hljs-number">1024</span>; // Array-Größe <br/> int iRet; <br/> <span class="hljs-built_in">float</span>* a = <span class="hljs-built_in">new</span> <span class="hljs-built_in">float</span>[n * n]; <br/> <span class="hljs-built_in">float</span>* b = <span class="hljs-built_in">new</span> <span class="hljs-built_in">float</span>[n * n]; <br/> <span class="hljs-built_in">float</span>* c = <span class="hljs-built_in">new</span> <span class="hljs-built_in">float</span>[n * n]; <br/><br/> // Arrays initialisieren <br/> <span class="hljs-keyword">for</span> (int i = <span class="hljs-number">0</span>; i &lt; n; i++) { <br/> <span class="hljs-keyword">for</span> (int j = <span class="hljs-number">0</span>; j &lt; n; j++) { <br/> int <span class="hljs-literal">ind</span> = i * n + j; <br/> a[<span class="hljs-literal">ind</span>] = (<span class="hljs-built_in">float</span>)(i + j); <br/> b[<span class="hljs-literal">ind</span>] = (<span class="hljs-built_in">float</span>)(i - j + <span class="hljs-number">1</span>); <br/> c[<span class="hljs-literal">ind</span>] = <span class="hljs-number">0.</span><span class="hljs-number">0f</span>; <br/> } <br/> } <br/> // Auf der CPU rechnen <br/> //start = clock(); <br/> //iRet = cpuTest(a, b, c, n); <br/> //ende = clock() - start; <br/><br/> // Kernel einmal vorübersetzen <br/> iRet = cudaMethod(a, b, c, n); <br/> // Auf der GPU rechnen <br/> start = clock(); <br/> iRet = cudaMethod(a, b, c, n); <br/> ende = clock() - start; <br/><br/> // Fehler <span class="hljs-literal">und</span> Zeit ausgeben <br/> <span class="hljs-keyword">if</span> (iRet == <span class="hljs-number">0</span>) <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">"Kein Fehler.\n"</span>); <br/> <span class="hljs-keyword">else</span> <br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">"Error!!!!\n"</span>); <br/><br/> <span class="hljs-built_in">printf</span>(<span class="hljs-string">"Zeit: %.3f msek\n"</span>, ende); <br/> cudaDeviceReset(); <br/> <span class="hljs-built_in">return</span> <span class="hljs-number">0</span>; <br/>} <br/><br/>cudaError_t cudaMethod(<span class="hljs-built_in">float</span>* a, <span class="hljs-built_in">float</span>* b, <br/> <span class="hljs-built_in">float</span>* c, int n) <br/>{ <br/> <span class="hljs-built_in">float</span> *d_a = <span class="hljs-number">0</span>; <br/> <span class="hljs-built_in">float</span> *d_b = <span class="hljs-number">0</span>; <br/> <span class="hljs-built_in">float</span> *d_c = <span class="hljs-number">0</span>; <br/> cudaError_t cudaStatus; <br/> cudaStatus = cudaSetDevice(<span class="hljs-number">0</span>); <br/> <span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) goto Error; <br/><br/> // Speicher allokieren <br/> cudaStatus = cudaMalloc((void**)&amp;d_c, n * n * sizeof(<span class="hljs-built_in">float</span>)); <br/> <span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) goto Error; <br/><br/> cudaStatus = cudaMalloc((void**)&amp;d_a,<br/> n * n * sizeof(<span class="hljs-built_in">float</span>)); <br/> <span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) goto Error; <br/><br/> cudaStatus = cudaMalloc((void**)&amp;d_b, <br/> n * n * sizeof(<span class="hljs-built_in">float</span>)); <br/> <span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) goto Error; <br/><br/> // Daten kopieren <br/> cudaStatus = cudaMemcpy(d_a, a, <br/> n * n * sizeof(<span class="hljs-built_in">float</span>), cudaMemcpyHostToDevice); <br/> <span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) goto Error; <br/><br/> cudaStatus = cudaMemcpy(d_b, b, <br/> n * n * sizeof(<span class="hljs-built_in">float</span>), cudaMemcpyHostToDevice); <br/> <span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) goto Error; <br/><br/> // Grid- <span class="hljs-literal">und</span> Blockgröße ermitteln <br/> int nBlocks = n / BLOCK_WIDTH; <br/> <span class="hljs-keyword">if</span> (n <span class="hljs-symbol">%</span> BLOCK_WIDTH != <span class="hljs-number">0</span>) <br/> nBlocks++; <br/> // Kernel aufrufen <br/> dim3 dimGrid(nBlocks, nBlocks, <span class="hljs-number">1</span>); <br/> dim3 dimBlock(BLOCK_WIDTH, BLOCK_WIDTH, <span class="hljs-number">1</span>); <br/> matMultKernel &lt;&lt;&lt; dimGrid, dimBlock &gt;&gt;&gt;(d_a, <br/> d_b, d_c, n); <br/> // Warten, bis Kernel fertig ist <br/> cudaStatus = cudaDeviceSynchronize(); <br/> <span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) goto Error; <br/> // Ergebnis kopieren <br/> cudaStatus = cudaMemcpy(c, d_c, <br/> n * n * sizeof(<span class="hljs-built_in">float</span>), cudaMemcpyDeviceToHost); <br/> <span class="hljs-keyword">if</span> (cudaStatus != cudaSuccess) goto Error; <br/><br/> // Speicher freigeben <br/> cudaFree(d_c); <br/> cudaFree(d_a); <br/> cudaFree(d_b); <br/><br/> <span class="hljs-built_in">return</span> cudaStatus; <br/><br/>Error: <br/> cudaFree(d_c); <br/> cudaFree(d_a); <br/> cudaFree(d_b); <br/><br/> <span class="hljs-built_in">return</span> cudaStatus; <br/>} <br/><br/>int cpuTest(<span class="hljs-built_in">float</span>* a, <span class="hljs-built_in">float</span>* b, <span class="hljs-built_in">float</span>* c, int n) <br/>{ <br/> // Einfache Matrixmultiplikation <br/> <span class="hljs-keyword">for</span> (int i = <span class="hljs-number">0</span>; i &lt; n; i++) <br/> { <br/> <span class="hljs-keyword">for</span> (int j = <span class="hljs-number">0</span>; j &lt; n; j++) <br/> { <br/> <span class="hljs-built_in">float</span> <span class="hljs-built_in">sum</span> = <span class="hljs-number">0.</span><span class="hljs-number">0f</span>; <br/> <span class="hljs-keyword">for</span> (int k = <span class="hljs-number">0</span>; k &lt; n; k++) <br/> { <br/> <span class="hljs-built_in">sum</span> += a[i * n + k] * b[k * n + j]; <br/> } <br/> c[i * n + j] = <span class="hljs-built_in">sum</span>; <br/> } <br/> } <br/> <span class="hljs-built_in">return</span> <span class="hljs-number">0</span>; <br/>}
Das Beispiel beginnt nach den Header-Dateien wieder mit der Kernel-Funktion matMultKernel. Da hier mit quadratischen Arrays gerechnet wird, ist es sinnvoll, auch zweidimensionale Indizes für die Threads zu benutzen. Aus dem Block-Index blockIdx und dem Thread-Index threadIdx werden dann die Indizes für Zeile und Spalte errechnet, die anschließend in der Schleife mit der Laufvariablen k multipliziert und summiert werden.Die main-Methode arbeitet genauso wie in den vorhergehenden Beispielen: Daten bereitstellen, Speicher anlegen, Daten kopieren, Kernel aufrufen, Ergebnisse kopieren.Es lohnt sich den Kernel-Aufruf genauer anzuschauen: Ganz am Anfang des Beispielprogramms wurde der Platzhalter BLOCK-WIDTH auf den Wert 32 gesetzt. Die Arrays bestehen aus n Zeilen und Spalten. Sind die Arrays sehr groß (im Beispiel: 1 024 * 1 024 Elemente), kann die Berechnung nicht in einem einzigen Thread-Block ausgeführt werden, denn in einem Block sind üblicherweise nur 1 024 Threads erlaubt. Man muss also mehrere Blöcke verwenden. In der Variablen nBlocks wird die Anzahl der benötigten Blöcke in horizontaler und vertikaler Richtung berechnet und abgelegt. Das heißt, für die Berechnung werden nBlocks * nBlocks Thread-Blöcke benutzt. Jeder Thread-Block wiederum enthält BLOCK_WIDTH * BLOCK_WIDTH Threads. Mit diesen Informationen wird nun die Kernel-Funktion aufgerufen. In diesem Zusammenhang ist es nicht ganz unwichtig, in der Kernel-Funktion mit einem if-Statement zu prüfen, ob die benutzen Werte für Zeile und Spalte immer kleiner als n sind. Wenn n %BLOCK_WIDTH nämlich nicht null ist, wird nBlocks um eins erhöht.In den äußeren Blöcken können dann aber nicht alle Threads mit sinnvollen Daten rechnen. Diese Thread-Indizes werden durch den if-Befehl ignoriert.Das Beispiel in Listing 4 enthält auch normalen C-Code, um die Matrixmultiplikation auf der CPU in einem einzigen Thread auszuführen. In der main-Methode ist dieser Code als Kommentarzeilen eingebaut. Dieser Teil kann bei Bedarf zum Zeitvergleich aktiviert werden.Die Zeiten für die CUDA-Berechnung werden inklusive der erforderlichen Datenübertragungszeiten in und aus der GPU gemessen.Weiterhin wird beim CUDA-Aufruf der Kernel zweimal ausgeführt. Der zweite Aufruf ist normalerweise schneller, denn der Code muss nicht mehr für die Ausführung auf der GPU vorbereitet werden.Die Zeitmessung erfolgt mit der einfachen C-Funktion clock(), die in der Datei time.h deklariert ist. Unter Windows liefert die Funktion Millisekunden zurück. Unter Linux liefert diese Funktion die Zeit in Mikrosekunden zurück.Ein Zeitvergleich der CPU- und der CUDA-Implementierung ergibt für eine 1 024 * 1 024-Matrix ein deutliches Ergebnis:
<span class="hljs-string">CPU :</span> <span class="hljs-number">6</span>,<span class="hljs-number">396</span> Sekunden,
<span class="hljs-string">CUDA:</span> <span class="hljs-number">0</span>,<span class="hljs-number">044</span> Sekunden
Die CUDA-Routine ist also etwa 145-mal so schnell wie der Code auf der CPU. Allerdings muss man natürlich erwähnen, dass man den CPU-Code ebenfalls parallelisieren und entsprechend kürzere Rechenzeiten erhalten kann – zum Beispiel mit OpenMP.Der Code aus dem zuletzt gezeigten Beispiel ist aber noch nicht optimal. Man kann das noch etwas schnellere Shared Memory für die Berechnungen nutzen (Bild 2). Allerdings ist dieser Speicherbereich begrenzt und die drei Arrays passen nicht komplett hinein. Die wichtigsten Änderungen für den Einsatz des schnelleren Speichers finden Sie in Listing 5.
Architekturdes GPU-Speichers(Bild 2)
Autor
Listing 5: Verbesserte Matrixmultiplikation
<span class="hljs-strong">__global__</span> void matMultKernel(float<span class="hljs-bullet">* d_a, </span><br/><span class="hljs-bullet"> float*</span> d<span class="hljs-emphasis">_b, float* d_</span>c, int n) <br/>{ <br/> // Hilfsarrays im schnellen Speicher <br/> <span class="hljs-strong">__shared__</span> float ads[<span class="hljs-string">TILE_WIDTH</span>][<span class="hljs-symbol">TILE_WIDTH</span>]; <br/> <span class="hljs-strong">__shared__</span> float bds[<span class="hljs-string">TILE_WIDTH</span>][<span class="hljs-symbol">TILE_WIDTH</span>]; <br/><br/> int bx = blockIdx.x; <br/> int by = blockIdx.y; <br/><br/> int tx = threadIdx.x; <br/> int ty = threadIdx.y; <br/><br/> int iRow = by <span class="hljs-bullet">* TILE_WIDTH + ty; </span><br/><span class="hljs-bullet"> int iCol = bx *</span> TILE<span class="hljs-emphasis">_WIDTH + tx; </span><br/><br/><span class="hljs-emphasis"> float sum = 0.0f; </span><br/><br/><span class="hljs-emphasis"> // Hilfsarrays mit Daten füllen </span><br/><span class="hljs-emphasis"> for (int m = 0; m &lt; n / TILE_</span>WIDTH; m++) { <br/> ads[<span class="hljs-string">ty</span>][<span class="hljs-symbol">tx</span>] = <br/> d<span class="hljs-emphasis">_a[iRow * n + m * TILE_</span>WIDTH + tx]; <br/> bds[<span class="hljs-string">ty</span>][<span class="hljs-symbol">tx</span>] = <br/> d<span class="hljs-emphasis">_b[(m * TILE_</span>WIDTH + ty) <span class="hljs-bullet">* n + iCol]; </span><br/><span class="hljs-bullet"> __syncthreads(); </span><br/><span class="hljs-bullet"> // Multiplikation </span><br/><span class="hljs-bullet"> for (int k = 0; k &lt; TILE_WIDTH; k++) { </span><br/><span class="hljs-bullet"> sum += ads[ty][k] *</span> bds[<span class="hljs-string">k</span>][<span class="hljs-symbol">tx</span>]; <br/> } <br/> <span class="hljs-emphasis">__syncthreads(); </span><br/><span class="hljs-emphasis"> } </span><br/><br/><span class="hljs-emphasis"> d_</span>c[iRow * n + iCol] = sum; <br/>}
In dieser Variante werden Teile der GPU-Arrays d_a und d_b in das schnellere Shared Memory kopiert, dann wird die Multiplikation mit den kleinen Hilfs-Arrays ausgeführt und die Ergebnisse werden an die richtige Stelle im Ergebnis-Array d_c kopiert. Dabei müssen die Threads aber an zwei Stellen mit syncthreads() synchronisiert werden. Beim ersten Aufruf wird sichergestellt, dass alle erforderlichen Daten in die kleinen Arrays kopiert wurden. Beim zweiten Aufruf sollten dann alle Multiplikationen ausgeführt sein, und das Ergebnis kann korrekt im Ziel-Array abgelegt werden.Statt des Platzhalters BLOCK_WIDTH wird nun der Name TILE_WIDTH benutzt. Dies muss auch in der main-Funktion entsprechend angepasst werden. Setzt man die TILE_WIDTH im Beispiel auf die Werte 4, 8, 16 oder 32 und misst die Ausführungszeiten, so stellt man fest, dass die Zeiten mit einer größeren TILE_WIDTH besser werden. Je größer TILE_WIDTH ist, desto mehr Rechenoperationen können mit den kleinen Arrays pro Kopiervorgang durchgeführt werden. Bei einem Wert von 32 ist jedoch Schluss, da ein Block nur 1 024 (also 32 * 32) Threads ausführen kann.Mit einer TILE_WIDTH von 32 wurde eine Rechenzeit von 0,016 Sekunden gemessen, das ist also noch einmal fast dreimal schneller als die Variante aus Listing 4. Diese Werte variieren von Rechner zu Rechner. Einen besonderen Einfluss hat sicherlich die Grafikkarte. Je mehr CUDA-Kerne zur Verfügung stehen, desto kürzer sind die Rechenzeiten.
Zusammenfassung
Durch die Berechnungen auf der GPU können Sie bestimmte Code-Teile einer Anwendung deutlich beschleunigen. Dafür müssen jedoch einige Bedingungen erfüllt werden:- Die Datenmengen, die in die GPU kopiert werden müssen, dürfen nicht zu groß sein.
- Mit den kopierten Daten müssen möglichst viele Rechenoperationen ausgeführt werden.
- Die Daten müssen möglichst in der gleichen Weise verarbeitet werden.
- Der Rechenalgorithmus muss an die Grafikhardware angepasst werden, um optimale Performance-Steigerungen zu erzielen.
Fussnoten
- CUDA-Download, http://www.dotnetpro.de/SL1904CUDA1
- CUDA (Wikipedia), http://www.dotnetpro.de/SL1904CUDA2
