GPU Computing and CUDA: How GPUs Run Code

GPU Computing and CUDA

A CPU is built to finish one thread's work as quickly as possible, with deep pipelines, branch predictors, and large caches. A GPU is built for throughput: it runs tens of thousands of lightweight threads, each relatively slow, and keeps its arithmetic units busy by switching between them whenever one waits for memory. When a problem applies the same operation to millions of independent elements — pixels, matrix entries, tokens — that design wins by an order of magnitude or more.

GPUs started as graphics hardware and became general-purpose processors through programming models such as NVIDIA's CUDA. Today they power deep-learning training and inference, scientific simulation, video processing, and data analytics. Most developers use them through libraries such as PyTorch, but knowing how the hardware executes code explains why some workloads fly, why others crawl, and why GPU memory is usually the binding constraint.

TL;DR

GPUs execute threads in groups of 32 (warps) that run the same instruction in lockstep; branches that split a warp run both paths.
A CUDA program launches kernels over a grid of thread blocks; each thread computes its own index and handles a slice of the data.
Memory bandwidth, not raw arithmetic, limits most GPU code: adjacent threads should read adjacent addresses (coalescing).
Moving data between CPU and GPU over PCIe is slow relative to on-device memory, so keep data resident and batch work.
CUDA is NVIDIA-only; ROCm/HIP, SYCL, Metal, and WebGPU target other hardware, and frameworks such as PyTorch hide most of the difference.
For ML, GPU memory capacity decides what fits: weights, activations, optimizer state, and the KV cache all compete for it.

Quick Example

A complete CUDA program: add two vectors of 16 million floats, one GPU thread per element.

Each of the 16 million threads runs the same tiny function on a different element. The launch syntax <<<blocks, threads>>> is how CUDA expresses that parallelism.

Core Concepts

CPU Versus GPU

The CUDA Execution Model

A kernel is a function marked __global__ that runs on the GPU.
A launch creates a grid of thread blocks; each block holds up to 1,024 threads.
Blocks are scheduled onto streaming multiprocessors (SMs). Threads in a block can share fast on-chip shared memory and synchronize with __syncthreads(); threads in different blocks cannot synchronize within a kernel.
Each block is split into warps of 32 threads that execute instructions together (SIMT — single instruction, multiple threads).

The GPU Memory Hierarchy

Coalescing and Divergence

When the 32 threads of a warp read 32 consecutive 4-byte values, the hardware combines them into a few wide memory transactions — a coalesced access. If they read scattered addresses, the same warp needs many transactions and effective bandwidth collapses. Similarly, if threads in a warp take different branches (divergence), the warp executes each path in turn with some threads idle.

Occupancy and Latency Hiding

GPUs hide memory latency by keeping many warps resident per SM and switching to a ready warp when another stalls. Occupancy — active warps relative to the maximum — depends on registers and shared memory per thread. More is not always better, but very low occupancy often means the hardware sits idle waiting for memory.

Arithmetic Intensity

Arithmetic intensity is the number of operations per byte moved from memory. Element-wise operations such as vector addition do one operation per 12 bytes and are bandwidth-bound; large matrix multiplications reuse each loaded value many times and can saturate the compute units. The roofline model plots this trade-off and is the standard way to tell whether a kernel is limited by memory or compute.

Programming Options

Most teams should start at the highest level that works and write custom kernels only for profiled hotspots.

How Machine Learning Uses GPUs

Training

Training is dominated by large matrix multiplications, which GPUs accelerate with tensor cores — units that multiply small matrices in reduced precision (FP16, BF16, TF32, FP8) and accumulate in higher precision. Mixed-precision training cuts memory and roughly doubles throughput with little accuracy loss. Larger models spread across GPUs with data, tensor, and pipeline parallelism, exchanging gradients over fast interconnects such as NVLink — see Distributed Training.

Inference

LLM inference has two phases: prefill processes the prompt in parallel and is compute-bound, while decode generates one token at a time and is memory-bandwidth-bound because every step reads all the weights. Batching many requests, keeping the KV cache on the GPU, and quantizing weights (8-bit or 4-bit) are the main levers — see LLM Inference.

Estimating Memory

A rough lower bound for inference is parameters × bytes per parameter: a 7-billion-parameter model needs about 14 GB in FP16 or roughly 4 GB at 4-bit precision, plus the KV cache, which grows with context length and batch size. Training needs several times more for gradients, optimizer state, and activations.

Best Practices

Keep Data on the Device

Transfer inputs once, run as many operations as possible on the GPU, and copy back only results. Pinned (page-locked) host memory and asynchronous copies on CUDA streams let transfers overlap with computation.

Design for Coalesced Access

Lay out data so consecutive threads touch consecutive addresses — structure-of-arrays rather than array-of-structures — and use shared memory to stage tiles that are reused.

Batch Small Work

Every kernel launch costs microseconds. Fuse small operations (compilers such as torch.compile and Triton do this), batch requests, and avoid per-element launches.

Check Errors and Synchronize Deliberately

Kernel launches are asynchronous, so errors surface later. Check cudaGetLastError() after launches in debug builds, and synchronize only where results are needed — each synchronization stalls the pipeline.

Profile Before Tuning

Use Nsight Systems for timeline-level problems (idle gaps, transfers, synchronization) and Nsight Compute for kernel-level ones (occupancy, memory throughput). In PyTorch, torch.profiler shows the same information at the operator level.

Common Mistakes

Copying Data Every Iteration

Strided, Uncoalesced Memory Access

Using a GPU for Small or Branchy Work

FAQ

Do I need to learn CUDA to use GPUs for machine learning?

No. Frameworks such as PyTorch and JAX, and libraries such as cuDF, already contain optimized GPU kernels; you move tensors to the device and call the API. CUDA knowledge helps when you profile, write custom fused kernels, or need to understand why a model runs out of memory or underuses the GPU.

Why is my GPU utilization low?

Common causes are a slow input pipeline (the GPU waits for data), small batches, frequent host–device copies or synchronizations such as .item() in a loop, and many tiny kernel launches. A profiler timeline shows the idle gaps and what the GPU was waiting for.

How do CUDA, ROCm, SYCL, and WebGPU differ?

CUDA targets NVIDIA GPUs and has the largest ecosystem. ROCm and its HIP API target AMD GPUs with a CUDA-like model. SYCL is a vendor-neutral C++ standard, used heavily on Intel hardware. WebGPU and Vulkan compute are portable APIs, with WebGPU also available in browsers. Frameworks such as PyTorch support several of these backends.

How much GPU memory do I need to run an LLM?

Start with parameters × bytes per parameter for the weights — about 2 bytes in FP16, 1 byte at 8-bit, half a byte at 4-bit — then add the KV cache, which grows with context length, batch size, and the number of layers. A 7B model fits comfortably on a 16–24 GB GPU when quantized; serving long contexts to many users needs considerably more.

When is a GPU slower than a CPU?

When the problem is small, sequential, or full of data-dependent branches, or when moving data to the GPU costs more than the computation saves. Latency-critical single requests can also be faster on a CPU because they avoid transfer and launch overhead.