---

name: analyze-kernel-bottleneck description: > Systematically identify whether a GPU kernel is compute-bound, memory-bound, or latency-bound using roofline analysis, occupancy calculations, compute/load ratio per tile, and SASS instruction inspection. Produces a decision matrix for optimization strategy selection (cp.async, warp interleaving, tiling, double-buffering, or CuAssembler hand-tuning). license: MIT allowed-tools: Read Grep Glob Bash metadata: author: Philipp Thoss version: "1.0" domain: gpu-optimization complexity: advanced language: CUDA tags: gpu, roofline, occupancy, sass, tensor-core, bottleneck-analysis, compute-load-ratio locale: de source_locale: en source_commit: 11edabf5 fence_basis_commit: 11edabf5 translator: "Claude + human review" translation_date: "2026-05-03"

Kernel-Engpass analysieren

Systematisch identifizieren ob ein GPU-Kernel rechen-, speicher- oder latenzbegrenzt ist, indem die Baseline-Performance gemessen, auf der Roofline klassifiziert, Occupancy und Compute-Load-Ratio pro Tile berechnet, SASS-Instruktionsmix und Stall-Codes inspiziert, der Shared-Memory-Cliff geprueft und eine Entscheidungsmatrix angewendet wird um die richtige Optimierungsstrategie zu waehlen.

Wann verwenden

• Vor der Optimierung jedes CUDA-Kernels -- Baseline etablieren und Engpasstyp klassifizieren
• Nachdem eine erste funktionierende Version eines Kernels geschrieben wurde, um den Optimierungspfad zu identifizieren
• Wenn ein Kernel die Erwartungen relativ zum theoretischen Peak unterschreitet
• Bei der Entscheidung zwischen cp.async, groesseren Tiles oder algorithmischer Restrukturierung

Eingaben

• Erforderlich: Kompilierter Kernel (.cubin oder .cu-Quelle mit Build-Befehl)
• Erforderlich: Benchmark-Harness der den Kernel mit CUDA-Event-Timing startet
• Erforderlich: Problemdimensionen (z.B. M, N, K fuer GEMM; seq_len, heads, head_dim fuer Attention)
• Optional: Ziel-GPU-Architektur (Standard: GA104 / sm_86 / RTX 3070 Ti)
• Optional: Erwarteter Peak-Auslastungsprozentsatz fuer Vergleich
• Optional: Frueher Profiling-Daten (Nsight-Compute-Berichte)

Vorgehensweise

Schritt 1: Baseline-Performance messen

Den Kernel mit CUDA-Events (BenchTimer) ausfuehren, Zeit in Millisekunden aufzeichnen. Effektive Durchsatzmetriken berechnen:

1. Kompilieren des Kernels falls noch nicht gebaut:

   nvcc --cubin -arch=sm_86 -O2 -o kernel.sm_86.cubin kernel.cu
   nvcc -arch=sm_86 -O2 -o bench bench.cu -lcuda -I../../phase2/common
   

2. Ausfuehren mit repraesentativen Problemgroessen, sicherstellen dass Warmup-Laeufe der Messung vorausgehen:

   ./bench 4096 4096 4096
   

3. Aufzeichnen der Kernel-Zeit in ms aus CUDA-Events (nicht Wall-Clock).
4. Berechnen effektiver GFLOPS und effektiver Bandbreite:
   • GEMM: effective_gflops = (2 * M * N * K) / (time_ms / 1000) / 1e9
   • Bandbreitenbegrenzte Kernels: effective_bw = total_bytes / (time_ms / 1000) / 1e9
   • Flash Attention: effective_gflops = (4 * batch * heads * seq_len^2 * head_dim) / (time_ms / 1000) / 1e9

Erwartet: Baseline-Zahlen: Kernel-Zeit in ms, effektive GFLOPS und effektive Bandbreite.

