2.1 GPU vs. CPU Architecture & Parallel Execution

GPU vs. CPU Architecture & Parallel Execution

Modern artificial intelligence and high-performance computing (HPC) workloads rely fundamentally on hardware accelerated computing. Understanding the foundational design philosophies that separate Central Processing Units (CPUs) from Graphics Processing Units (GPUs) is essential for deploying, managing, and optimizing enterprise AI infrastructure.

---

1. Architectural Philosophies: Latency Optimization vs. Throughput Optimization

At the core of processor design lies a fundamental trade-off between minimizing latency for individual tasks and maximizing throughput across massive volumes of parallel data.

CPU Architecture: Latency-Oriented Design

CPUs are engineered to execute complex, sequential streams of instructions with the absolute lowest possible latency per instruction. A modern enterprise CPU (such as an Intel Xeon or AMD EPYC) features between 16 and 128 sophisticated physical cores running at high clock frequencies (3.0 GHz to 4.5+ GHz).

To ensure single-threaded sequential performance remains high, CPUs dedicate approximately 70% to 80% of their silicon die area to control logic and cache hierarchies rather than raw arithmetic compute units:

• Large Multi-Level Caches (L1, L2, L3): Megabytes of ultra-low-latency on-chip Static Random-Access Memory (SRAM) keep instructions and data physically close to execution units, avoiding slow trips to system DRAM.
• Out-of-Order (OoO) Execution Engines: Dedicated hardware logic analyzes instruction dependency graphs at runtime, reordering and executing instructions speculatively to keep internal pipelines saturated.
• Sophisticated Branch Predictors: Advanced branch prediction tables and algorithms (including neural branch predictors) anticipate the direction of conditional if-else branches, speculatively executing subsequent code to prevent pipeline bubbles.
• Deep Execution Pipelines: Pipelining allows high clock rates by breaking instruction execution into 14 to 20+ fine stages.

GPU Architecture: Throughput-Oriented Design

GPUs were originally developed for graphics rendering—a domain characterized by millions of independent pixels requiring identical geometric and color transformations. Modern GPUs (such as the NVIDIA Hopper H100 or Blackwell B200) apply this paradigm to deep learning, matrix algebra, and tensor operations.

Rather than accelerating a single thread, a GPU is designed to maximize aggregate computational throughput across tens of thousands of concurrent threads. GPUs dedicate 80% to 90% of their silicon die area to arithmetic logic units (ALUs), including standard 32-bit floating-point (FP32) cores, integer (INT32) cores, 64-bit floating-point (FP64) cores, and specialized matrix calculation units known as Tensor Cores.

To accommodate thousands of compute cores on a single die, GPUs intentionally omit complex control hardware:

• No Out-of-Order Execution: Instructions are issued in-order per thread.
• Minimal Branch Prediction: GPUs do not maintain speculative execution pipelines for divergent branches.
• Modest Cache Capacity per Core: While modern GPUs include substantial shared L2 caches (e.g., 50 MB on H100), the cache capacity per active thread is tiny compared to a CPU.

Instead of preventing memory stalls through huge caches, GPUs use Thread-Level Parallelism (TLP): when one group of threads stalls waiting for data from high-bandwidth memory (HBM), the hardware instantly switches execution to another ready group of threads, keeping ALUs continuously utilized.

---

Comparison: CPU vs. GPU Architectural Profiles

2. SIMD vs. SIMT Execution Paradigms

Parallel hardware architectures utilize distinct programming and hardware execution abstractions to process multiple data elements concurrently.

Single Instruction, Multiple Data (SIMD)

SIMD is a data-parallel hardware model traditionally implemented in CPU vector extensions (such as Intel AVX-512, AMD AVX2, or ARM NEON).

In SIMD:

• A single CPU thread controls a vector register with a fixed width (e.g., 512 bits, capable of holding sixteen 32-bit floating-point numbers).
• A single instruction (e.g., _mm512_add_ps) executes across all elements of the vector register simultaneously in hardware.
• Programmer Burden: The programmer or compiler must explicitly pack data into contiguous vector registers (vectorization). Conditional logic (if/else) requires manual bit-masking and blend operations across vector lanes.

Single Instruction, Multiple Threads (SIMT)

SIMT is the architectural paradigm developed by NVIDIA for CUDA computing. SIMT bridges the gap between scalar programming simplicity and vector execution efficiency.

In SIMT:

• The programmer writes code for a single scalar thread, expressing operations from the perspective of an individual data element (threadIdx.x).
• The GPU hardware automatically groups 32 scalar threads into a single hardware execution unit called a Warp.
• The hardware warp scheduler issues a single instruction to all 32 threads in the warp simultaneously. Each thread executes the instruction on its own private registers and independent data.
• Key Distinction: Unlike SIMD, each thread in SIMT maintains its own Program Counter (PC), register state, and call stack abstraction. This allows threads within a warp to follow independent execution paths when necessary (handled via hardware execution masking).

+-------------------------------------------------------------------------+
|                               SIMD vs SIMT                              |
+------------------------------------+------------------------------------+
| SIMD (Vector Model)                | SIMT (NVIDIA CUDA Model)           |
+------------------------------------+------------------------------------+
| Single thread executes vector ops  | 32 independent threads form a warp |
| Vector registers (e.g., 512-bit)   | Scalar registers per thread        |
| Compiler/programmer packs vectors  | Hardware groups scalar threads     |
| Explicit mask registers for branch | Automatic hardware warp masking    |
| Rigid lockstep on single core      | Flexible thread-level parallelism  |
+------------------------------------+------------------------------------+

---

3. Warp Architecture, Scheduling & Latency Hiding

The 32-Thread Warp

