CUDA

Hinweis: Bitte beachten Sie, dass dieser Artikel / diese Kategorie sich entweder auf ältere Software/Hardware Komponenten bezieht oder aus sonstigen Gründen nicht mehr gewartet wird. Diese Seite wird nicht mehr aktualisiert und ist rein zu Referenzzwecken noch hier im Archiv abrufbar.

Im Jahr 2006 stellte Nvidia ihre "General Purpose Parallel Computing Architecure" für Anwendungen auf Nvidia Grafikkarten (GPUs) vor. Unter dem Namen CUDA (Compute Unified Device Architecture) wurde ein neues Programmiermodell und ein Set von Werkzeugen vorgestellt. CUDA ermöglicht die Entwicklung von Applikationen auf GPUs, die nicht auf grafische Operationen spezialisiert sein müssen ("General Computing"). Erst durch CUDA erhält ein Programmierer direkten Zugriff auf das Instruction Set und die Speicherverwaltung der GPU.

Die GPUs erreichen vor allem bei der Durchführung von rechenintensiven, hoch parallelen Anwendungen hervorragende Geschwindigkeiten.^[1] Genauer gesagt sind GPUs besonders für das Lösen von Problemen geeignet, bei denen die gleichen Algorithmen parallel auf verschiedenen Daten ausgeführt werden müssen. Dieses Konzept wird auch unter dem Namen "Single Instruction, Multiple Data" (SIMD) wiedergefunden.

Für die Entwicklung von CUDA Applikationen wird C mit einzelnen Erweiterungen als Programmiersprache verwendet. Für die Übersetzung der Quelldateien wird ein eigener Compiler mitgeliefert (nvcc^[2]).

Historische Entwicklung

Zu Beginn stellten OpenGL und DirectX die einzige Möglichkeit dar mit einer GPU zu interagieren. Diese APIs sind vor allem auf die Entwicklung von Multimedia-Anwendungen ausgelegt. Aus diesem Grund mussten zu Beginn Wege gefunden, einfache Berechnungen auf graphische Operationen umzusetzen, um die Rechenkapazität der GPUs nutzen zu können. Die Berechnung waren aber stets limitiert auf die Funktionen, die von der jeweiligen API zur Verfügung gestellt wurden. Doch nicht nur die eingeschränkte Funktionalität sondern auch die erhöhte Komplexität, die die APIs mit sich brachten, waren Grund für die mäßige Verbreitung von GPU-Computing.^[3]

Die oben genannten Probleme wurden durch den Release von CUDA durch Nvidia beseitigt. Zum einen wurde ein Programmiermodell vorgestellt, dass eine einfache und effiziente parallele Entwicklung auf GPUs ermöglicht. Zum anderen wurden die Grafikchips um Hardware erweitert.

CUDA Architektur

Typischerweise stehen einem Programmierer für eine Applikation der Host (CPU) und ein oder mehrere Devices (GPUs) zur Verfügung. Mit diesen beiden Komponenten gilt es nun einen Algorithmus zu implementieren, der ein gegebenes Problem auf dem schnellst möglich Weg löst. Auf diesem Wer wird es Abschnitte geben, die seriell ausgeführt werden müssen (geeignet für die CPU). Jedoch werden auch Abschnitte auftauchen, wo dieselben arithmetischen Operationen gleichzeitig auf Datenstrukturen angewendet werden. Genau diese Teile sind geeignet für eine Ausführung auf der GPU.

Somit besteht die Applikation aus Teilen, die auf der CPU ausgeführt werden und aus Teilen, die für die GPU vorgesehen sind. Dabei müssen alle Quelldateien, die CUDA-spezifische C-Erweiterungen enthalten, mit dem CUDA-Compiler "nvcc" kompiliert werden. Der CUDA-Compiler "nvcc" kompiliert dabei die Teile für die GPU in PTX oder Object Code und übergibt den Host-Teil an den spezifizierten Host-Compiler (unter Linux z.B. der gcc).^[4] Dadurch wird das Separieren von Host- und Device-spezifischen Code automatisch übernommen.

CUDA Konzepte

Um die maximale Performance mit einer CUDA-Applikation erreichen zu können ist es wichtig, dass verstanden wird, wie die GPU arbeitet. Erst dann ist es möglich die Multiprozessoren der GPU so auszulasten, dass deren Leistungsfähigkeit optimal ausgenutzt wird.

Das Ziel ist die Applikation optimal für die GPU auszulegen:

Auslastung möglichst vieler Threads, um Latenzen zu kaschieren
Anwendung der gleichen Logik auf möglichst viele Threads ("Single Instruction, Multiple Thread" und Vermeidung von "Divergent Branches")
Vermeidung von komplexen arithmetischen Operationen
Vermeidung von Abhängigkeiten zwischen Threads
Minimierung der Speichertransfers zwischen Karte und Host
Nutzung der verschiedenen Speicher, die die GPU zur Verfügung stellt (Global Memory, Shared Memory, Constant Memory etc.). ^[5]
Optimierung der Speicherzugriffe ("Coalesced Access")

