CUDA Concepts and GPU Profiling

A GPU is fast only when you keep its thousands of cores fed with data, and most slow training or serving code is slow because of memory traffic, small kernels or the CPU getting in the way - not because of the math. This note covers the CUDA execution model (threads, warps, blocks and SMs), the memory hierarchy, why kernel fusion matters, what Triton and torch.compile do, and a profiling workflow with torch.profiler, Nsight Systems and Nsight Compute for finding the real bottleneck.

The Execution Model

flowchart TD
    G["🧩 Grid<br/>one kernel launch"] --> B1["📦 Block 0"]
    G --> B2["📦 Block 1"]
    G --> B3["📦 ... Block N"]
    B1 --> W1["🧵 Warp = 32 threads<br/>same instruction"]
    B1 --> W2["🧵 Warp"]
    B1 -.->|"scheduled onto"| SM["🖥️ Streaming multiprocessor<br/>registers · shared memory · tensor cores"]
    B2 -.-> SM2["🖥️ Another SM"]

    style G fill:#d8dfe8,stroke:#b0bac8
    style SM fill:#e8e0d4,stroke:#c8b89a
    style SM2 fill:#e8e0d4,stroke:#c8b89a

Three consequences worth remembering:

• Divergence costs time. If threads in a warp take different branches, the warp runs both paths with some threads masked off.
• Occupancy - how many warps are resident per SM relative to the maximum (64 on Hopper) - limits how well latency is hidden. It is capped by registers per thread and shared memory per block; high occupancy helps memory-bound kernels, though well-tuned matmul kernels often run at low occupancy on purpose, using registers and shared memory heavily instead.
• Small kernels underuse the GPU. A kernel with fewer blocks than there are SMs leaves processors idle, and every launch costs a few microseconds of CPU-side overhead.

---

The Memory Hierarchy

Level	Size (H100)	Speed	Scope
Registers	64K 32-bit registers per SM (256 KB)	Fastest	One thread
Shared memory / L1	Up to 228 KB shared memory per SM (256 KB combined L1 + shared)	Very fast, on-chip	One thread block
L2 cache	50 MB	Fast, on-chip	Whole GPU
HBM (global memory)	80 GB at 3.35 TB/s (SXM)	Slowest on the GPU, yet far faster than CPU memory	Whole GPU

Data used once must come from HBM; the game is to load each byte from HBM as few times as possible and reuse it from registers and shared memory. Tiled matrix multiplication loads tiles of the inputs into shared memory and reuses each loaded value for many multiply-adds. FlashAttention applies the same idea to attention: it computes the softmax in tiles in shared memory and never writes the T x T score matrix to HBM (GPU Memory & Hardware).

Arithmetic intensity, from the roofline model: FLOPs per byte moved. The H100's ridge point is about 295 FLOP/byte for dense BF16.

Operation (BF16)	FLOPs / byte	Bound
4096 x 4096 @ 4096 x 4096 matmul	≈ 1,365	Compute
One token's 1 x 4096 @ 4096 x 4096 (decode, batch 1)	≈ 1	Memory bandwidth
Elementwise ops (GELU, add, residual, most norms)	≈ 1 or less	Memory bandwidth

---

Kernel Fusion

Every unfused elementwise operation reads its inputs from HBM and writes its output back. A chain of them moves the same data repeatedly.

Worked example. Three elementwise operations on a 4096 x 16384 BF16 activation (134 MB):

• Unfused: each reads and writes the tensor - 3 x 2 x 134 MB ≈ 805 MB, about 0.24 ms at 3.35 TB/s
• Fused into one kernel: one read and one write - 268 MB, about 0.08 ms

A 3x speedup with identical math - and fusion also removes two kernel launches. This is what FlashAttention, fused optimizers, fused RMSNorm and torch.compile all exploit.

