🟩 Visuelle Erklärung

Eine GPU programmieren – visuell und interaktiv

Wie tausende Threads gleichzeitig rechnen: vom Streaming Multiprocessor über Warps und Speicherhierarchie bis zum eigenen CUDA-Kernel. Mit Grid-Visualisierer, Occupancy-Rechner, Coalescing- und Bank-Konflikt-Labor und einem Schritt-Debugger, der Reduktion und Matrixmultiplikation im Browser nachrechnet.

🧵 Grid & Threads ausprobieren🐞 Kernel durchsteppen
vecadd.cu – das „Hallo Welt“ der GPU
1__global__ void add(const float* a, const float* b, float* c, int n) {
2    int i = blockIdx.x * blockDim.x + threadIdx.x;   // „Wer bin ich?“
3    if (i < n) c[i] = a[i] + b[i];                   // ein Element pro Thread
4}
5 
6add<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n);    // Tausende Threads starten

🆚CPU oder GPU?

Eine CPU minimiert die Zeit für einen einzelnen Befehlsstrom (Latenz). Eine GPU maximiert die Menge an Arbeit pro Zeit (Durchsatz) – vorausgesetzt, es gibt genug unabhängige Arbeit.

🧠 CPU – latenzoptimiert

SteuerungALUSteuerungALUSteuerungALUSteuerungALUgroßer Cache

Wenige, komplexe Kerne: Sprungvorhersage, Out-of-Order, große Caches – ein einzelner Befehlsstrom soll schnell fertig werden.

🟩 GPU – durchsatzoptimiert

L2-Cache + breite Speicheranbindung

Viele einfache Rechenspuren in Multiprozessoren (SMs): Wartezeiten werden durch tausende Threads überdeckt statt durch große Caches.

⏱️ Spielmodell: wer ist schneller fertig?

Alle Werte sind einstellbare Annahmen, keine Messwerte. CPU: 1 Zeiteinheit je Aufgabe; GPU: jede Spur ist langsamer, aber es gibt viel mehr davon.

CPU
256 ZE
GPU
8 ZE
Sieger
GPU
Auslastung GPU
50 %

Bei wenigen Aufgaben gewinnt die CPU (die meisten GPU-Spuren bleiben leer). Erst bei massiver Parallelität zahlt sich die GPU aus – ohne die Kosten für das Kopieren der Daten über PCIe, die hier noch fehlen.

🧭Kapitel

🧩

Architektur

GPC, TPC, SM, CUDA-Kerne, Tensor Cores, Warp-Scheduler, Caches und Generationen von Pascal bis Blackwell.

🧵

CUDA-Programmiermodell

Grid, Block, Thread live visualisiert, Warps à 32, Warp-Divergenz und Occupancy-Rechner.

💾

Speicherzugriffe

Coalescing, Bank-Konflikte, cudaMemcpy, Unified Memory, Streams als Zeitleiste und __syncthreads.

🐞

Schritt-Debugger

Vektoraddition, Reduktion (atomic → shared → Shuffle) und Matrixmultiplikation naiv vs. Tiling.

📜

Kommentierter Code

Vollständige CUDA-Programme mit Fehlerprüfung, Shuffle-Reduktion, Tiling, Streams und Unified Memory.

🧰

Ökosystem & WebGPU

nvcc, PTX/SASS, Bibliotheken, Python-Welt, Profiling, Container – und ein Compute-Shader live im Browser.

🗺️Der Weg eines Kernels

  Host (CPU)                                   Device (GPU)
  ─────────                                    ────────────
  cudaMalloc   ──────────────────────────────▶ Speicher im VRAM reservieren
  cudaMemcpy   ═══ PCIe / NVLink ════════════▶ Eingaben kopieren (Host → Device)
  kernel<<<grid, block>>>  ──────────────────▶ Grid aus Blöcken; jeder Block landet auf einem SM
                                                 └─ Block wird in Warps à 32 Threads zerlegt
                                                    └─ Warp-Scheduler gibt Befehle für 32 Lanes aus
  (CPU läuft weiter – der Start ist asynchron)
  cudaMemcpy   ◀══ PCIe / NVLink ═════════════ Ergebnis zurück (wartet auf den Kernel)