Ein Paradebeispiel für den sinnvollen Einsatz von GPUs ist die Multiplikation zweier Matrizen. Diese wird auch oft in Büchern als Beispiel für die Programmierung und den richtigen Einsatz der verschiedenen Speicher der GPU hergenommen.

Werden zum Beispiel zwei 1000x1000 Matrizen miteinander multipliziert, müssen insgesamt 1000000 Skalarprodukte berechnet werden. Jedes dieser Produkte setzt sich aus 1000 Multiplikationen und 1000 Additionen zusammen. Hierbei lässt sich gut erkennen, dass eine Matrizenmultiplikation mit Matrizen großer Dimensionen sehr viele Rechenoperationen erfordert.

Mit CUDA können die Berechnungen der Skalarprodukte parallelisiert werden, da alle auf die gleiche Weise berechnet werden und Daten-Parallelität genutzt werden kann (voneinander unabhängige Datensätze).

Kernels

CUDA C besitzt gegenüber traditionellem C mehrere Erweiterungen. Der wohl wichtigste Punkt ist die Möglichkeit sogenannte "Kernels" definieren zu können. Ein Kernel kann mit einer normalen C-Funktion verglichen werden. Die normale C-Funktion wird jedoch bei einem Aufruf nur einmal ausgeführt. Ein Kernel hingegen wird bei seinem Aufruf von mehreren Threads parallel auf der GPU ausgeführt. Die Anzahl der Threads, die den Kernel ausführen sollen, werden beim Aufruf als Parameter mit angeführt:

Add<<<2,512>>>(A,B,C);

In diesem Beispiel wird der Kernel mit dem Namen "Add" von zwei CUDA Blöcken mit jeweils 512 Threads parallel ausgeführt. Das heißt 1024 Threads führen den Kernel zugleich auf der GPU aus. Zusätzlich werden dem Kernel noch die Parameter A,B und C übergeben.

Im vorigen Beispiel der Matrix-Multiplikation wäre es z.B. möglich, dass für die Berechnung eines Skalarprodukts ein Thread gestartet wird. Dies würde als Resultat einen Kernel ergeben, der ein Skalarprodukt berechnen kann und von 1000000 Threads gleichzeitig auf der GPU ausgeführt wird.

Dabei ist zu beachten, dass ein CUDA-Thread nicht mit einem CPU-Thread verglichen werden kann. CUDA Threads sind leicht-gewichtiger und können schneller erzeugt und vom Scheduler zugeteilt werden, da sie sehr auf die darunterliegende Hardware zugeschnitten sind.

Beispielhafte Anordnung von zwei Blöcken mit 512 Threads

Blöcke

CUDA-Blöcke können als eine Gruppierung von Threads aufgefasst werden. Diese Gruppierung ist insofern interessant, da Threads und Blöcke unterschiedlich auf die verschiedenen Speicher der GPU zugreifen können. Wie viele Blöcke und Threads gestartet werden können ist von GPU zu GPU verschieden. Wurde bereits das CUDA Toolkit und das SDK auf einem Rechner installiert, so bietet das Programm "deviceQuery" einen schnellen Weg detaillierte Informationen über die eigene Grafikkarte zu erhalten.

Blöcke bestehen also aus mehreren Threads.
Mehrere Blöcke zusammen ergeben ein Grid.
Alle Threads eines Grids führen immer denselben Kernel aus.

Für die Unterscheidung der einzelnen Threads und Blöcke stehen unterschiedliche Systemvariablen ("Built-in variables") zur Verfügung. Diese Variablen ermöglichen es den Threads und Blöcken sich zur Laufzeit voneinander zu unterscheiden. So können für die Grid-Dimension zweidimensionale (2D) Werte und für die Block-Dimension dreidimensionale (3D) Werte definiert werden. Dies kann insofern von Vorteil sein, wenn mit mehrdimensionalen Datenstrukturen gearbeitet wird und Threads und Blöcke darauf zugreifen müssen. Dann kann nämlich direkt auf diese Strukturen zugegriffen werden und eine Umrechnung auf eindimensionale Indices wird nicht benötigt.

Block- und Thread-Hierarchie eines Grids

Zum Beispiel kann ein Kernel gestartet werden, der insgesamt aus 4 (2x2) Blöcken und 8 (4x2) Threads pro Block besteht. Die Grafik rechts verdeutlicht den schematischen Aufbau des Grids!

dim3 dimBlock(4,2,1);
dim3 dimGrid(2,2,1);
Add<<<dimGrid,dimBlock>>>(...);