Bei Fehler: Pruefen dass der Kernel ohne Fehler startet (CHECK_CU-Macro). Verifizieren dass Warmup-Laeufe der Messung vorausgehen. Sicherstellen dass Problemdimensionen gross genug sind um die GPU zu saettigen (kleine Probleme koennen am Launch-Overhead haengen bleiben).

Schritt 2: Auf der Roofline klassifizieren

Arithmetische Intensitaet berechnen und mit dem Maschinen-Balancepunkt vergleichen um den Kernel zu klassifizieren:

1. Arithmetische Intensitaet berechnen: AI = FLOPs / bytes_loaded_from_global_memory. Nur einzigartige Bytes aus DRAM zaehlen (nicht Shared Memory oder Register-Reuse).
2. Maschinen-Balancepunkt nachschlagen: balance = peak_compute / peak_bandwidth.
3. Klassifizieren: Wenn AI < balance, ist der Kernel speicherbegrenzt. Wenn AI > balance, ist der Kernel rechenbegrenzt.

GA104 (RTX 3070 Ti) Referenzwerte:

Resource	Peak	Unit
FP32 FFMA	21.7	TFLOPS
FP16 Tensor Core (HMMA)	174	TFLOPS
INT8 Tensor Core (IMMA)	696	TOPS
DRAM Bandwidth	608	GB/s
L2 Cache	4	MB
SMs	48

Abgeleitete Balancepunkte:

Precision	Balance Point (FLOP/byte)
FP32 FFMA	21700 / 608 = 35.7
FP16 TC	174000 / 608 = 286.2
INT8 TC	696000 / 608 = 1144.7

4. Erreichten Anteil berechnen: attained = effective_throughput / peak_throughput. Bei Speicherbegrenzung: effektive Bandbreite mit 608 GB/s vergleichen. Bei Rechenbegrenzung: effektive GFLOPS mit dem relevanten Peak vergleichen.

Erwartet: Klassifikation als rechen-, speicher- oder latenzbegrenzt (geringe Occupancy verursacht weder Rechen- noch Speichersaettigung) mit numerischer Begruendung.

Bei Fehler: Byte-Zaehlung erneut pruefen. Auf redundante Reads achten (z.B. 9x bei direkter conv2d ohne im2col). Wenn weder Rechnung noch Speicher gesaettigt sind, ist der Kernel wahrscheinlich latenzbegrenzt (siehe Schritt 3).

Schritt 3: Occupancy berechnen

Aktive Warps pro SM aus der Launch-Konfiguration und dem Ressourcenverbrauch ermitteln:

1. Ressourcenverbrauch extrahieren:

   nvcc --cubin -arch=sm_86 -O2 --resource-usage -o kernel.sm_86.cubin kernel.cu 2>&1 | grep -E 'registers|smem'
   

2. Aus Launch-Konfig: warps_per_block = threads_per_block / 32.
3. Bloecke/SM berechnen aus jedem begrenzenden Faktor:
   • Register-Limit: floor(65536 / (registers_per_thread * threads_per_block))
   • Smem-Limit: floor(available_smem_per_SM / smem_per_block) -- siehe Schritt 6 fuer Cliff
   • Warp-Limit: floor(48 / warps_per_block) (GA104 max: 48 Warps/SM)
   • Block-Limit: 16 Bloecke/SM max auf GA104
4. Tatsaechliche Bloecke/SM = min(register_limit, smem_limit, warp_limit, block_limit).
5. Aktive Warps/SM = blocks_per_SM * warps_per_block.
6. Schluessel-Schwellwert: 8 Warps/SM sind ausreichend fuer Latency Hiding auf GA104. Unterhalb 8 = strukturelles Problem das latenzbegrenztes Verhalten verursacht.

Erwartet: Occupancy-Tabelle die Bloecke/SM, aktive Warps/SM und den begrenzenden Faktor (Register, smem oder Warps) zeigt.

