GPU Hardware and Software (Georgia Tech MSCS)

July 20, 2026

Lecture notes from my time in Georgia Tech’s MSCS program. Reference text: Programming Massively Parallel Processors.

Module 1: Introduction to GPU’s
Module 2: Parallel Programming
Module 3: GPU Programming Introduction
Module 4: GPU Architecture
Module 5: Advanced GPU Programming
Module 6: GPU Architecture Optimizations - I
Module 7: GPU Architecture Optimizations - II
Module 8: GPU Simulation
Module 9: Multi-GPU
Module 10: Compiler Background - I
Module 11: Compiler Background - II
Module 12: ML Accelerations on GPUs
Papers

Module 1: Introduction to GPU’s

CPU’s are optimized for single threaded latency sensitive applications and have an open Instruction Set Architecture and follow SISD/SIMD model vs GPU’s optimized for throughput, may have Closed Instruction Set Architecture and follow SPMD model
5 Stage pipeline: Fetch first pulls the instruction out from cache and figures out next instruction, Decode turns an instruction into micro operations, Schedule decides when instructions execute (in order strictly in order vs out of order instructions execute when dependencies ready), Execution does the actual computation, Write-back updates memory
Can increase parallelism in 2 ways: Superscalar Processing and Multi-threading
- Superscalar processing allows every pipeline stage to handle more than 1 instruction increasing instruction level parallelism and IPC > 1 ie within an instruction stream, find instructions that don’t depend on each other and execute them early out of order
  - Can improve performance further with:
    - Large caches reduce amount of cache misses reducing memory latency
    - Deeper pipelines spread out the work allowing for short clock cycles but you need better branch predictors for this
      - Branches (if/else) make it hard to know what instruction comes next in the instruction stream so utilize branch predictor to foresee next instruction or flush entire pipeline
      - We need better branch predictors here since we can execute more cpu cycles but can lose more cpu cycles if prediction goes wrong
GPU’s have much more cores than CPU’s but much less cache since GPU’s hide latency with parallelism and multithreading and CPU’s hide latency with cache
- You don’t want big caches for GPU’s since uses much more power as they grow, does not scale to thousands of threads, expect regular memory accesses, can run other threads if something is stalled
Amdahl’s law shows that the more serial a program gets, the infinitely less speedup happens
Assume we want to write a program to add elements of an array + find min/max value
- Task decomposition will assign 1 core for add elements and 1 core for min/max
- Data decomposition will split data by core and reduce partial results for each task
CPUs can be SISD, SIMD, and MIMD and GPU’s can be SIMT
SPMD is what GPU’s use where all cores perform same work but on different data so we use data decomposition here - the hardware’s execution model here however is SIMT
GPU cores are not like CPU cores they have multiple stream processors that each handle the warps threads
- Warps are the basic unit of execution on a GPU and refer to a group of threads (standard is 32) that execute the same instruction in SIMT style one warp at a time
- The GPU pipeline is similar to a CPU pipeline but fetch/instruction happens once per warp (why we have a warp scheduler) and execution is SIMT across multiple stream processors

SM ("core")
 ├── Warp scheduler
 ├── Instruction fetch/decode
 └── SP SP SP SP SP SP ... (ALU lanes)
        ↑  ↑  ↑  ↑
       T0 T1 T2 T3   (threads of one warp)

Module 1: Introduction to GPU’s diagram

Module 2: Parallel Programming

To parallel program, discover concurrency, structure the algorithm to harness concurrency, implement algorithm, execute and fine tune the algorithm
There are multiple parallel programming patterns
- Master/Worker Pattern has a master process that manages thread pool via a task queue where workers execute tasks concurrently by dequeuing tasks from the task queue suitable for embarrassingly parallel problems
- Loop Parallelism Pattern allows you to parallelize the loop if every loop iteration is independent
- SPMD Pattern makes all processing elements execute same program in parallel with its dataset used in GPU programming
- Fork/Join Pattern makes a thread fork a task and wait for its completion via join before continuing used in programs with a single entry point
- Pipeline Pattern makes every parallel processor handle different stages of a task used for processing data streams
In shared memory all threads can see the same memory (OpenMP) while in distributed memory all processes have their own address space must request memory from other processes explicitly if needed (MPI)

Module 3: GPU Programming Introduction

Host code on CPU, execute kernel with <<<numBlocks, threadsPerBlock>>>, kernel code on GPU
Grids contain blocks and blocks contain threads where block and thread execution have no order (SPMD model + SIMT execution)
If given 4 blocks and 4 threads per block use blockIdx.x * blockDim.x + threadIdx.x for 1d access for a specific element
Shared memory is local to blocks and faster than accessing global memory due to it being stored on chip
Can do barrier synchronization for all threads within a block with __syncthreads() with a typical pattern being load → compute → store
We can run multiple blocks per Streaming Multiprocessor based on the number of threads, register size, and shared memory bounding the blocks
- Compiler will set the number of registers per CUDA block
- We care about occupancy because means higher occupancy leads to greater hardware utilization leads to more parallelism
In order to program on the device/GPU with CUDA we must first
1. Allocate memory on the device/GPU
2. Copy memory to/from host/device

