GPU Architecture and CUDA: A Low-Level Enthusiast's Overview

Introduction

Modern Graphics Processing Units (GPUs) have evolved from simple graphics accelerators into powerful parallel processors driving high-performance computing and machine learning. As a low-level computing enthusiast, I find GPU architecture fascinating because it challenges traditional CPU design assumptions and offers massive data-parallel throughput. In this article, we will explore the fundamentals of GPU architecture and the CUDA programming model, highlighting how GPUs achieve high performance through parallelism and low-level optimizations. The discussion is technical yet it provides an intuitive insight into how GPUs work under the hood.

CPU vs GPU: Latency vs Throughput

Central Processing Units (CPUs) and GPUs are designed with different goals in mind. A typical CPU is optimized to minimize the latency of single-thread execution – it's built to complete a task as quickly as possible and to rapidly switch between tasks. This makes CPUs excellent for sequential, varied workloads and operating system tasks. They feature a handful of very powerful cores with large caches and sophisticated control logic (branch predictors, out-of-order execution) to maximize single-thread performance. In contrast, a GPU prioritizes throughput: it processes as many tasks or data elements in parallel as possible. To achieve this, GPUs consist of thousands of smaller cores that are individually simpler than CPU cores, but together can handle a vast number of simultaneous operations. In essence, a CPU aims for low latency on a few threads, whereas a GPU aims for high throughput across many threads.

The fundamentally different philosophies can be summarized as follows: a CPU behaves like a "sprinter," running a few tasks with great speed and agility (handling complex code and branching), whereas a GPU is more of a "workhorse," moving a huge volume of data through the same operation pipeline simultaneously. This makes GPUs ideal for data-parallel problems – scenarios like graphics rendering, deep learning, or scientific simulations where the same operation (e.g., matrix multiplication) must be applied to many data elements at once. By matching hardware design to workload characteristics, GPUs deliver phenomenal performance on these parallel tasks that would bottleneck a CPU.

The CUDA Programming Model: Threads, Blocks, and Warps

To harness GPU parallelism, NVIDIA's CUDA model introduces a clear hierarchy of threads. When a GPU kernel (a parallel function) is launched, it spawns a grid of thread blocks. Each thread block contains a number of threads (up to thousands) that execute the same kernel code concurrently. All threads have their own registers and local memory, but threads in the same block can also cooperate by synchronizing and sharing data in a fast on-chip memory. Crucially, a block is the unit of work scheduled on one GPU hardware unit called an SM (Streaming Multiprocessor), and a block never splits across multiple SMs. This means all threads in a block execute on the same multiprocessor and can efficiently share data through the on-chip shared memory. Meanwhile, the grid (collection of blocks) can span the entire GPU, utilizing all SMs in parallel.

To clarify the terminology of CUDA's hierarchy and its correspondence to hardware, here are the key concepts:

• Thread: the fundamental execution unit in CUDA, analogous to a CPU thread. Each thread executes the kernel code on a different data element and has its own index (e.g., thread ID). Threads are organized into blocks and are the finest granularity of parallelism.

• Thread Block: a group of threads (e.g., 256 or 512 threads, up to a hardware-defined max) that can cooperate. All threads in a block run on the same SM and can synchronize with each other and share data via shared memory. A block is a scheduling unit for the GPU (it will not be split across SMs).

• Warp: a subgroup of threads that represents the basic unit of execution on the hardware. On NVIDIA GPUs, a warp consists of 32 threads that execute instructions in lockstep (SIMT – Single Instruction, Multiple Threads). The GPU's scheduler dispatches work to warps, and all threads in a warp perform the same operation simultaneously on different data. (If threads within a warp diverge due to a branch, the warp will serialize the different paths, as discussed later.)

• Grid: the collection of all thread blocks launched for a kernel. The grid may consist of tens or hundreds of blocks, which are distributed across the GPU's many SMs. In effect, a grid represents the entire problem being solved in parallel.

When a CUDA kernel runs, each thread can determine its indices (block ID and thread ID within the block) to figure out what portion of data to process. For example, a 1D vector addition kernel might use an index formula like i = blockIdx.x * blockDim.x + threadIdx.x to assign each thread to a unique element of an array. This simple calculation gives each thread a distinct i for data access, demonstrating how the grid/block/thread indexing maps the problem domain onto many parallel threads.

