# CS 211 - Lecture 22

## Computer Architecture

### Parallel and GPU Architecture

Bernhard Firner

2026-04-15

---

## Review: Cache

* Terminology
  * Block, Line, Set
* Looking up where a block goes with its memory address
  * Set and tag
  * Bits for block memory denote an offset in the block
  * Middle bits must be used for set selection so contiguous memory doesn't collide
  * Upper bits used to distinguish between blocks in a set

---

## Cache Misses in Loops

* The first read of anything causes a compulsory miss
* When indexing a matrix by column and the row doesn't change
  * Contiguous memory combined with sequential prefetching will prevent cache misses
* Indexing by a matrix row instead can cause misses

---

## Miss Example

* If a matrix is 10x10 doubles, it is 80 bytes per row
* With a block size of 64 bytes, reading the 0,0 location reads in 64 bytes (8 doubles)
  * The first value is used for the request
* Then we go to row 2, getting another cache miss and skipping the other 7 doubles in the first block
  * The rest of the block is unused
  * Unless it remains in the cache until we loop around again, we wasted that bandwidth

---

## Cache Simulator

* There is a nifty online simulator
  * [https://courses.cs.washington.edu/courses/cse351/cachesim/](https://courses.cs.washington.edu/courses/cse351/cachesim/)
* If you aren't getting some cache concepts, try playing with it

---

## GPUs

* Graphics Processing Units
* Two major makers today
  * NVIDIA
    * [https://www.nvidia.com/en-us/technologies/](https://www.nvidia.com/en-us/technologies/)
  * AMD
    * [https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html](https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html)

---

## Original Uses

* GPUs were originally exclusively for graphics
  * Not for training AI or mining bitcoin
* Began with simple screen drawing of 2D graphics
* Eventually moved on to rendering of 3D graphics

---

## Increased Capabilities

* Early forms *only* did basic math
* Complex rendering required more features
  * This eventually lead to more flexibility
  * Logic and looping, etc

---

## New Uses

* GPUs got fast at doing parallel computation
  * Generally matrix math
* Linear algebra gives us ways to use matrix math for lots of problems
  * This pushed some scientific computing onto GPUs
* Direct support for programming GPUs was introduced with CUDA and OpenCL

---

## What's in it?

* Like CPUs, GPUs also have cache
  * At varying levels, L0, L2, L3
* And the same split between instructions and memory
* Their own DRAM, independent of the system's memory

---

## What's the Difference?

* So how are GPU's different from CPUs?
  * They are optimized for parallel operations
  * In other words, you won't find the best branch prediction here
  * But you will find high bandwidth and throughput

---

## Execution Model

* Modern CPUs have multiple cores
  * Maybe 100!
* GPUs had 100s to 1000s... in 2017
* Modern GPUs architectures (Blackwell and Instinct MI300) have far more
* 100s of GB/s of bandwidth

---

## AMD Compute Units

* MI300 is the latest version
* 8 Compute Engines
  * Each with 38 compute units
    * Each CU has a 32KB L1 cache
  * Shared 4MB L2 cache
* 128 FP64 operations/clock/cu
  * 81.7 TFLOPS
* Goes up to 2614.9 TFLOPS 8-bit

---

## NVIDIA Cores

* Blackwell architecture
* Compute split across Tensor cores and CUDA cores
* 8 Graphical Processor Clusters
  * Shared L2 cache between pairs
  * 256KB L1 data cache
  * 160 streaming processors per cluster
    * 4 tensor cores per streaming processor
* 5 (10 sparse) Peta FLOPS 8 bit

---

## Bandwidth

* How do we keep all of these FLOPS fed?
* 8 TB/s per GPU on Blackwell
* Direct memory access
  * Bypass CPU and dump data directly from Ethernet onto the GPU

---

## Tensor Core Vs CPU Cores

* In 2017 NVIDIA added Tensor cores alongside existing CUDA cores
* Tensor cores operate directly on a small matrix
  * Remember how we could optimize our matrix multiply to work on blocks?
* The SIMD CPU instructions we talked about last time are also going in this direction

---

## Cost of Compute

* In current architectures, large amounts of memory are dedicated just to tensor cores
* The more "general purpose" CUDA cores get squeezed out
  * And where does rendering happen?
  * This is still used for graphics, right?

---

### Architecture

<small> Source: [Inside NVIDIA Blackwell Ultra](https://developer.nvidia.com/blog/inside-nvidia-blackwell-ultra-the-chip-powering-the-ai-factory-era/) </small>

---

## Silicon Tradeoffs

* The focus on ML workloads squeezes rendering
* When Tensor Cores were first introduced, they were a risk
  * But ML was used to compensate
* Deep Learning Super Scaling (DLSS)
  * Render the image at quarter resolution
  * Then use ML to upscale!
* Is it a worse graphics card for gaming? Maybe.

---

## Trends

* Sometimes these hardware decisions push an industry
  * Now we are seeing matrix units moving into CPUs
* This pushes us to rewrite our algorithms to run quickly on this hardware
  * Here's [quicksort on a GPU](https://github.com/NVIDIA/cuda-samples/blob/master/Samples/3_CUDA_Features/cdpSimpleQuicksort/cdpSimpleQuicksort.cu)
* Once a technology is widely adopted, it is unlikely to go away

---

## Other Side

* GPUs have been expanding their roles from graphics into scientific computing and ML
  * But one piece of hardware can't do everything
  * What about special purpose hardware just for ML?
* One example is Google's Tensor Processing Unit
  * [https://blog.google/products/google-cloud/ironwood-tpu-age-of-inference/](https://blog.google/products/google-cloud/ironwood-tpu-age-of-inference/)

---

## Google TPUs

* No graphics, just math
* 4614 TFLOPS (8-bit)/chip in the latest iteration
* Around 30 TFlops/Watt

---

## Systolic Arrays

* The TPU also wants data to be cut into blocks
* In fact, the TPU *needs* data tiled into 128x128 blocks
* The compute takes advantage of something called a *[systolic array](https://en.wikipedia.org/wiki/Systolic_array)*
  * The compute nodes perform a single operation independently
  * Each node is triggered into activity by the arrival of data
  * Thus there are no instructions required

---

## Advantages

* Google's [TPU documentation](https://docs.cloud.google.com/tpu/docs/system-architecture-tpu-vm#chips) has some nifty animations showing off
* The advantage of the systolic array architecture is that the different operations of a matrix multiply are purely in hardware
  * No sequence of steps, no storing and reading of operands
* The disadvantage? It can only do one type of operation
  * So data has to be sent to a different piece of hardware if it isn't a matrix multiply

---

## More Details

* Those systolic arrays can consume a ton of data, so they are kept busy with a high bandwidth memory
  * This operates [up to 2.8TB/s](https://jax-ml.github.io/scaling-book/tpus/)
* The memory is transferred into a vector memory unit
  * Similar to an L1/L2 cache, but larger and controlled by the programmer
* The HBM could be 16 GiB compared to VMEM of 128MiB

---

## Work In Parallel

* On the GPU we transfer data first before operating upon it
* On a TPU, matrices are chunked and transferred
* Work usually begins after the first chunks arrive rather than waiting for the transfer to complete
* TPU v6e has reported FLOPs/second/chip of 1.84e15 for int8 values, 9.20e14 for 16-bit float

---

## But there's more!

* Notice that everything is being measured on 8-bit values
  * That isn't even a type in x86_64
  * 16-bit floats aren't either
* And now NVIDIA is pushing for [FP4](https://developer.nvidia.com/blog/nvfp4-trains-with-precision-of-16-bit-and-speed-and-efficiency-of-4-bit/)!
* We'll get to that

---

## FP TFLOPS

* From AMD MI300x datasheet
* It's obvious *why* there's a push to smaller data types

---

## Execution Model

* Let's discuss about this parallel execution model
* We had a brief experience with threads during our fun with matrix multiply
  * Threads didn't develop with CPUs; they were introduced later
  * SIMD instructions are too new to be represented in language syntax
* Parallelism has been with GPUs for a long time now
  * So the execution model has threads at its core

---

## Parallel Execution

* Two GPU languages are OpenCL and CUDA 
  * I only know CUDA, so that's what we'll use
* Execution is done by *threads* that are inside of *blocks*

---

## Memory Hierarchy

* <p style='font-size:13pt'> From [https://developer.nvidia.com/blog/cuda-refresher-cuda-programming-model/](https://developer.nvidia.com/blog/cuda-refresher-cuda-programming-model/) </p>

---

## Streaming Processors

* Recall that NVIDIA GPUs have clusters of streaming multiprocessors (SM)
* Each SM supports a block of up to 1024 threads
* Parallelism can be spread across blocks and threads
  * So the 256KB L1 cache is only shared per 1024 threads
  * This brings up some locality concerns when spread over multiple blocks
    * Although we've seen that brute force application of threads leads to speedups

---

## Grids

* CUDA threads run on a CUDA core, which reside in the SMs
* When a calculation is spread over multiple SMs it is called a grid
* So to specify a threaded GPU function call, you must also specify the grid size and threads per block

---

## Cuda Syntax

```c
__global__ void gpu_func(inputs...) {
   int idx = blockDim.x * blockIdx.x + threadIdx.x;
}
```

* Functions look like C++
* Functions on the GPU can access their block and thread location
  * These would be like a loop counter in an outer loop
* Some explicit commands to synchronize threads and share memory

---

## Calling CUDA

```C
dim3 threadsPerBlock(16, 16)
dim3 numBlocks(10, 10)
someFun<<<numBlocks, threadsPerBlock>>>(arg1, arg2);
```

* The above code would run 16x16 (256) threads on 100 different blocks
* Each thread is independent, and can check its index to determine its task
  * Notice from that earlier figure that there are 256KB of registers
  * In an A100, which is 5 years old

---

## Breaking Into Blocks

* 1024 threads can only handle a 32x32 matrix
  * If we are multiplying out an entire row by column per thread
* So we'll process blocks of 32x32 in the output matrix
  * Each output value comes from an input row and column

---

## Code

```c [35-49 | 51-58 | 60-67 | 69-83 | 7-13 | 15-20 | 22-28 ]
// For the CUDA runtime routines (prefixed with "cuda_")
#include <cuda_runtime.h>
#include "helper_cuda.h"

#include "hw3_matrix_lib.h"

/**
 * CUDA Kernel Device code
 *
 * Computes the vector multiply of A and B into C.
 * We must ensure that all three have the same size.
 */
__global__ void matrixMultiply(const double *left, const double *right, double *out, int size)
{
    // The function is called by an arbitrary execution thread
    // We can deduce our index by inspecting our current block and thread
    // blockDim.x is the number of threads per block
    // blockIdx.x is the block that we are in
    int column = blockDim.x*blockIdx.x + threadIdx.x;
    int row = blockDim.y*blockIdx.y + threadIdx.y;

// Compute this element if we are inside of the matrix
    // Note that we likely have more threads than required, so some won't do anything
    if (row < size) {
        for (int i = 0; i < size; ++i) {
            out[row*size + column] += left[row*size + i] * right[i*size + column];
        }
    }
}

/**
 * Declare with extern "C" for the linker.
 * nvcc builds with the assumption that it will be linked to c++
 */
extern "C" Matrix matMul(Matrix left, Matrix right, Matrix product) {
    // CUDA memory is on the GPU, so we need special functions to allocate memory over there.
    double *gpu_left = NULL;
    checkCudaErrors(cudaMalloc((void **)&gpu_left, left.width * left.height * sizeof(double)));

// checkCudaErrors is a macro in helper_cuda.h that will print the file and
    // line number of an error.

// Right matrix GPU memory
    double *gpu_right = NULL;
    checkCudaErrors(cudaMalloc((void **)&gpu_right, right.width * right.height * sizeof(double)));

// Output matrix GPU memory
    double *gpu_out = NULL;
    checkCudaErrors(cudaMalloc((void **)&gpu_out, left.width * right.height * sizeof(double)));

// Copy the left and right matrices onto the GPU
    // We call the Async version so these happen in the background as we keep going
    // The stream is how we synchronize between threads
    cudaStream_t stream;
    cudaStreamCreate(&stream);

checkCudaErrors(cudaMemcpyAsync(gpu_left, left.data, left.width * left.height * sizeof(double), cudaMemcpyHostToDevice, stream));
    checkCudaErrors(cudaMemcpyAsync(gpu_right, right.data, right.width * right.height * sizeof(double), cudaMemcpyHostToDevice, stream));

dim3 grid((left.height+31) / 32, (left.height+31) / 32, 1);
    dim3 blocks(32, 32, 1);

// Make sure the copying is complete
    checkCudaErrors(cudaStreamSynchronize(stream));

// Launch the matrix multiplication kernel
    matrixMultiply<<<grid, blocks, 0, stream>>>(gpu_left, gpu_right, gpu_out, left.width);

// Copy back to the host
    checkCudaErrors(cudaMemcpyAsync(product.data, gpu_out, left.width * left.height * sizeof(double), cudaMemcpyDeviceToHost, stream));

// Make sure things went okay
    checkCudaErrors(cudaStreamSynchronize(stream));

// Free device global memory
    checkCudaErrors(cudaFree(gpu_left));
    checkCudaErrors(cudaFree(gpu_right));
    checkCudaErrors(cudaFree(gpu_out));

// Done with the stream
    cudaStreamDestroy(stream);

return product;
}
```

---

## Optimal?

* There are better ways to do matrix multiplication
  * We didn't try to transpose
    * Will that be okay? Will locality fail?
  * Who knows how prefetching works on the GPU!
* Let's just try it out

---

## Compiling

* Compiling is weird
* Need to compile our C code with gcc
  * `gcc -c lec22_matrix.c -o lec22_matrix.o`
  * `gcc -c hw3_matrix.c -o hw3_matrix.o`
* Then use NVIDIA's compiler for the cuda code:
  * `nvcc -c gpu_matrix.cu -o gpu_matrix.o`
* Then link together:
  * `gcc -o lec22_matrix lec22_matrix.o hw3_matrix.o gpu_matrix.o -lcudart -lm`

---

## Results

---

## Discussion

* We pay a price to transfer data from host to GPU and back
  * Only see a break-even at around 1000x1000 matrix
  * Although we've haven't spent a moment optimizing the GPU code
* The GPU line is incredibly flat though
  * What if we go farther?

---

## Results 2

---

## Discussion 2

* About as fast with a 5000x5000 matrix as our optimized CPU code with 2000x2000
  * keep in mind, the larger matrix is 6.25 times larger
* Both the CPU and the GPU are old, but this trend has held up
  * GPUs make it easier to handle big data because they are optimized for it

---

## Impacts

* GPU optimizations are creeping into CPUs
  * Without the SIMD instruction, even small matrix math would be better with a GPU
* That allows CPUs to nibble on the low end of the workloads
* Meanwhile, TPUs are being optimized for machine learning (and only for machine learning)
* This is just the latest in a long history of competition between special purpose and general purpose hardware

---

## Thoughts

* Does it make sense for the ML hardware to be the same as the graphics hardware?
  * That's not clear
* The graphics are arguably worse because of the increase in ML-specific hardware
  * But on-chip ML has enabled things like low-res ray tracing followed by AI upscaling
  * So maybe it is both good and bad

<!--
TODO: Put in the a graph of speedups given data sizes
Would like to intro GPUs and changes over time
Also talk about TPUs, mention competition in the space.

Would like to show nvcc making a parallel matrix multiply
  I guess with transposed matrix and with block multiply
-->

<!--
https://docs.cloud.google.com/tpu/docs/intro-to-tpu#TPU
The XLA compiler performs code transformations, including tiling a matrix multiply into smaller blocks, to efficiently execute computations on the matrix unit (MXU). The structure of the MXU hardware, a 128x128 systolic array, and the design of TPUs memory subsystem, which prefers dimensions that are multiples of 8, are used by the XLA compiler for tiling efficiency.

Consequently, certain layouts are more conducive to tiling, while others require reshapes to be performed before they can be tiled. Reshape operations are often memory bound on the Cloud TPU.

https://docs.cloud.google.com/tpu/docs/system-architecture-tpu-vm
192 GB of RAM

Also, TPUs use bfloat16

https://docs.cloud.google.com/tpu/docs/ironwood-performance
>  A common and effective approach is to use a mix of FP8 formats. For example, use E4M3 for weights and activations in the forward pass to take advantage of the higher precision of E4M3, and use E5M2 for gradients in the backward pass to accommodate for the wider dynamic range of the gradients.

And for number formats:
https://research.google/blog/turboquant-redefining-ai-efficiency-with-extreme-compression/
The paper: PolarQuant: Leveraging Polar Transformation of Efficient Key Cache Quantization and Decoding Acceleration
https://arxiv.org/abs/2502.00527
snazzy website: https://turbo-quant.com

-->