Bei Fehler: cuFuncSetAttribute fuer dynamisches Shared Memory pruefen. Verifizieren dass --resource-usage-Berichte mit der tatsaechlichen Launch-Konfiguration uebereinstimmen. Wenn Register-Anzahl unerwartet hoch ist, --maxrregcount=N versuchen um Register zu deckeln (Register-Spills gegen Occupancy tauschen).

Schritt 4: Compute-Load-Ratio pro Tile berechnen

Compute-Instruktionen und Load-Bytes pro K-Tile aus SASS zaehlen (nicht Quellcode):

1. Disassemblieren:

   cuobjdump -sass kernel.sm_86.cubin > kernel.sass
   

2. Compute-Instruktionen pro Tile zaehlen (die innere Schleife ueber ein K-Tile):
   • grep -c 'HMMA' kernel.sass -- FP16-Tensor-Core-Ops
   • grep -c 'IMMA' kernel.sass -- INT8-Tensor-Core-Ops
   • grep -c 'FFMA' kernel.sass -- FP32-Fused-Multiply-Add
3. Globale Loads pro Tile zaehlen:
   • grep -c 'LDG' kernel.sass -- Global-Memory-Loads
   • Mit Bytes pro Load multiplizieren (typischerweise 16 Bytes fuer LDG.128)
4. Ratio berechnen: compute_ops / load_ops pro Tile.
5. Klassifizieren mit dem cp.async-Entscheidungsschwellwert (aus gpu_reflections.md Insight 2):
   • Hoch (>20:1): cp.async ist netto-negativ; Warp-Interleaving versteckt bereits DRAM-Latenz. Auf algorithmische Aenderungen fokussieren. Referenz: Flash Attention hat 64 HMMA pro Tile = hohes Ratio, cp.async gemessen -5%.
   • Mittel (5-20:1): cp.async kann helfen, beide Pfade benchmarken.
   • Niedrig (<5:1): cp.async stark vorteilhaft; Loads dominieren und Async-Copy versteckt Latenz. Referenz: IGEMM hat 8 IMMA pro Tile = niedriges Ratio, cp.async gemessen +35%.

Erwartet: Compute-Load-Ratio mit Klassifikation (hoch/mittel/niedrig) und cp.async-Empfehlung.

Bei Fehler: Aus SASS-Disassembly zaehlen, nicht Quellcode -- der Compiler kann fusionieren, eliminieren oder Instruktionen umordnen. Sicherstellen dass nur Instruktionen innerhalb der inneren Schleife gezaehlt werden (die K-Tile-Iteration), nicht der ganze Kernel.

Schritt 5: SASS-Instruktionen inspizieren

Den vollstaendigen SASS-Instruktionsmix und Stall-Codes pruefen:

1. Disassemblieren (falls in Schritt 4 nicht erfolgt):

   cuobjdump -sass kernel.sm_86.cubin > kernel.sass
   

2. Schluesselinstruktionstypen zaehlen:

   grep -c 'HMMA.16816' kernel.sass      # FP16 Tensor Core
   grep -c 'IMMA.16816' kernel.sass      # INT8 Tensor Core
   grep -c 'FFMA' kernel.sass            # FP32 fused multiply-add
   grep -c 'LDGSTS' kernel.sass          # cp.async (global->shared)
   grep -c 'LDG' kernel.sass             # Global load
   grep -c 'STS' kernel.sass             # Shared store
   grep -c 'LDS' kernel.sass             # Shared load
   grep -c 'BAR.SYNC' kernel.sass        # Barrier synchronization
   grep -c 'SHFL' kernel.sass            # Warp shuffle (reductions)
   grep -c 'MUFU' kernel.sass            # Special function unit
   

3. Stall-Codes pruefen auf kritischen Instruktionen:

   grep 'HMMA' kernel.sass | head -5     # Expect S08 minimum (hardware constraint)
   grep 'IMMA' kernel.sass | head -5     # Compiler emits S04, reducible to S02 via CuAssembler
   grep 'FFMA' kernel.sass | head -5     # Check for S04 (reducible to S01 on independent FFMAs)
   

