The Three Fundamental GPU Constraints: Compute, Memory, and Parallelism

Core Concept

Every GPU kernel is first constrained by compute throughput, data-movement bandwidth, and the amount of parallel work available to hide latency. Classify the bottleneck before optimizing.

Why It Affects Performance

Misclassification leads to wasted effort: removing arithmetic rarely helps a memory-bound kernel, higher occupancy may not help a compute-bound kernel, and a grid with too little work cannot reach either hardware ceiling.

Example or Pseudocode

__global__ void saxpy(int n, float a, const float* x, const float* y, float* out) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) out[i] = a * x[i] + y[i];
}

FP32 SAXPY performs about 2 FLOPs and transfers 12 primary bytes per element, for an arithmetic intensity of roughly 0.167 FLOP/byte.

Profiling and Performance Judgment

Measure steady-state latency with CUDA Events, compute effective bandwidth as 12N/T, and use Nsight Compute SpeedOfLight and MemoryWorkloadAnalysis to determine which roof is relevant.

Five-Minute Exercise

For N=2^24 and 0.8 ms runtime, calculate GFLOP/s and effective GB/s. What is the likely bottleneck?