# Testat - CUDA
## Fragenkatalog 1 (leicht)
---
**Wie kann man konstanten Speicher kopieren?**
- Via `cudaMemcpyToSymbol`
- Alternative: `cudaGetSymbolAddress` + `cudaMemcpy`
---
**Erklären Sie den Unterschied zwischen CPU und GPU.**
- GPU -> Graphics Processing Unit.
- Eine GPU ist ebenfalls ein Prozessor wie die CPU, allerdings ist sie optimiert für die Grafik-Berechnung bei PCs, Konsolen, Smartphones, Tablets und Co. Eine GPU ist also der Grafikprozessor eines Geräts.
<img src="https://i.imgur.com/PYM6JsU.png" alt="CPU vs GPU" width="80%" style="display: block; margin-left: auto; margin-right: auto;">
CPU
- Kleine Anzahl an (komplexen) Kernen O(10)
- Komplexe Kerne mit Steuerlogik zur Maximierung der Parallelität auf Befehlsebene
- Out-of-Order Ausführung
- Verzweigungsvorhersage
- Relativ große Menge an Hauptspeicher pro Kern
- Effiziente Ausnutzung von Caches
GPU
- Weniger Fokus auf Support-Strukturen (Control & Cache), sondern eher auf Processing Units
- Mehr Transistoren für Data Processing
- Hohe Anzahl an Kernen O(1000)
- Wenig Steuerlogik
- Sehr große Anzahl an Threads
- realisiert Parallele Datenberechnungen
- Code mit hoher Berechnungsintensität
- Hilft Speicherlatenz zu verbergen
- Kleinere caches
---
**Wird alles in der Raytracing Lab mit einem Kernel ausgeführt?**
Ja.
---
**Erklären Sie die Funktionalität von `cudaMemcpy` inklusive Argumente.**
```c
cudaError_t cudaMemcpy(void *dst, const void *src, size_t count
enum cudaMemcpyKind kind)
```
- Kopiert Daten vom Speicher auf den src zeigt in den Speicher auf den dst zeigt
- `dst`: Ziel-Speicheradresse
- `src`: Quell-Speicheradresse
- `count`: Größe in Bytes zum Kopieren
- `kind`: Art der Übertragung
- cudaMemcpyHostToHost
- cudaMemcpyHostToDevice
- cudaMemcpyDeviceToHost
- cudaMemcpyDeviceToDevice
Die Speicherbereiche dürfen sich nicht überschneiden
---
## Fragenkatalog 2 (mittel)
---
**Erklären Sie die Funktionalität von `global`.**
- Aufrufbar von Host und wird im Device ausgeführt
- signalisiert Kernel Code - Kernel wird ausgeführt (auf der GPU ?)
- `__global__` void KernelFunc()
---
**Erklären Sie die Funktionalität von `cudaMalloc` inklusive Argumente und warum muss der erste Parameter ein Pointer sein?**
```c
cudaError_t cudaMalloc(void **devPointer, size_t size)
```
- Alloziert Device Speicher
- `**devPointer`: Pointer auf dem zugewiesenen Device Speicher
- `size`: Angefragte Größe in Bytes
- Der zugewiesene Speicher ist für jede Art von Variablen geeignet
- (Begründung: Weil der zugewiesene Speicher für jede Art von Variablen geeignet ist und sein soll, ein void Pointer lässt sich problemlos zur benötigten Art casten) - weil ein Pointer auf eine Speicheradresse verweist
- TODO
---
**Erläutern Sie die Speicherhierachie.**
- Die Speicherhierarchie ist eine Organisationsform des Speichers in Computersystemen.
- Sie basiert auf der Grundlage, dass ein Speicher nicht gleichzeitig alle Anforderungen an Zugriffszeit, Datenübertragungsrate, Preis, Größe oder Energieverbrauch erfüllen kann.
---
## Fragenkatalog 3 (schwer)
---
Gegeben sei folgende Abbildung
(Originalbild)
Ungefähre Skizze (Pfeile sind unvollständig oder inkorrekt can't remember)
Ergänzen Sie in den orangen Boxen die jeweiligen Speicherstellen.
Antwort: Siehe Abbildung (Global, Constant?, Lokal, Geteilt)
Fortführende Frage: Wie kann man shared Daten synchronisieren?
- `__syncthreads();`
---
Gegeben ist folgender Code-Abschnitt mit einem Fehler. beschreiben Sie wie der Fehler zustande kommt.
- Unbekannter Code
- fehlende Synchronisation, ohne die ein Read after Write entstanden ist
---
**Gegeben sei folgender Code: Wo fehlt sync?**
```c
const size_t N 1024; // Array size
int host_array[N]; // Array allocated to host
int *dev_array; // Pointer to a device array
// Device code
__global__ void partialSum(int *array) {
__shared__ long s_array[N];
int idx = threadIdx.x;
s_array[idx] = array[idx];
// HIER MUSS __syncthreads()
for(unsigned int stride = blockDim.x / 2; stride > 0; stride /= 2) {
if(idx < stride) {
s_array[idx] += s_array[idx + stride];
}
__syncthreads();
}
if(idx == 0) {
array[0] = s_array[0];
}
}
```
```c
__global__ void stencil_1d(int *in, int *out) {
__shared__ int tempB[BLOCK_SIZE + 2 * RADIUS];
int gindex = threadIdx.x + blockIdx.x * blockDim.x;
int lindex = threadIdx.x + RADIUS;
// Read input elements into shared memory
temp[lindex] = in[gindex];
if(threadIdx.x < RADIUS) {
temp[lindex - RADIUS] = in[gindex - RADIUS];
temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];
}
// HIER MUSS __syncthreads()
// Apply the stencil
int result = 0;
for(int offset = -RADIUS; offset <= RADIUS; offset++) {
result += temp[lindex + offset];
}
// Store the result
out[gindex] = result;
}
```
- `__syncthreads();` nach der `if`, um zu gewährleisten, dass alle Daten vorhanden sind
- `temp` kann ansonsten nicht fertig bearbeitet sein
- gibt also race condition zwischen Threads in verschiedenen Warps
---
**Nennen Sie die Speicherstelle, Scope, Lebenszeit für die Variablen `a`, `b`, `c`, `d` und `array`.**
```c
#include "stdio.h"
// Memory location, scope, lifetime of a, b, c, d, array
__shared__ int a;
__device__ int b;
__constant__ int c;
__device__ void kernel() {
float d;
float array[100];
}
int main() {
// ...
}
```
- `a`: Shared, Block, Kernel
- `b`: Global, Grid, Application
- `c`: Constant, Grid, Application
- `d`: Register, Thread, Kernel
- `array`: Local, Thread, Kernel
**Fortführende Frage: Welche Art Speicherstelle würde man für eine Matrixmultiplikation verwenden?**
- Shared: Weil Effizienter. Nachfolgende Einträge können schneller gelesen werden
- (in Vorlesung war das Beispiel mit den Fliesen (Tiles)) -> Threads können sich nämlich auf einen kleinen Ausschnitt der Input Matrix fokussieren und diese Werte im shared memory kollaborativ benutzen.
---
**Irgendwas mit: Erklären Sie den Aufbau von GPU o.Ä.**
- Multiprozessor von Multiprozessoren
- die individuellen Multiprozessoren sind SIMT Architekturen
- einige SM (= Streaming Multiprocessor) und SMC (= Streaming Multiprocessor Controller)
- viele SPs (= Streaming Prozessorkerne)
- viele Transistoren für die Datenverarbeitung (?)
- kleine Caches
**Erklären wie ein Block aufgebaut ist...**
Block
- Blöcke sind Elemente innerhalb eines Grids können parallel berechnet werden
- Blöcke sind unabhängig voneinander
- Skalierbarkeit gut
- Block wird von einem SM ausgeführt
- Auf Threads abgebildet
- Threads organisiert in einer Hierarchie von Blöcken
- Menge von Threads, die untereinander kooperieren durch Barrieresynchronisation und Zugriff auf blockeigenen Speicherplatz
- Maximal 1024 Threads pro Block auf aktuellen GPUs
---
## Zusatz (Potentiell möglich?)
### Einfach
---
**Wofür steht CUDA / Was macht CUDA?**
- Compute Unified Device Architecture
- CUDA execution model: Ausführung eines Kernel als grid von thread blocks
- CUDA virtualisiert Multiprozessoren
- folgt SPMD (Single Program Multiple Data) model
- durch die Architektur wird Parallelisierung realisiert
- Vorteil: Produktivität geseigert --> automatisches Thread management
- CUDA device memories: registesr, shared memory, global memory / constant memory
- (Nvidia entwickelte API, mit der Programmteile durch den Grafikprozessor (GPU) abgearbeitet werden können)
**Was versteht man unter GPU?**
- Graphics Processing Unit
**Verwendung:**
- Generierung von 2D & 3D Grafiken, Bilder und Videos
- Anwendungen in GUIs, Videospielen, visuellen Bildgebungsanwendungen
- Visual Computing, lineare Algebra, Vektoroperationen, ML, ...
**Moderne GPU:**
- hochparallel, mutithreadingfähiger Multiprozessor
- optimiert für visual computing
- Ermöglicht Echtzeit Interaktionen mit berechneten Objekten
- Einheitliche Grafik- und Computerarchitektur
- Programmierbarer Grafikprozessor
- Skalierbare parallele Rechenplattform
---
**Erklären Sie, was man unter Grid, Block und Kernel versteht?**
Grid
- Logische Sicht auf das gesamte Problem
- Partitioniert in Blöcke
- Blöcke können parallel berechnet werden
- Auf Multiprozessoren abgebildet
- Menge von Thread-Blöcken, die jeweils unabhängig voneinander ausgeführt werden können und somit parallel ausgeführt werden können
Block
- Blöcke sind Elemente innerhalb eines Grids und können parallel berechnet werden
- Blöcke sind unabhängig voneinander
- Skalierbarkeit gut
- Block wird von einem SM ausgeführt
- Auf Threads abgebildet
- Threads organisiert in einer Hierarchie von Blöcken
- Menge von Threads, die untereinander kooperieren durch Barrieresynchronisation und Zugriff auf blockeigenen Speicherplatz
- Maximal 1024 Threads pro Block auf aktuellen GPUs
Kernel
- Ein Eintrag in einem Block
- Berechnet einen Grid
- Aufgerufen aus einem seriellen Hauptprogramm
- Einfache Funktion oder vollständiges Programm
- Parallele Ausführung über eine Reihe von parallelen Threads
Jeder Thread berechnet ein oder mehrere Ergebniselemente
Thread
- ausgeführt von einem SP
- Threads eines Blocks führen gleichzeitig aus und können bei einer Barriere synchronisieren
<img src="https://i.imgur.com/nycrLib.png" alt="Visualisierung Grid, Block, Kernel" width="60%" style="display: block; margin-left: auto; margin-right: auto;">
---
**Wie haben sich GPUs über die Zeit entwickelt?**
- Mehr programmierbar
- Feste Funktiosnlogik durch programmierbare Prozessoren ersetzt
- Präziser
- Ganzzahl $\Rightarrow$ Festkomma, Übergang zu $\Rightarrow$ Single precision $\Rightarrow$ Double precision
- Besser skalierbar
- Vereinheitlichung verschiedener programmierbare Pipline-Elemente aus verschiedenen programmierbaren Prozessoren
- Verbesserter Load Balancing
- Vollständig programmierbar
- Unterstützung für allgemeine (general-purpose) Programmiersprachen wie C und C++
---
**Beschreiben Sie den Aufbau einer logischen Grafikpipline.**
<img src="https://i.imgur.com/d9url8v.png" alt="Logischen Grafikpipline" width="70%" style="display: block; margin-left: auto; margin-right: auto;">
- Vereinheitlichung verschiedener programmierbare Pipline-Elemente aus verschiedenen programmierbaren Prozessoren
- Keine verschiedenen Recheneinheiten, sondern alles in einem
- "Unified processor array"
<img src="https://i.imgur.com/zgMcs7j.png" alt="Vereinheitlichung Pipeline-Elemente" width="60%" style="display: block; margin-left: auto; margin-right: auto;">
---
**Welche GPU Programming Interfaces gibt es?**
- OpenGL
- Plattformübergreifende API für das Schreiben von Anwendungen zur Erstellung von 2D und 3D-Computergrafiken
- Direct3D
- Microsoft DirectX Multimedia-Schnittstelle
- Vulkan
- Plattformübergreifende 3D-Grafik- und Rechen-API mit geringem Overhead
- CUDA
- (Compute Unified Device Architecture) - C/C++ Erweiterung von NVIDIA
- OpenCL
- C-Framework zum Schreiben von Programme, die auf heterogenen Plattformen bestehend aus CPUs, GPUs und anderen Prozessoren ausgeführt werden können
- OpenACC & OpenMP
- Direktiven baiserte Interfaces
---
**Wie kann man sich vorstellen, wie CUDA parallele Berechnungen durchführt**?
<img src="https://i.imgur.com/lhxQbCt.png" alt="CUDA Visualisierung" width="80%" style="display: block; margin-left: auto; margin-right: auto;">
---
Erklären Sie die Funktionalität von `cudaFree` inklusive Argumente.
```c
cudaError_t cudaFree(void *devPointer)
```
- Gibt den Speicher im Device frei
- `*devPointer`: Device Pointer auf freizugebenden Speicher
---
**Erklären Sie die Funktionalität von `cudaGetErrorString` inklusive Argumente.**
```c
const char* cudaGetErrorString(cudaError_t error)
```
- Gibt die Beschreibung als String für einen Fehlercode zurück
- `error`: Fehlercode zur Umwandlung in einen String
---
### Mittel
---
**Welche Arten von Speicher gibt es?**
Globaler Speicher
- R/W von Host und Device
- R/W per Grid
- Hohe Latenz und Bandbreite
- Groß
Geteilter Speicher
- R/W von Device
- R/W per Block
- Niedrige Latenz
Register
- Wie Register auf der CPU
- Zugriff pro Thread
- Zugriff in einem Zyklus
- Klein
Constant
Local
---
**Erklären Sie den Begriff Warp.**
- Streaming-Multiprozessor erstellt, verwaltet, plant und führt Threads in Gruppen von 32 parallelen Threads (auch "Warps" genannt) genannt
---
**Was versteht man unter Warp Scheduling?**
- TODO
- Latenz verbergen
---
**Wie wird ein Kernel ausgeführt?**
- Beim Booten, nach Hardwarecheck ausgeführt
- Kernel wird mithilfe eines Bootloaders in Speicher geladen und gestartet
- initialisiert die Geräte vollständig und startet einen weiteren Prozess.
(wikipedia) reicht das?
---
**Welchen Arten von Kompatibilität gibt es?**
Binäre Kompatibilität:
- Binärcode ist architekturspezifisch (d. h. rechenfähig)
- Garantierte Kompatibilität zwischen kleineren Revisionen
PTX-Kompatibilität:
- Einige Befehle werden nur auf Geräten mit höherer Rechenleistung unterstützt
Anwendungskompatibilität:
- Die Anwendung muss Binär- oder PTX-Code laden, der mit der Rechenleistung des Geräts kompatibel ist.
- Um Code auf zukünftigen Architekturen auszuführen, für die noch keine Binärdateien erstellt werden können, muss eine Anwendung PTX-Code laden, der just-in-time kompiliert wird
C/C++ Kompatibilität:
- Vollständige C++-Unterstützung für Host-Code
- Nur eine Teilmenge von C++ ist auf Devicecode unterstützt
64-Bit-Kompatibilität:
- 64-Bit-Gerätecode wird nur mit 64-Bit-Hostcode unterstützt
- 32-Bit-Gerätecode wird nur mit 32-Bit-Hostcode unterstützt
---
**Welche CUDA-Erweiterungen zu funktionalen C-Deklarationen gibt es?**
`__device__`: Aufrufbar von Device und wird im Device ausgeführt
`__global__`: Aufrufbar von Host und wird im Device ausgeführt
`__host__`: Aufrufbar von Host und wird im Host ausgeführt
`__constant__`: Aufrufbar von Device
---
**Erklären Sie die Funktionalität von `cudaDeviceSynchronize` inklusive Argumente.**
```c
cudaError_t cudaDeviceSynchronize(int *device)
```
- wird gewartet, bis das Rechengerät fertig ist
- Blockiert, bis das Gerät alle vorhergehenden angeforderten Aufgaben abgeschlossen hat
- `*device`: Device Pointer auf freizugebenden Speicher
---
**Erklären Sie die Funktionalität von `syncthreads` inklusive Argumente.**
```c
void __syncthreads()
```
- Barrier Synchronisation
- Wartet, bis alle Threads im Thread-Block diesen Punkt erreicht haben und alle globalen und gemeinsamen Speicherzugriffe, die von diesen Threads vor `__syncthreads()` vorgenommen wurden, für alle Threads im Block sichtbar sind
- Hilft, potenzielle Gefahren durch Lesen-nach-Schreiben, Schreiben-nach-Lesen oder Schreiben-nach-Schreiben zu vermeiden.
---
### Schwer
---
**Worin unterscheiden sich SIMD und SIMT?**
SIMD:
- Wendet eine Anweisung auf mehrere Datenspuren an
- Erfordert, dass die Software die Datenparallelität explizit in jedem Vektorbefehl ausdrückt
SIMT:
- Wendet eine Anweisung auf mehrere unabhängige Threads in parallel
- Ermittelt Datenparallelität zwischen unabhängigen Threads
- Programme und Anweisungen beschreiben das Verhalten eines einzelnen unabhängigen Threads
SIMT kombiniert die Leistung von SIMD mit der Produktivität von Multithreading
---
**Wie ist ein Kompilierungs Workflow aufgebaut und welche Arten gibt es?**
<img src="https://i.imgur.com/cbGtwAY.png" alt="Compilation Workflow" width="70%" style="display: block; margin-left: auto; margin-right: auto;">
(PTX = Parallel Thread Execution)
Offline:
- Compiler trennt Host- von Devicecode
- Kompiliert Gerätecode in Assemblerform (PTX) und/oder Binärform (Cubin-Objekt)
- Modifizierung des Host-Codes durch Ersetzen der <<<>>>-Syntax (Cuda-Kernel-Aufruf) durch Laufzeit-Funktionsaufrufe
- Ausgabe des modifizierten Host-Codes entweder als C-Code, der von einem anderen Tool kompiliert werden kann einem anderen Tools kompiliert werden kann oder als Objektcode
Just-in-time:
- PTX-Code, der von einer Anwendung zur Laufzeit geladen wird, wird vom Gerätetreiber weiter zu Binärcode kompiliert
- Erhöht die Ladezeit der Anwendung
- Ermöglicht der Anwendung die Nutzung der neuesten Compiler-Verbesserungen
- Gerätetreiber speichert Kopie des Binärcodes im Cache
Sign in with Wallet
Connect another wallet