GPU Fundamentals: Profiling Metrics & Roofline

Advanced Optimization via Evidence (Theory)

At an advanced level, optimization becomes:

• Evidence-driven: measure first
• Metric-guided: interpret profiling signals
• Model-based reasoning: use mental models like “roofline”

Roofline mental model

A simple view:

• Operational intensity (OI) = FLOPs / Bytes moved
• If OI is low → kernel tends to be memory-bound
• If OI is high → kernel tends to be compute-bound

Even if you increase occupancy, you won’t help much if you’re fundamentally limited by bandwidth or by too many memory transactions.

Nsight Workflow (Theory)

1. Nsight Systems

   • Identify CPU↔GPU gaps
   • Confirm kernel launch frequency and sizes
   • Check for synchronization that serializes the timeline

2. Nsight Compute

   • Focus on your hot kernels
   • Inspect:
     • Achieved occupancy (and why it’s limited)
     • Global memory throughput/transactions
     • Stall reasons (e.g., memory dependency, execution dependency)
     • Cache behavior (L1/L2)
     • Divergence and instruction mix

Code Example 1 — FLOPs + Bytes Reasoning (Host-Side)

This example is not a full profiler; it shows how you can compute the “shape” of OI.

# roofline_reasoning.py
def operational_intensity(flops: float, bytes_moved: float) -> float:
    return flops / bytes_moved

# Example: elementwise add
N = 1_000_000
# flops: 1 add per element (roughly)
flops = N * 1

# bytes moved: read A and B (2 floats) + write C (1 float)
# float = 4 bytes => (2 reads + 1 write) * N * 4
bytes_moved = N * (3 * 4)

oi = operational_intensity(flops, bytes_moved)
print("Operational intensity:", oi, "FLOPs/byte")

Interpretation: elementwise ops usually have low OI → memory-bound.

Code Example 2 — Improve OI via Fusion (CUDA C++ Kernel)

Fusion often increases OI by reducing intermediate global memory traffic.

Before: two kernels

1. y = a[i] * x[i]
2. out[i] = y[i] + b[i]

After: one fused kernel

#include <cuda_runtime.h>

__global__ void fused_scale_add(const float* x, const float* a, const float* b, float* out, int N) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i >= N) return;

  // Intermediate lives in registers (no global store of y)
  float y = a[i] * x[i];
  out[i] = y + b[i];
}

Why it matters

If you removed a global read/write for y, you likely:

• reduce bytes moved
• increase operational intensity
• improve achieved bandwidth utilization

Common Gotchas

• Optimizing away symptoms: occupancy may improve but time doesn’t if you’re still memory-bound.
• Ignoring launch/timeline gaps: Systems-level issues can dominate despite good kernel metrics.
• Misinterpreting bytes moved: profiling metrics may include cache effects; confirm with “transaction” and “throughput” data.
• Overfitting to one GPU: metrics vary across architectures; validate across targets.

Quick Checklist

• Start with Nsight Systems to locate macro bottlenecks
• Use Nsight Compute to inspect your hot kernels
• Ask: memory-bound or compute-bound?
• If memory-bound: improve coalescing, reduce transactions, fuse operations
• If compute-bound: reduce instruction overhead, optimize math patterns