# CS 211 - Lecture 22

## Computer Architecture

### Parellel and GPU Architecture

Bernhard Firner

2025-11-18

---

## Review: Cache

* Terminology
  * Block, Line, Set
* Looking up where a block goes with its memory address
  * Set and tag
  * Bits for block memory aren't used
  * 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
* If the block is 32 bytes, then reading the 0,0 location reads in 32 bytes (4 doubles)
  * The first value is used
* Then we go to row 2, getting another cache miss and skipping the other 9 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 the other 24 bytes

---

## 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)

---

## Originally 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

<p style='font-size:13pt'>$^*$ Source: [Inside NVIDIA Blackwell Ultra](https://developer.nvidia.com/blog/inside-nvidia-blackwell-ultra-the-chip-powering-the-ai-factory-era/)</p>

---

## 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

---

## 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
* Parellelism 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
* 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's 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
-->