The fundamental atomic unit of execution on an NVIDIA Streaming Multiprocessor (SM) is the Warp. A warp consists of exactly 32 parallel threads.

• When a CUDA kernel launches a thread block containing 256 threads, the hardware SM immediately decomposes that block into $256 / 32 = 8$ warps (Warps 0 through 7).
• Threads in a warp always have consecutive thread IDs (e.g., Lane IDs 0 through 31).
• All active threads in a warp execute the same instruction simultaneously on different data elements.

Zero-Overhead Warp Scheduling and Latency Hiding

Memory access to off-chip Global Memory (DRAM/HBM) takes between 400 and 800 clock cycles. If a traditional processor stalls for 400 cycles waiting for memory, it attempts to mitigate the stall through deep out-of-order execution buffers or falls idle.

GPUs solve memory and arithmetic latency through zero-overhead hardware warp scheduling:

1. Hardware Context Storage: An SM contains a massive physical register file (e.g., 64K 32-bit registers per SM on Ampere and Hopper). The register state of every active warp currently allocated to the SM is held permanently in on-chip SRAM registers.
2. Zero-Cycle Context Switching: Because the architectural state of all allocated warps already resides in hardware registers, switching execution between warps requires zero clock cycles. No state is pushed or popped from a stack, and no operating system interrupt or kernel context switch is triggered.
3. Latency Hiding in Action:
   • At Cycle $t$, Warp 0 issues a global memory read (LDG) and stalls, waiting for data from HBM3.
   • At Cycle $t+1$, the SM's warp scheduler checks its pool of active warps, finds that Warp 1 is ready with arithmetic operands in registers, and dispatches Warp 1's instruction to the FP32 ALUs.
   • At Cycle $t+2$, the scheduler dispatches Warp 2 to the Tensor Cores.
   • By the time all other ready warps have executed their instructions, Warp 0's memory load from HBM3 has completed, and Warp 0 returns to the ready pool without the compute pipelines ever sitting idle.

This principle is governed by Little's Law applied to parallel computing: $\text{Concurrency Required} = \text{Throughput} \times \text{Latency}$

To fully saturate the arithmetic pipelines of an SM, a kernel must maintain sufficient active warps (high SM Occupancy) to hide both memory fetch latencies and multi-cycle arithmetic instruction latencies.

4. Warp Divergence (Branch Divergence) & Mitigation

While SIMT provides the abstraction of independent scalar threads, the underlying hardware executes instructions across the 32 threads of a warp in lockstep.

How Warp Divergence Occurs

When code contains a conditional branch (such as an if-else statement or a data-dependent while loop) where different threads within the same warp evaluate different branch conditions, warp divergence occurs.

Consider the following CUDA code:

__global__ void processArray(float *data) {
    int tid = threadIdx.x;
    if (tid % 2 == 0) {
        // Path A: Even threads perform multiplication
        data[tid] = data[tid] * 2.0f;
    } else {
        // Path B: Odd threads perform square root
        data[tid] = sqrtf(data[tid]);
    }
}

In this kernel, threads with even IDs (0, 2, 4, ... 30) evaluate the condition to true, while threads with odd IDs (1, 3, 5, ... 31) evaluate it to false.

The Hardware Execution Penalty

Because the SM's warp dispatch unit can only broadcast one instruction type to the warp's execution lanes at any single clock cycle, the hardware resolves divergence through serialization and execution masking:

1. Pass 1 (Path A): The SM activates an internal 32-bit active mask, enabling lanes 0, 2, 4, ..., 30 and disabling (masking off) lanes 1, 3, 5, ..., 31. The warp executes the multiplication instruction. The disabled odd threads sit idle, consuming power and execution time without performing useful work.
2. Pass 2 (Path B): The SM inverts the active mask, enabling odd lanes 1, 3, 5, ..., 31 and disabling even lanes. The warp executes the square root instruction while the even threads sit idle.
3. Reconvergence: Once both branch paths reach the reconvergence point, the active mask is restored to full 32-thread enablement, and the warp resumes full-throughput lockstep execution.

In this scenario, total execution time is the sum of Path A plus Path B, cutting effective computational throughput in half (50% efficiency).

Thread Lane ID:  0  1  2  3  4  5 ... 30 31
--------------------------------------------
Step 1 (Path A): [X][ ][X][ ][X][ ] ... [X][ ]  (Even active, Odd masked)
Step 2 (Path B): [ ][X][ ][X][ ][X] ... [ ][X]  (Odd active, Even masked)
Reconverged:     [X][X][X][X][X][X] ... [X][X]  (All 32 lanes active)

Volta+ Independent Thread Scheduling (ITS)

Starting with the NVIDIA Volta architecture (V100) and continuing through Ampere (A100), Hopper (H100), and Blackwell (B200), NVIDIA introduced Independent Thread Scheduling (ITS):

• Each thread retains its own dedicated Program Counter (PC) and call stack rather than sharing a single PC per warp.
• Threads within a divergent warp can interleave execution at a finer granularity and synchronize explicitly using warp-level primitives such as __syncwarp().
• However, lockstep execution remains the underlying mechanism for peak arithmetic throughput: divergent code paths still incur execution serialization penalties.

Best Practices to Avoid Divergence

1. Branch at Warp Boundaries: Ensure all 32 threads in a warp evaluate to the same branch path (e.g., branching on threadIdx.x / 32 or blockIdx.x rather than threadIdx.x % 2).
2. Use Data Sorting / Binning: Pre-sort input data so items requiring similar processing paths are grouped into contiguous memory locations processed by the same warps.
3. Utilize Warp Shuffle Primitives: Use __shfl_sync(), __ballot_sync(), and __any_sync() to exchange data and vote across warp lanes without branch divergence.