Triton is a Python-embedded language for writing GPU kernels at the level of blocks of data rather than individual threads; the compiler handles the thread-level details. torch.compile's default backend (Inductor) generates Triton kernels that fuse chains of operations automatically. A fused SwiGLU gate (silu(gate) * up), which otherwise needs two kernels and five tensor passes, in one kernel with three:

import torch
import triton
import triton.language as tl

@triton.jit
def swiglu_kernel(gate_ptr, up_ptr, out_ptr, n, BLOCK: tl.constexpr):
    pid = tl.program_id(axis=0)                      # which block of elements this program handles
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < n                                  # guard the ragged last block
    g = tl.load(gate_ptr + offs, mask=mask).to(tl.float32)
    u = tl.load(up_ptr + offs, mask=mask).to(tl.float32)
    out = g * tl.sigmoid(g) * u                      # silu(g) * u, computed in registers
    tl.store(out_ptr + offs, out.to(out_ptr.dtype.element_ty), mask=mask)

def swiglu(gate: torch.Tensor, up: torch.Tensor) -> torch.Tensor:
    assert gate.is_contiguous() and up.is_contiguous()
    out = torch.empty_like(gate)
    n = gate.numel()
    swiglu_kernel[(triton.cdiv(n, 1024),)](gate, up, out, n, BLOCK=1024)
    return out

This kernel follows the pattern of Triton's official tutorials but was not GPU-run for this note. In practice, try torch.compile first - it usually produces this fusion for you - and write a custom kernel only when profiling shows a hot spot it misses.

---

Timing GPU Code Correctly

Kernel launches are asynchronous: the CPU queues work and returns immediately. Timing with time.time() around GPU code measures how long it took to queue the work, not to run it.

import torch

start, end = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
for _ in range(3):                  # warm up: compilation, caches, allocator
    step()
torch.cuda.synchronize()
start.record()
for _ in range(20):
    step()
end.record()
torch.cuda.synchronize()            # wait for the GPU before reading the timer
print(f"{start.elapsed_time(end) / 20:.2f} ms per step")

The flip side: an accidental synchronization - .item(), .cpu(), printing a tensor, torch.cuda.synchronize() inside the loop - stalls the CPU until the GPU drains, so the GPU then idles while the CPU prepares the next launch. Hidden syncs in a training loop are a common cause of low utilization.

Launch overhead. When a step is many tiny kernels (small batches, decode), CPU launch time dominates. CUDA Graphs record a sequence of launches once and replay it with a single launch; torch.compile(mode="reduce-overhead") uses them, and serving engines such as vLLM capture CUDA graphs for decode.

---

A Profiling Workflow

flowchart TD
    M["📏 Measure step time + MFU<br/>is there a problem?"] --> TP["🔬 torch.profiler<br/>which ops, CPU vs GPU time"]
    TP --> NS["🧭 Nsight Systems<br/>timeline: gaps, syncs, data loader,<br/>NCCL overlap"]
    NS --> Q{"Where is the time?"}
    Q -->|"GPU idle between kernels"| CPU["⚙️ CPU-bound: data loading,<br/>launch overhead, syncs<br/>→ more workers, CUDA graphs, remove syncs"]
    Q -->|"many small elementwise kernels"| FU["🔗 Memory-bound chains<br/>→ torch.compile, fused kernels"]
    Q -->|"long collectives"| CO["🌐 Communication<br/>→ overlap, bucket sizes, topology"]
    Q -->|"one big kernel dominates"| NC["🔎 Nsight Compute<br/>roofline, memory vs compute,<br/>occupancy for that kernel"]

    style M fill:#d8dfe8,stroke:#b0bac8
    style Q fill:#e8e0d4,stroke:#c8b89a
    style NC fill:#ddd8e4,stroke:#b8b0c8

1. Start with numbers. Step time, tokens per second and MFU - achieved model FLOPs per second divided by peak (PyTorch for LLMs). Well-tuned large dense training reaches roughly 35-50% MFU on H100s; far below that, something specific is wrong.