Streaming Multiprocessors: The GPU's Core Execution Units

The Streaming Multiprocessor (SM) is the fundamental hardware unit in an NVIDIA GPU that executes the threads. You can think of an SM as a many-core processor within the GPU: it contains numerous arithmetic logic units (ALUs, often called CUDA cores), special function units, load/store units, and other hardware responsible for executing instructions in parallel. Each SM typically manages and executes threads from several warps at once, interleaving their execution to keep the hardware busy. Unlike a CPU core, an SM is designed for throughput and zero-overhead thread scheduling – meaning it can rapidly switch between warps each clock cycle to hide latency (for instance, if one warp is waiting for a memory access, another warp can execute in the next cycle).

Architecturally, a single SM includes: multiple execution cores (for integer and floating-point math), a large register file, one or more warp schedulers, and several types of on-chip memory caches. Notably, each SM has a programmable shared memory (a.k.a. L1/shared memory) that acts as a user-managed cache for threads in the same block, plus caches for constants and textures. The register file and shared memory are critical resources that allow an SM to keep thousands of threads in flight. For example, because an SM might run up to (say) 1024 or more threads concurrently, it needs a very large number of registers to hold the context (variables) of all those threads. Indeed, GPUs dedicate significant silicon area to registers and simple cores rather than huge caches or complex control – this is the "many-lightweight-cores" approach in action.

It's important to note that GPU cores are simpler and slower clocked compared to high-end CPU cores. A GPU SM typically forgoes heavy speculation or deep out-of-order execution. For instance, GPU hardware does not perform advanced branch prediction like CPUs do. If there is a branch (conditional code) within a GPU kernel, and some threads in a warp take one path while others take another, the warp will execute both paths sequentially (masking out threads that aren't on that path) to ensure correctness. This phenomenon is known as warp divergence, and it can reduce performance because not all threads are active for the duration of the divergent code. In essence, divergence breaks the lockstep execution model of a warp, forcing serialization of the different branch paths and thus under-utilizing the GPU's ALUs. GPU programmers strive to minimize warp divergence (for example, by structuring algorithms so that threads in a warp follow the same control flow) to keep the hardware fully utilized.

Instead of sophisticated per-thread speculation, GPUs rely on massive parallelism to tolerate latency. An SM employs fine-grained multithreading: it can rapidly context-switch between warps every cycle, with zero scheduling overhead in hardware. If one warp stalls (e.g., waiting for data from memory), the SM has many other warps ready to execute, thereby keeping the pipelines busy. This strategy is often described by Little's Law from queueing theory – if the latency of an operation is fixed, the throughput can be increased by having more operations (threads) in flight concurrently. GPUs exemplify this by running tens of thousands of threads simultaneously; the large number of concurrent threads ensures there is always useful work to do, effectively hiding memory and execution latencies and maximizing throughput.

However, there are practical limits to how many threads a single SM can run at once, determined by hardware resources. Each SM has a finite pool of registers and a fixed size of shared memory. If each thread uses a lot of registers or shared memory, the occupancy (number of active threads per SM) will be lower because the resources are exhausted with fewer threads. Conversely, if threads are lightweight in resource usage, more of them can fit on the SM concurrently. For example, a kernel that uses a large shared memory tile might only allow a few thread blocks resident on each SM at a time, whereas a lightweight kernel could support a larger number of blocks/threads simultaneously. Optimizing GPU performance often involves balancing this trade-off – using enough resources to do useful work per thread, but not so much that it drastically limits parallelism. In summary, a GPU's performance is maximized when each SM has a high occupancy (many warps active) and those warps are efficiently executing without stalling.

Memory Hierarchy and Optimization Techniques

Feeding thousands of parallel execution units with data is a non-trivial challenge. GPUs address this with a memory hierarchy designed for bandwidth. The off-chip global memory (typically GDDR6/GDDR6X or HBM on modern GPUs) has very high bandwidth, but also relatively high latency (hundreds of clock cycles to access). To mitigate this, GPUs rely on parallelism and locality: many threads can perform memory accesses concurrently, and if those accesses are to contiguous addresses, the hardware will combine them into larger, efficient transfers. This is known as memory coalescing. In simple terms, if all 32 threads in a warp access addresses that lie in the same 128-byte aligned segment of memory, the GPU can service that with a single memory transaction, rather than 32 separate ones. Coalesced memory accesses thus maximize the utilization of the memory bus, achieving much higher effective bandwidth. On the other hand, if threads in a warp access scattered addresses, the accesses may be split into multiple transactions (or serialized), hurting throughput. An important optimization for CUDA programmers is to arrange data structures and thread access patterns such that threads in the same warp access adjacent memory locations whenever possible.

