🧵Das CUDA-Programmiermodell
Ein CUDA-Programm besteht aus Host-Code (CPU) und Kernels (GPU). Ein Kernel wird einmal geschrieben und von tausenden Threads gleichzeitig ausgeführt – jeder Thread findet über seine Indizes heraus, welches Datenelement ihm gehört.
🖥️Host und Device
Host (CPU + RAM) Device (GPU + VRAM)
┌──────────────────┐ PCIe ┌──────────────────────────┐
│ main() │ ════════▶ │ Grid │
│ cudaMalloc │ │ ┌───────┐┌───────┐ │
│ cudaMemcpy │ │ │Block 0││Block 1│ … │
│ kernel<<<…>>> │ │ │▓▓▓▓▓▓▓││▓▓▓▓▓▓▓│ │
│ cudaMemcpy │ ◀════════ │ └───────┘└───────┘ │
└──────────────────┘ │ ▓ = Thread │
└──────────────────────────┘__syncthreads() abstimmen · Thread = eine Ausführung des Kernels mit eigenen Registern.1// Qualifizierer: wo läuft eine Funktion, von wo wird sie aufgerufen?2__global__ void kernel(float* x, int n); // läuft auf der GPU, Aufruf vom Host3__device__ float helper(float v); // läuft auf der GPU, Aufruf nur von GPU-Code4__host__ __device__ float both(float v); // für beide Seiten übersetzt56// Start-Konfiguration: <<<Blöcke, Threads pro Block, dyn. Shared Memory, Stream>>>7dim3 block(256); // bis 1024 Threads pro Block8dim3 grid((n + block.x - 1) / block.x); // genug Blöcke für alle n Elemente9kernel<<<grid, block, 0, stream>>>(d_x, n);1011cudaError_t err = cudaGetLastError(); // Startfehler (z. B. zu viele Threads)12cudaDeviceSynchronize(); // auf das Ende warten (Laufzeitfehler)
🔢Grid-, Block- und Thread-Visualisierer
int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) out[i] = f(in[i]); // Start: kernel<<<3, 16>>>
Graue Threads haben kein Element (i ≥ n). Rot = Element ohne Thread (Grid zu klein, keine Grid-Stride-Schleife). Blassere Zellen = zweiter oder späterer Schleifendurchlauf. Abwechselnd helle/dunkle Threads = unterschiedliche Warps (je 32).
👯Warps: 32 Threads im Gleichschritt
Aufteilung
Der SM zerlegt jeden Block in Warps zu je 32 Threads (warpSize): Threads 0–31 = Warp 0, 32–63 = Warp 1 usw. Blockgrößen sollten daher Vielfache von 32 sein – ein Block mit 48 Threads belegt 2 Warps, der zweite ist nur halb voll.
Ausführung
Der Warp-Scheduler gibt einen Befehl für den ganzen Warp aus (SIMT). Lanes, die gerade nicht mitmachen sollen, werden maskiert. Warpwechsel kosten nichts, weil alle Register fest zugeteilt sind.
Warp-Funktionen
Threads eines Warps können direkt Register austauschen: __shfl_down_sync, __ballot_sync, __any_sync, __syncwarp – ohne Shared Memory (siehe Reduktion im Debugger).
🔀Warp-Divergenz
if (threadIdx.x % 2 == 0) {// if-Zweig: 4 Befehle} else {// else-Zweig: 4 Befehle}// wieder alle 32 Lanes gemeinsam
Divergenz gibt es nur innerhalb eines Warps. Ist die Bedingung für alle 32 Lanes gleich (z. B. abhängig von threadIdx.x / 32), läuft jeder Warp nur einen Pfad. Seit Volta hat jeder Thread einen eigenen Programmzähler (Independent Thread Scheduling) – die beiden Pfade werden dennoch nicht gleichzeitig ausgeführt.
📊Occupancy-Rechner
Hohe Occupancy hilft, Speicherlatenz zu verbergen – ist aber kein Selbstzweck: Manche schnelle Kernels (z. B. getiltes GEMM) laufen bewusst mit wenigen Warps und vielen Registern. Die Register-Zahl zeigt nvcc --resource-usage; begrenzen lässt sie sich mit __launch_bounds__ oder -maxrregcount.
Vereinfachung: angenommen wird die größte Shared-Memory-Aufteilung (Carveout) des SM; der Treiber wählt sie in echt je nach Kernel.