2. torch.profiler - per-operator CPU and CUDA time, memory, and a Chrome/Perfetto trace. Sort by CUDA time to find the operators that matter; compare CPU and CUDA timelines to spot a CPU-bound loop.

   from torch.profiler import profile, ProfilerActivity, schedule
   
   with profile(activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
                schedule=schedule(wait=1, warmup=1, active=3), record_shapes=True) as prof:
       for _ in range(5):
           step()
           prof.step()
   print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=15))
   prof.export_chrome_trace("trace.json")          # open in https://ui.perfetto.dev
   

3. Nsight Systems (nsys profile --trace=cuda,nvtx,osrt -o run python train.py) - a whole-system timeline of CPU threads, kernel launches, memory copies and NCCL communication. It answers "why is the GPU idle?" - waiting on the data loader, on a sync, or on communication that isn't overlapped with compute. Add NVTX ranges (torch.cuda.nvtx.range_push("attention")) to label regions.

4. Nsight Compute (ncu --set full -o kernel python train.py, filtered to the kernel you care about) - per-kernel detail: achieved memory bandwidth and FLOP/s against the roofline, occupancy, and the reasons warps stall. Use it last, on the one kernel that dominates.

---

Common Bottlenecks and Fixes

Symptom	Likely cause	Fix
GPU utilization low, gaps between kernels	Data loading on the CPU, synchronizations, Python overhead	More DataLoader workers, pinned memory, prefetch; remove .item() in the loop; CUDA graphs
Many short elementwise kernels	Unfused memory-bound operations	torch.compile; fused kernels (RMSNorm, optimizer, attention)
Matmuls slow for their size	Odd shapes not multiples of 8 or 64, FP32 instead of BF16	Pad dimensions; use BF16 or FP8 tensor cores
Attention dominates and memory explodes at long context	Unfused attention materializing T x T	FlashAttention / SDPA backends
Long NCCL kernels, compute waiting	Communication not overlapped	Overlap gradient all-reduce with backward; tune bucket sizes; check interconnect topology (Distributed Training at Scale)
Decode tokens/s low at batch 1	Memory-bandwidth bound by design	Batch more requests; quantize weights and KV cache (Inference & Serving)

---

Check Yourself

Exercises

Study Notes

Must-know:

• Kernel → grid of blocks → warps of 32 threads on SMs (132 on H100 SXM); divergence serializes; occupancy hides latency
• Memory hierarchy: registers → shared memory/L1 (228 KB/SM) → L2 (50 MB) → HBM (80 GB, 3.35 TB/s); reuse on-chip, minimize HBM traffic
• Arithmetic intensity vs ridge point (~295 FLOP/byte on H100 BF16): big matmuls compute-bound; elementwise ops and batch-1 decode memory-bound
• Fusion cuts memory traffic and launches; FlashAttention, fused norms and optimizers, torch.compile (Inductor → Triton)
• Launches are async: time with CUDA events after synchronizing; hidden syncs (.item()) starve the GPU; CUDA graphs cut launch overhead
• Workflow: step time and MFU → torch.profiler → Nsight Systems timeline → Nsight Compute for the one hot kernel

References

• NVIDIA, CUDA C++ Programming Guide (2026) and Hopper Tuning Guide (2026)
• NVIDIA, NVIDIA Hopper Architecture In-Depth (2022)
• Williams, Waterman and Patterson, Roofline: An Insightful Visual Performance Model (CACM, 2009)
• Tillet, Kung and Cox, Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations (MAPL 2019); Triton fused-softmax tutorial
• Ansel et al., PyTorch 2: Faster Machine Learning Through Dynamic Python Bytecode Transformation and Graph Compilation (ASPLOS 2024)
• He, Making Deep Learning Go Brrrr From First Principles (2022)
• PyTorch, torch.profiler (2026); NVIDIA, Nsight Systems and Nsight Compute (2026)

Last reviewed: 2026-10