💾Speicherzugriffe – wo die Leistung entschieden wird

Die meisten Kernels warten mehr auf Daten, als sie rechnen. Wer Zugriffe bündelt, Bank-Konflikte vermeidet und Kopien überlappt, holt oft ein Vielfaches an Geschwindigkeit heraus.

🏔️Speicherhierarchie

Register und Shared Memory liegen im SM, L2 und globaler Speicher sind für die ganze GPU da. Latenzen sind als Größenordnungen gekennzeichnet.
▲ schneller, kleiner, näher am Rechenwerkgrößer, langsamer ▼

Register

Sichtbar für
je Thread
Größe
64 K × 32 Bit = 256 KB je SM
Latenz
≈ 1 Takt Größenordnung
Hinweis
schnellster Speicher; max. 255 Register pro Thread

Latenzen sind grobe Größenordnungen und hängen stark von Generation, Takt und Zugriffsmuster ab (vgl. Mikrobenchmark-Studien wie Jia et al., „Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking“, 2018). Größen: NVIDIA-Whitepaper bzw. CUDA Programming Guide.

🚚Coalescing: Zugriffe eines Warps bündeln

Wähle ein Zugriffsmuster und sieh, wie viele 32-Byte-Sektoren und 128-Byte-Cache-Lines ein einzelner Warp-Befehl anfordert.
Zugriffsmuster
Elementgröße (Byte)
int i = blockIdx.x * blockDim.x + threadIdx.x;
x = a[i];
🧪 Modell: 32-Byte-Sektoren, 128-Byte-Lines
32-B-Sektoren
4
128-B-Lines
1
übertragen / benötigt
128 / 128 B
Effizienz
100 %
Warp: 32 Lanes → Adresse (Byte)
0
0
1
4
2
8
3
12
4
16
5
20
6
24
7
28
8
32
9
36
10
40
11
44
12
48
13
52
14
56
15
60
16
64
17
68
18
72
19
76
20
80
21
84
22
88
23
92
24
96
25
100
26
104
27
108
28
112
29
116
30
120
31
124
Betroffene Cache-Lines (128 B = 4 Sektoren à 32 B)
Byte 0
✓
✓
✓
✓

Liegen die Adressen eines Warps dicht beieinander, reichen wenige Transaktionen (coalesced). Bei großem Stride holt die Hardware ganze 32-Byte-Sektoren, obwohl nur 4 Byte davon gebraucht werden – die Bandbreite geht verloren. Abhilfe: Datenlayout „Structure of Arrays“ statt „Array of Structures“, oder Daten gemeinsam über Shared Memory umsortieren.

🏦Shared Memory und Bank-Konflikte

Shared Memory ist in 32 Bänke à 4 Byte aufgeteilt. Jede Bank liefert pro Takt ein Wort – greifen mehrere Lanes auf verschiedene Wörter derselben Bank zu, wird serialisiert.
Zugriff
Zeilenbreite des Tiles
__shared__ float tile[32][32];
float x = tile[threadIdx.x][0];   // Spalte lesen
🧪 Modell: 32 Bänke à 4 Byte, 32-Bit-Zugriffe
Konfliktgrad
32-fach
Durchgänge (statt 1)
32
Lanes mit Broadcast
0
32 Bänke – jede Spalte zeigt die Wörter, die in diesem Zugriff aus der Bank gelesen werden
0
0
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31

Zahl = Lane · „×n“ = n Lanes lesen dasselbe Wort (Broadcast, kostenlos) · rot umrandet = zusätzlicher Durchgang wegen Konflikt.

Bank = (Wortadresse) mod 32. Liest ein Warp die Spalte eines float tile[32][32], liegen alle 32 Elemente in derselben Bank (Abstand 32 Wörter) → 32 Durchgänge. Mit einer Füllspalte (tile[32][33]) verschiebt sich jede Zeile um eine Bank → konfliktfrei. Allgemein gilt für Stride s: Konfliktgrad = ggT(s, 32).

🚧__syncthreads(): Barriere im Block

Warp 0
0
Warp 1
1
Warp 2
2
Warp 3
3
schreibt sdata[tid]__syncthreads()liest sdata[tid ± k]

Schnelle Warps warten an der roten Linie, bis der letzte Warp des Blocks angekommen ist.

Wichtig: __syncthreads() muss von allen Threads des Blocks erreicht werden – in einem if-Zweig, den nur manche Threads nehmen, führt es zu Hängern oder undefiniertem Verhalten. Es synchronisiert nur innerhalb eines Blocks, nicht blockübergreifend.

📦Speicher anlegen und kopieren

1float *h = (float*)malloc(bytes), *d;
2cudaMalloc(&d, bytes);                               // VRAM reservieren
3cudaMemcpy(d, h, bytes, cudaMemcpyHostToDevice);     // explizit kopieren
4kernel<<<grid, block>>>(d, n);
5cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost);     // zurück (wartet auf Kernel)
6cudaFree(d); free(h);
✅ Vorteile
  • volle Kontrolle, vorhersagbar
  • Kopien sichtbar im Profiler
⚠️ Nachteile
  • zwei Zeiger je Puffer (h/d)
  • Kopien vergisst man leicht – oder kopiert zu oft

🛤️Streams und asynchrone Kopien

Zerlege die Arbeit in Teile und verteile sie auf mehrere Streams: Kopieren und Rechnen überlappen sich.
Kopier-Engines der GPU
Absetzreihenfolge
🧪 vereinfachtes Zeitmodell, Zeiteinheiten frei gewählt
Kopier-Engine H2D
H2D0
H2D1
H2D2
H2D3
Rechnen (SMs)
Kernel0
Kernel1
Kernel2
Kernel3
Kopier-Engine D2H
D2H0
D2H1
D2H2
D2H3
seriell
100
mit Streams
55
Beschleunigung
1.82×

Innerhalb eines Streams läuft alles der Reihe nach, verschiedene Streams dürfen sich überlappen. So kopiert die GPU Teil 1, während sie Teil 0 schon rechnet. Mit nur einer Kopier-Engine blockiert die Rückkopie von Teil 0 die Hinkopie von Teil 1, wenn Teil für Teil abgesetzt wird – dann hilft die Reihenfolge „erst alle Kopien hin“.

Überlappung mit Streams
1// Host-Speicher muss „pinned“ sein, sonst ist cudaMemcpyAsync nicht asynchron
2cudaMallocHost(&h_buf, bytes);
3cudaStream_t s[NSTREAMS];
4for (int k = 0; k < NSTREAMS; ++k) cudaStreamCreate(&s[k]);
5 
6for (int c = 0; c < CHUNKS; ++c) {
7    cudaStream_t st = s[c % NSTREAMS];
8    size_t off = c * chunkElems;
9    cudaMemcpyAsync(d_buf + off, h_buf + off, chunkBytes, cudaMemcpyHostToDevice, st);
10    process<<<grid, block, 0, st>>>(d_buf + off, chunkElems);
11    cudaMemcpyAsync(h_buf + off, d_buf + off, chunkBytes, cudaMemcpyDeviceToHost, st);
12}
13cudaDeviceSynchronize();   // auf alle Streams warten
⚠️ Default-Stream
Ohne Stream-Angabe landet alles im Default-Stream. Im klassischen („legacy“) Verhalten synchronisiert er mit allen anderen Streams des Geräts – nvcc-Option --default-stream per-thread ändert das.