Module 4: GPU Architecture

CPU’s hide latency with out of order processing, cache, instruction level parallelism while GPU’s hide latency with multithreading via warp scheduling and execution
GPU’s are very fast in context switching since every warp has its own program counter and since multiple warps are present, a context switch only requires a pointer change vs CPU where must store registers/stack/other state in memory and bring next thread state into memory
Banking allows us to compress read/write ports leading to less wiring/hardware resources/power used and allows us to access multiple banks simultaneously (vertical bank slices)
However if we read multiple values from the same bank, then we get a bank conflict and serial accesses within a bank
- To avoid register bank conflict, optimize code layout at compile time ⇒ scoreboard values since register file accesses take more than 1 cycle and are buffered ⇒ execute warp when source operands are ready
Scoreboarding in action where we use the source operands to load in data from register files and on write back we write back to the register file
- Can store mask bits to tell us which threads are active so we know when to send warp for execution
If 1 warp generates up to 32 memory requests and there are 32 Streaming Multiprocessors then 32 * 32 = 1024 requests per cycle * 64kb on a 1Ghz CPU ~= 64 TB/s bandwidth so global memory accesses are very expensive
- Combine this with a small GPU cache and this is amplified
Memory coalescing happens when we can make all memory requests from 1 load fit into 1 request
- R1 addresses are close together so can fetch in 1 go vs R2 where strides of 128 necessitate multiple memory requests

Module 5: Advanced GPU Programming

Can profile to identify application bottlenecks
Execution time = data transfer time + compute time + memory access time where all 3 can be optimized

Module 6: GPU Architecture Optimizations - I

Within a warp threads can execute different instructions leading to divergent branches which execute groups of threads at a time with active masks
Can solve with predicated execution where we compute the results then choose the result needed per branch but wastes a lot of overhead since need to calculate everything and performance gets worse for complicated branches
Can instead use SIMT stack reconvergence for larger branches to place divergent branches in and pop from stacks once current branch reaches reconvene point and push in stack if multiple subbranches at current branch
Reconvergence point can be done with compile time analysis with control flow graph and hardware implements this with a stack
- Larger warps allow more instructions to be executed but higher chance of divergence
Within warps, can group divergent threads across them to execute for better utilization
Each SM has lots of threads which means lots of registers to track so the larger the register file, the more one must pay attention to bandwidth and latency
- Can reduce access latency or overall register file size
Note 3 observations: not all threads are active since GPU execution is async and threads can take different times to complete/execute, not all registers are live and the same register access between instructions can be short or far away, registers used immediately after access or read only
Hierarchical register files and partitioned register files reduce access latency and expand total register capacity cheaply
Can also implement virtualization on register files similar to OS virtualization on the CPU
Difference between register file and software managed cache is every instruction in register file has 2 source operands and memory has a singular address
Use unified structures ie the same hardware for shared memory and register files but you need high memory bandwidth consisting of sufficient ports and bank structures to enable flexible resource sharing

Module 7: GPU Architecture Optimizations - II

GPU handles GPU page allocations and CPU handles CPU page allocations but both collaborate with each other for allocations/transfers
GPU’s originally had 1:1 mappings but now have virtual memory
Since GPUs have multiple warps with multiple threads 1 uncoalesced memory request can generate multiple TLB misses so need a high bandwidth Page Table Walker here
- Can reduce address translation cost with segmentation, larger pages, multiple levels of TLB
Can do explicit data copying or have a Unified Virtual Address space where CPU and GPU’s share the same address space and dynamic memory transfers happen on page faults
IOMMU page transfers are expensive since includes PCI-E communication and CPU interrupt service handler so can prefetch pages to predict future memory addresses, GPU managed page mapping management, cooperative work between CPU (manages large chunks) and GPU(manages small pages) drivers
- Can prefetch addresses or pages within block or Nvidia does Tree Based Neighborhood Prefetching
CPUs choose instruction among queue of ready instructions
GPU warp scheduling implements in order scheduling within a warp
- Round robin scheduling leads to memory access stalls vs greedy then oldest which switches only when memory access from current warp occurs
- 2 level scheduler used to improve energy efficiency since choosing a warp takes significant amounts of energy
- Cache conscious scheduling limits working set of warps and only when all working set is done, add new ones + prefetching aware warp scheduler and CTA-aware scheduling is also available for use

Module 8: GPU Simulation