In addition to coalescing, GPUs have small caches and shared memory to exploit data reuse. Shared memory is a manually controlled scratchpad located on the SM, which acts like a very fast user-managed cache (on the order of tens of TB/s bandwidth). Developers can tile their algorithms to use shared memory: for example, in matrix multiplication, a block of threads can load a tile of the matrix into shared memory and each thread will reuse those values multiple times, rather than each thread fetching every element from slow global memory. This technique can dramatically improve performance. In fact, using shared memory and coalesced accesses, one can achieve orders-of-magnitude speedups on certain workloads. For instance, one report showed that an optimized matrix multiplication using shared memory increased performance from about 234 GFLOPS to 7490 GFLOPS on the same GPU – a testament to how effective on-chip memory and access pattern optimizations can be. The GPU's L1 cache and texture caches also play roles in optimizing memory access, but they are typically smaller and more specialized than a CPU's caches, since the primary strategy on GPUs is to tolerate cache misses by having many threads rather than to cache huge working sets.

Another important concept is arithmetic intensity (or operational intensity), which measures the ratio of computation to memory access (e.g., floating-point operations per byte of memory transferred). GPU kernels that have a high arithmetic intensity can more easily keep the many ALUs busy without being bottlenecked by memory throughput. In contrast, if a kernel performs only a few computations per data element (low FLOPs per byte) and has to stream a lot of data from memory, it may become memory-bound – the GPU's execution units end up waiting on data, and the achieved performance will be limited by memory bandwidth rather than compute power. Tools like the roofline model describe this by plotting performance versus operational intensity. A well-optimized GPU workload will try to increase operational intensity (for example, by reusing data from shared memory, using fused operations, etc.) so that the program is compute-bound and approaches the theoretical peak FLOPS of the GPU. For many numeric applications (dense linear algebra, deep learning matrix ops), GPUs can indeed operate near their peak throughput when optimally fed with data.

To summarize the optimization mindset: GPU programmers aim to keep the GPU busy and avoid stalls. That means launching enough threads (and using enough blocks to occupy all SMs), writing code that minimizes warp divergence, using fast memory (shared memory) to cache data and optimize memory access patterns (coalescing), and ensuring the kernel does a substantial amount of work per memory byte transferred. When these conditions are met, the GPU can reach its full potential, completing in seconds tasks that would take a CPU many minutes or hours by leveraging its massive parallel hardware.

Conclusion

GPU architecture is a marvel of parallel engineering – it abandons the single-thread speed of CPUs in favor of a design that shines for the right workloads. By understanding the low-level details of GPUs, from warps and SMs to memory coalescing and occupancy, we gain insight into how to write high-performance code that fully utilizes modern hardware. As an enthusiast without formal specialization in hardware architecture, I've found that studying these fundamentals greatly improves my ability to reason about performance and to optimize algorithms for parallel execution. This knowledge not only helps in GPU programming with CUDA, but also gives a deeper intuition about computing in general (for example, thinking about data locality, throughput vs. latency trade-offs, and parallel algorithm design).

In a professional context, this low-level awareness translates to writing more efficient and scalable software. Whether it's accelerating a machine learning inference pipeline or optimizing a game engine's rendering loop, the principles of GPU architecture guide us to make better use of the hardware. Enthusiasm for low-level details – like how registers are allocated or why a certain memory access pattern is slow – shows a mindset of continuous learning and optimization. I believe this passion for understanding "how things work under the hood" ultimately leads to robust engineering skills. As GPU technology continues to evolve (with newer cores, tensor accelerators, and advanced memory systems), I remain excited to keep learning and demystifying the architecture, ensuring that I can leverage these advancements fully in any project I undertake. The journey into GPU internals is challenging but rewarding – much like the GPU itself, it's all about embracing parallel efforts to achieve a greater outcome.