Lecture 7: GPU Architecture and CUDA Programming

What this lecture covers

This lecture answers a simple question: why does CUDA look the way it does, and what actually happens on the chip when you launch a CUDA kernel? It follows CS149's running theme of multi-core, SIMD, and multi-threading, showing that GPU programming reuses the same three ideas from earlier lectures at much larger scale rather than introducing new concepts.

The lecture sits after the class's treatment of ISPC and multi-core CPU parallelism, and builds directly on that foundation: CUDA is presented as ISPC's spmd model applied to a very differently structured processor. After watching, you should be able to explain how a CUDA kernel launch turns into thread blocks and warps on real hardware, why shared memory and __syncthreads() exist, and why CUDA's programming rules (no preemption, no ordering guarantees between blocks) follow from how the scheduler actually works.

Key ideas

• Graphics pipeline origins: early GPUs existed to compute pixel colors from triangle meshes, materials, and lighting, running a small per-pixel program over millions of pixels many times per second.
• The GPGPU hack: before general compute support, programmers tricked GPUs into doing non-graphics work by drawing two triangles covering the screen and repurposing the pixel-color program to do arbitrary math (as in the Brook stream-programming project).
• Compute mode and CUDA: in 2007 Nvidia added a direct interface where you write a kernel function and ask the GPU to run "n copies" of it, an spmd model similar to ISPC's gang-of-program-instances.
• Threads, blocks, and grids: a CUDA thread is like an ISPC program instance; threads are grouped into thread blocks, and a kernel launch creates many blocks at once, similar to ISPC tasks containing gangs.
• Separate address spaces: host (CPU) and device (GPU) memory are distinct; data must be explicitly copied with cudaMemcpy before a kernel can use it, unless the system uses newer unified memory.
• Shared memory: each thread block gets a fast, block-local memory region (backed by hardware like an L1 cache) that threads can use to cooperatively load and reuse data, coordinated with __syncthreads() as a barrier.
• Warps and implicit SIMD: the hardware groups 32 consecutive threads into a warp; when all threads in a warp share the same program counter, the GPU executes them together in SIMD, without the compiler generating explicit vector instructions.
• Resource-based scheduling: the GPU scheduler assigns thread blocks to cores (SMs) only when enough execution contexts and shared memory are free, which is why block resource requirements must be known up front.

Walkthrough

From graphics chips to general-purpose GPUs (2:09)

The lecture opens with a condensed history of computer graphics: GPUs were built to take a scene description (triangle meshes, camera position, lights) and compute an image, running one small program per pixel to determine its color from its material. Because chips needed to process millions of pixels dozens of times per second, GPU vendors kept adding cores and SIMD width, which by the early 2000s made these chips look like powerful parallel processors even though their only interface was "draw these triangles." Researchers began hijacking that interface: drawing two triangles to cover the screen and repurposing the color-computation program to run arbitrary per-pixel computations, such as physics simulation steps. A Stanford project using a language called Brook formalized this into a proper stream/data-parallel programming model, compiling data-parallel code down to the same triangle-drawing hack, which set the stage for Nvidia to build direct support for general-purpose computation.

The CUDA programming model: threads, blocks, and grids (14:16)

In 2007 Nvidia introduced compute mode: instead of drawing triangles, you write a kernel function and tell the GPU to run many copies of it, an spmd model comparable to ISPC. The lecture walks through a matrix-add example: you launch a kernel with a number of thread blocks and a block size (for example, 4x3 threads per block), and each CUDA thread uses built-in variables (blockIdx, blockDim, threadIdx) to compute its own array index and do its one piece of work. Thread IDs and block IDs can be multi-dimensional, which is convenient for image, tensor, and graphics workloads because it avoids extra divide operations when computing memory addresses.

CUDA's separate memory spaces (25:19)

CPU (host) code and GPU (device) code operate in separate address spaces. The lecture shows the standard pattern: allocate an array normally in C, allocate a matching array on the device with cudaMalloc, copy data across with cudaMemcpy, and pass only the device pointer into the kernel. Dereferencing a device pointer from host code, or vice versa, is invalid and can crash the program. cudaMemcpy is also identified as effectively a message-passing operation across the PCIe bus, which can be made asynchronous to hide latency.

Using shared memory: a 1D convolution example (33:29)

A 1D convolution (averaging neighboring input elements) is used to show why per-block shared memory matters. Neighboring threads read overlapping input elements from global memory, so a smarter version has each thread cooperatively load one element into a __shared__ array sized for the whole block, with a couple of threads loading the extra boundary elements. After a __syncthreads() barrier ensures all loads have completed, every thread computes its output using the fast shared array instead of repeatedly hitting global memory. The lecture stresses that without the barrier, some threads could read shared data before it was written, producing an incorrect result.

How the GPU executes threads: warps and SIMD (52:43)

CUDA threads each have their own scalar registers and program counter, unlike SIMD lanes on a CPU where the compiler emits explicit vector instructions. The hardware groups 32 consecutive threads into a warp, and whenever all threads in a warp share the same program counter, the GPU dynamically executes them together as one SIMD instruction; this is called implicit SIMD. On the Volta-era chip described, a warp's 32 threads actually run across only 16 physical ALUs, spread over two clock cycles, which frees up instruction fetch/decode bandwidth to issue other warp instructions in between.

Inside the SM: warps, multithreading, and scheduling (59:53)

Zooming out, a streaming multiprocessor (SM) is shown containing four such cores, together holding 64 warps worth of execution context (2,048 CUDA threads), with four separate instruction fetch/decode units that each pick a warp to run per clock. This heavy multithreading lets the SM hide memory and execution latency by switching between many resident warps. Multiplying across the roughly 80 SMs on the described chip gives the total floating-point throughput and the huge number of CUDA threads (over 160,000) that can be resident on the chip simultaneously, though only a fraction execute at once.

Rules of the programming model: barriers, atomics, and thread block independence (1:13:06)

The lecture closes with why CUDA enforces its scheduling rules. A thread block's declared thread count and shared memory must fit entirely on one SM at once, because threads in a block may synchronize with barriers; running only part of a block could deadlock the rest. Thread blocks, by contrast, can safely use atomic operations on global memory even though the runtime gives no guarantee about their relative execution order, but code that assumes one block finishes before another starts is unsafe and can deadlock on hardware with few cores.

Before you watch

• Be comfortable with ISPC's spmd model, gangs, and tasks from earlier CS149 lectures, since CUDA terminology (threads, blocks) is explained by direct analogy.
• Review the basics of the traditional graphics pipeline (triangles, per-pixel shading) if unfamiliar, since the lecture uses it to motivate why GPUs became fast parallel processors.
• Recall how multi-core CPU SIMD and multi-threading work, since the lecture repeatedly contrasts CPU SIMD (compiler-generated) with GPU implicit SIMD (hardware-detected).

Check your understanding

1. Why does CUDA require a program to declare its thread block size and shared memory usage before it runs, rather than letting the GPU schedule threads freely?
2. Explain the difference between how SIMD execution arises on a CPU versus how a warp achieves SIMD execution on a GPU.
3. Why does the 1D convolution example need a __syncthreads() call, and what could go wrong without it?
4. Why is it safe for two independent thread blocks to update the same variable with an atomic operation, but unsafe for one thread block to wait on a value written by another?
5. Why did the lecturer choose two-dimensional thread and block indices for the matrix-add example, even though it was not strictly necessary?