There are many different performance modeling techniques: cycle level simulation, event driven simulation, analytical model, sampling based techniques, etc
Cycle level simulations are common with a global clock and every cycle is modeled
- Execute driven simulation at fetch vs execute instruction executes at latter stages respectively
- Trace driven simulation collects traces and you run traces in simulation which decouples simulation and execution which makes them simpler and easier to develop
Queue based modeling moves instructions between stages of pipeline
- Number of cycles in every pipeline stage is depth of the queue and pipeline width is how many can move between queues
GPU cycle level modeling is similar to CPU but simulation modeling unit is warp and a warp instruction going through the pipeline
- In order scheduling within a warp and out of order scheduling across warps
- Simulator models divergent branches with SIMT stack with all paths, stores active warp bits, should coalesce memory requests with a cache hierarchy, sectored cache (bring part of cache block instead of entire)
Analytical models don’t need to execute the entire program and hardware constraints for GPU determined by warp width * warp depth
- Do not consider GPU as multi threading processor as excludes entire gpu architecture, branch divergence, memory divergence, etc
Roofline model measures arithmetic intensity vs performance and if performance is flatlining after a point then the kernel(s) are compute bounded
- CPI = CPI steady state + CPI event 1 + CPI event 2 + …
- Multithreading CPI = CPI single thread / Warp depth OR CPI ideal multithreading + CPI resource contention
Cycle simulation takes too long at 10 IPS, 1B instructions (1s execution) takes 28 hrs and 10 hrs of workload takes 100 years
- Can parallelize simulator where every pthread is a core but need to handle memory/network accesses/transfers + most kernels are memory bound
- Can do event driven simulation with event queue instead of by cycle
- Can simplify the model by replacing non critical components with average throughput/ latency of them
  - Instead of doing fetch/decode/execution just do IPC = issue width and model only does caching and memory accesses
  - Can simplify memory model by assuming fixed latency for memory system
- With warm up time to get cache/branch predictors ready, can randomly choose where to simulate by fast forwarding to point (execution driven) or generate traces for only that pt in time (trace driven)
- Can reduce resources simulated on GPU by reducing amount of CUDA blocks, kernels, warps simulated OR reducing iteration count OR reduce input sizes OR simulating common kernels only

Module 9: Multi-GPU

Multi GPU setups have been introduced to deal with parallel computation at scale
Memory for multiple GPU setups utilize PCI-E/NVLINK for high speed connections with shared IO with a programmer seeing multiple SM’s
- Should treat this as a NUMA setup where want to allocate and access memory close to your GPU
NVLINK preferred over PCI-E for faster bandwidth due to unified memory, direct GPU access, NVSwitch/Switch chip
RDMA uses to communicate on multi GPU boards
GPU vs GPU Stream vs MPS vs MIG architecture
MIG architecture allows multiple jobs to safely and concurrently run on 1 GPU by physically partitioning them into multiple isolated instances
MPS architecture removes the isolation and partitioning requirement which reduces overhead but a crash in 1 CUDA instance can affect others

Module 10: Compiler Background - I

GPU code compiled into multiple Intermediate Representation stages into an executable
- Clang used as the FE parser and generates the AST and C++ preprocessors substitute text before compiling (think replacing DEFINE variables)
- IR allows multiple compiler optimizations to happen
- PTX virtual ISA (architecture independent) which then gets translated to SASS machine code (architecture dependent)
PTX instruction format
IR format is standard 3 opcode with dest ← op src1 src 2
Basic block is a sequence of instructions that must be executed with only 1 entry and exit point
- Can find basic blocks by finding instructions + grouping them by leader that are either
  - Target of conditional or unconditional jump
  - Follows a conditional or unconditional jump
Global code optimization is across basic blocks while local code optimization is within a basic block
- Global code optimization relies on data flow analysis and must preserve semantics of program
- Optimization examples: redundant instructions, copy propagation, dead code elimination, cod emotion, induction variable detection, reduction strength
Transfer function notation
Definition d reaches point p if path exists from d to p
Gen and kill sets define d
Reaching definitions example

Module 11: Compiler Background - II

Live variable analysis determines which variables are live throughout the program and tracks registers in this process
Static Single Assignment optimizes the def use change and enforces that variables can only be defined once in SSA form
- If variables defined in multiple places then need to use phi function
Compiler performs multiple optimizations such as loop unrolling, function inlining, dead code elimination, constant propagation, strength reduction, calculating loop invariant dependencies
Can detect divergent branches at compile time by evaluating expression dependencies

Module 12: ML Accelerations on GPUs

GPUs good for ML because ML has lots of parallelism and a high number of floating point operations with high memory bandwidth which GPUs support with flexible data types
- Deep Neural Network operation categories are element-wise (activation), reduction (pooling), and dot product (convolution, GEMM) operations
Arithmetic intensity used to compare machine’s FLOP’s/B
Tensor cores hardware primitive suited for executing matrix multiply/accumulate operations in 1 instruction vs multiple instructions for CUDA cores which leads to massive performance boost
Can do matrix multiply with systolic arrays by streaming input and doing a prefix operation
Can use smaller floating point types to transfer/access less data which increases arithmetic intensity and quantization (maps long values to shortened fix set of values) also reduces storage size and increases arithmetic intensity and throughput
There are many factors to consider when designing ML accelerators