4. Optimierungsziele identifizieren:
   • HMMA-S08-Stalls: Hardware-Minimum auf Ampere, kann nicht reduziert werden. Anderswo fokussieren.
   • IMMA-S04-Stalls: Compiler ist konservativ. CuAssembler kann auf S02 verengen (gemessen 15-20% Gewinn).
   • FFMA-S04-Stalls: bei Unabhaengigkeit auf S01 reduzierbar via CuAssembler.
   • Exzessives BAR.SYNC: kann auf Ueber-Synchronisation zwischen Pipeline-Stufen hinweisen.

Erwartet: Instruktionsanzahltabelle und Stall-Code-Zusammenfassung mit identifizierten Optimierungszielen.

Bei Fehler: Sicherstellen dass die cuobjdump-Architektur mit dem Kernel-Kompilierungsziel uebereinstimmt (beide muessen sm_86 sein). Wenn SASS-Ausgabe leer ist, koennte das cubin korrupt sein -- neu kompilieren.

Schritt 6: Den Smem-Cliff pruefen

Ermitteln ob der Shared-Memory-Verbrauch den architekturspezifischen Occupancy-Cliff ueberschreitet:

1. Smem/Block lesen aus --resource-usage-Ausgabe (Schritt 3) oder cuobjdump --res-usage kernel.sm_86.cubin.
2. Mit Cliff-Schwellwert vergleichen:
   • GA104 (sm_86): 100 KB max smem/SM. Cliff bei 50 KB/Block.
   • Empirisch bestaetigt: 48 KB/Block -> 2 Bloecke/SM (gut), 56 KB/Block -> 1 Block/SM (2x Regression).
3. Wenn ueber Cliff (smem > 50 KB/Block):
   • Bloecke/SM faellt auf 1, aktive Warps fallen auf warps_per_block (typischerweise 4).
   • 2x Performance-Regression erwartet durch exponierte DRAM-Stalls.
4. Auswirkung von Double-Buffering pruefen: Double-Buffering verdoppelt smem-Verbrauch. Bei aktuellen 30 KB smem ist double-buffered = 60 KB, was den Cliff ueberschreitet. Bewerten ob der Async-Vorteil den Occupancy-Verlust ueberwiegt.
5. Aufzeichnen smem/Block, Bloecke/SM und ob der Cliff ueberschritten wird.

Erwartet: Smem/Block-Wert mit Bloecke/SM-Anzahl und expliziter Aussage ob der 50-KB-Cliff ueberschritten wird.

Bei Fehler: Wenn ueber Cliff und Occupancy der Engpass ist, muss sich die Optimierungsstrategie aendern: Tile-Groesse reduzieren um smem unter 50 KB zu bekommen, oder 1 Block/SM akzeptieren und mit hoeherer Compute-Load-Ratio pro Tile kompensieren (mehr Register-Reuse, laengere K-Tiles).

Schritt 7: Die Entscheidungsmatrix bauen

Befunde aus Schritten 2-6 zu einer Optimierungsstrategie synthetisieren:

Bedingung	Strategie
Speicherbegrenzt + niedrige Compute-Load-Ratio (<5:1) + smem unter Cliff	Software-Pipelining mit cp.async (LDGSTS). Globale Loads mit Compute ueberlappen.
Speicherbegrenzt + hohe Compute-Load-Ratio (>20:1) + 8+ Warps	Warp-Interleaving versteckt bereits Latenz. Auf algorithmische Aenderungen fokussieren: implicit GEMM, split-Q, im2col.
Rechenbegrenzt + FFMA-lastig	CuAssembler-Stall-Code-Verengung: S04 -> S01 auf unabhaengigen FFMAs.
Rechenbegrenzt + HMMA-lastig	S08 ist Hardware-Minimum, nicht reduzierbar. Tile-Reuse erhoehen (groessere M/N-Tiles, laengere K-Schleife).
Rechenbegrenzt + IMMA-lastig	CuAssembler: S04 -> S02 auf IMMA-Instruktionen (Compiler ist konservativ).
Latenzbegrenzt (niedrige Occupancy, weder gesaettigt)	Smem oder Register reduzieren um mehr Bloecke/SM zu erhalten. Ueber 8 Warps/SM kommen.
Smem ueber Cliff	Tile-Groesse reduzieren oder restrukturieren um smem/Block unter 50 KB zu bekommen (GA104).