Zur Laufzeit kann dann ein Block bzw. ein Thread auf seinen Index zugreifen, der ihm eindeutig zugewiesen ist. In der Grafik besitzt z.B. Thread (2,1,0) die Variablen threadIdx.x=2 und threadIdx.y=1. Der Block (1,0) hätte blockIdx.x=1 und blockIdx.y=0. Zu beachten ist, dass die Thread Indices nur pro Block gültig sind.

CUDA-Memories

Die GPU stellt Hardware-seitig verschiedene Speicher zur Verfügung, deren Aufbau sich auch Software-seitig widerspiegelt^[6]:

Lesen/Schreiben in Register pro Thread
Lesen/Schreiben in Local Memory pro Thread
Lesen/Schreiben in Shared Memory pro Block
Lesen/Schreiben in Global Memory pro Grid
Nur Lesen in Constant Memory pro Grid

Im Kernel legt der Programmierer fest in welchem Speicher die variablen verweilen sollen:

int var;  //Variable in einem Register
int var[100];  //Feld-Variable im Local Memory
__shared__ int s_var;  //Variable im Shared Memory
__device__ int g_var;  //Variable im Global Memory
__constant__ int c_var;  //Variable im Constant Memory

Diese verschiedenen Speicher unterscheiden sich auch in ihrer Zugriffsgeschwindigkeit:

Register: On-Chip Register (performant)
Shared Memory: On-Chip Memory (performant)
Local und Global Memory: Uncached off-chip memory (langsam)
Constant Memory: Cached off-chip memory (performant)

Common Practice

Für die effiziente Abarbeitung von Daten, die im Global Memory verweilen, hat sich folgende Abarbeitungsreihenfolge als vorteilhaft herausgestellt:

Teile die Daten im Global Memory in Portionen, die in den Shared Memory passen.
Jeder Thread-Block wird nun 1 dieser Teile bearbeiten.
Dazu lädt jeder Thread eines Blocks einen Teil des Global Memories in den Shared Memory. Dies hat den Vorteil, dass Daten parallel von mehreren Threads geladen werden.
Die Threads führen ihre Berechnungen auf dem vorher geladenen Teil durch.
Nach Beendigung werden die Ergebnisse zurück in den Global Memory geschrieben.

Vor allem bei Speicherbereichen, in denen mehrere Threads gleichzeitig Schreiben können, muss auf Race Conditions geachtet werden. Es wichtig, dass ohne expliziter Synchronisation undefiniert ist, in welcher Reihenfolge die Threads auf den Speicher zugreifen. Für diese Probleme stehen atomare Operationen und Synchronisationsbarrieren ("Barriers") zur Verfügung.

Solche Barrieren werden zum Beispiel schon beim Laden der Daten in den Shared Memory benötigt. Hier muss mit einem expliziten "__syncthreads" sicher gestellt werden, dass alle Threads ihre Daten in den Shared Memory geschrieben haben. Ansonsten kann es dazu kommen, dass auf bereits auf einen Speicherbereich zugegriffen wird, der noch nicht von einem Thread beschrieben wurde.

CUDA Applikationen

Auflistung verschiedener Applikationen

http://www.geforce.com/#/Hardware/Technologies/cuda/applications

CUDA Mathematica

Matlab plugin

http://www.nvidia.com/object/numerical-packages.html

CUDA Showcase

http://www.nvidia.com/object/cuda-apps-flash-new.html#

Einzelnachweise

Literatur

CUDA by Example

http://developer.nvidia.com/object/cuda-by-example.html

Programming Massively Parallel Processors: A Hands-on Approach

http://www.nvidia.com/object/io_1264656303008.html Einzelne Kapitel dieses Buches sind auch im Rahmen einer Vorlesung des Autors zum Download verfügbar(Gpu courses in Illinois (impact.crhc.illinois.edu)). Auf dieser Webseite könne außerdem die Präsentationen und Übungsaufgaben heruntergeladen werden: http://courses.engr.illinois.edu/ece498/al/Syllabus.html

Weitere Informationen

Aktuelle News

Informationen zu GPUs

Applikationen

Nvidia HPC Tesla

Nvidia Developer Webseiten

CUDA Trainingsmaterial und Vorlesungen

http://developer.nvidia.com/object/cuda_training.html

Tutorien zur Programmierung

CUDA Software

CUDA SDK Code Beispiele

http://developer.nvidia.com/cuda-cc-sdk-code-samples

CUDA Wrapper für andere Programmiersprachen

[1] CUDA Programming Guide Version 3.1.1

[2] Nvcc Compiler Manual

[3] Textbook Chapter 1 von Wen-mei W. Hwu und David Kirk

[4] CUDA Programming Guide Version 3.1.1, S. 16

[5] Stanford University - CUDA Memories

[6] Textbook Chapter 4 von Wen-mei W. Hwu und David Kirk Textbook Kapitel 4

[1]

[2]

[3]

[4]

[5]

[6]