1. Anwendbare Strategien rangieren nach erwartetem Gewinn, mit Compute-Load-Ratio und Occupancy-Daten.
2. Gewinn-Bereich schaetzen fuer jede Strategie basierend darauf wie weit der Kernel von der relevanten Decke entfernt ist.
3. Konflikte markieren: z.B. cp.async verdoppelt smem (kann Cliff ueberschreiten), groessere Tiles erhoehen Register-Druck (koennen Occupancy reduzieren).

Erwartet: Rangierte Liste empfohlener Optimierungen mit vorhergesagtem Gewinn-Bereich und potenziellen Konflikten.

Bei Fehler: Wenn kein klarer Gewinner auftaucht, Mikro-Benchmarks ausfuehren die jede Strategie isolieren (z.B. cp.async allein testen, reduzierte Tile-Groesse allein testen) um die tatsaechliche Wirkung vor dem Kombinieren zu messen.

Schritt 8: Befunde dokumentieren

Einen strukturierten Engpass-Bericht erzeugen:

1. Baseline: Kernel-Zeit, effektive GFLOPS, effektive Bandbreite, Problemdimensionen.
2. Roofline-Position: arithmetische Intensitaet, Klassifikation, erreichter Peak-Anteil.
3. Occupancy: Bloecke/SM, aktive Warps/SM, begrenzender Faktor.
4. Compute-Load-Ratio: Ratio-Wert, Klassifikation (hoch/mittel/niedrig), cp.async-Empfehlung.
5. SASS-Zusammenfassung: Instruktionsanzahltabelle, Stall-Code-Befunde, CuAssembler-Ziele.
6. Smem-Cliff: smem/Block, Bloecke/SM, Cliff-Status.
7. Empfehlung: rangierte Optimierungsstrategien mit Gewinn-Schaetzungen.

## Bottleneck Analysis Report: [kernel_name]

### Baseline
- Problem: [dimensions]
- Kernel time: [X] ms
- Effective GFLOPS: [Y] | Effective BW: [Z] GB/s

### Roofline Classification
- Arithmetic intensity: [AI] FLOP/byte
- Balance point: [BP] FLOP/byte ([precision])
- Classification: **[compute|memory|latency]-bound**
- Attained fraction: [X]% of peak

### Occupancy
| Resource | Per Block | Limit/SM | Blocks/SM |
|----------|-----------|----------|-----------|
| Registers | [N]/thread | 65536 | [B] |
| Shared mem | [X] KB | 100 KB (cliff: 50 KB) | [B] |
| Warps | [W] | 48 | [B] |
| **Limiting** | | | **[min(B)]** |
- Active warps/SM: [W] ([sufficient|insufficient] for latency hiding)

### Compute/Load Ratio
- Compute ops/tile: [N] [HMMA|IMMA|FFMA]
- Load bytes/tile: [N] bytes ([N] LDG x [N] bytes)
- Ratio: [X]:1 — **[high|medium|low]**
- cp.async recommendation: [beneficial|neutral|detrimental]

### SASS Instruction Mix
| Instruction | Count | Notes |
|-------------|-------|-------|
| HMMA.16816 | [N] | Stall: S08 (hardware min) |
| IMMA.16816 | [N] | Stall: S04 (reducible to S02) |
| FFMA | [N] | Stall: S04 (reducible to S01) |
| LDG | [N] | |
| LDGSTS | [N] | cp.async |
| BAR.SYNC | [N] | |

### Smem Cliff
- Smem/block: [X] KB — [under|over] 50 KB cliff
- Blocks/SM: [B] — [no occupancy loss|occupancy halved]

### Recommended Optimizations (ranked)
1. [Strategy] — estimated [X-Y]% gain
2. [Strategy] — estimated [X-Y]% gain
3. [Strategy] — estimated [X-Y]% gain

Erwartet: Vollstaendiger Markdown-Bericht der von einem Kernel-Optimizer-Agenten oder menschlichen Entwickler konsumierbar ist.

Bei Fehler: Mit unterschiedlichen Problemgroessen erneut ausfuehren (z.B. 1024, 2048, 4096, 8192) um zu bestaetigen dass Befunde nicht groessenspezifisch sind. Kleine Probleme koennen latenzbegrenzt erscheinen waehrend der echte Engpass im Massstab Speicherbandbreite ist.

Validierung

• Baseline mit CUDA-Events gemessen (nicht Wall-Clock)
• Roofline-Klassifikation ermittelt (rechen-/speicher-/latenzbegrenzt)
• Occupancy berechnet mit identifiziertem begrenzendem Faktor
• Compute-Load-Ratio pro Tile aus SASS berechnet
• SASS-Instruktionsmix und Stall-Codes dokumentiert
• Smem-Cliff gegen Architektur-Schwellwert geprueft
• Entscheidungsmatrix angewendet mit Strategie-Empfehlung
• Befunde in strukturiertem Bericht dokumentiert

Haeufige Stolperfallen

• Re-Read-Multiplikation: Direkte conv2d liest jedes Gewicht 9x ohne im2col, blaeht die Byte-Anzahl um 9x auf. Tatsaechliche einzigartige aus DRAM geladene Bytes nutzen, nicht gesamte Load-Instruktionen, beim Berechnen arithmetischer Intensitaet.
• FP16-Tensor-Core-Peak mit FP32-Peak verwechseln: FP16-TC-Peak ist 174 TFLOPS, FP32-FFMA-Peak ist 21,7 TFLOPS -- ein 8x-Unterschied. Mit dem falschen Peak wird Roofline-Klassifikation bedeutungslos.
• 64 KB als smem-Cliff statt 50 KB auf GA104 nutzen: GA104 (sm_86) hat 100 KB max smem/SM. Der Cliff ist bei 100/2 = 50 KB/Block, nicht 64 KB. Dies ist architekturspezifisch; andere GPUs unterscheiden sich.
• Warp-Interleaving beim Bewerten von cp.async ignorieren: 8 Warps mit langen Compute-Phasen (hohe Compute-Load-Ratio) verstecken bereits DRAM-Latenz durch Warp-Scheduling. cp.async in diesem Regime hinzuzufuegen fuegt smem-Druck und Barrier-Overhead ohne Nutzen hinzu (gemessen -5% auf Flash Attention).
• Instruktionen aus Quellcode statt SASS zaehlen: Der Compiler kann Operationen fusionieren, toten Code eliminieren, Schleifen anders unrollen oder Instruktionen umordnen. Immer aus cuobjdump -sass-Ausgabe zaehlen.
• Keine Warmup-Iterationen ausfuehren: Der erste Kernel-Launch enthaelt JIT-Compile-Overhead und Cold-Cache-Effekte. Immer 2-5 Warmup-Iterationen vor dem gemessenen Lauf ausfuehren.

Verwandte Skills

• pipeline-gpu-kernel -- Software-Pipelining mit cp.async implementieren wenn die Analyse einen speicherbegrenzten Kernel mit niedriger Compute-Load-Ratio identifiziert
• simulate-cpu-architecture -- ergaenzende Architekturanalyse fuer CPU-seitige Engpaesse in Host-Device-Workflows