Skip to content

All posts

Page 33 of 129

Choosing a Vocabulary, Emitting Tokens, and Why Speculative Decoding Is Fast | CMU 11-868 Lectures 8–9

L08 goes from BPE to VOLT, a method co-authored by the lecturer Lei Li: vocabulary size has both a cost and a value, and VOLT finds the sweet spot by asking how much normalized entropy each added token removes, then solves it as an optimal transport problem. The second half covers LLaMA 3 growing its vocabulary from 32k to 128k and the cost of byte-level BPE splitting one Chinese character into three tokens. L09 moves from greedy decoding, sampling, and beam search to speculative decoding: a small model guesses N tokens and the big model checks them in one forward pass, because checking is cheaper than generating. It ends with EAGLE, which predicts final-layer features instead of tokens.

TPU, JAX, and Pallas: One Attention Kernel on TPU, from XLA Fusions to Splash Attention | CMU 11-868 Lectures 12–13

Across two lectures and more than 200 slides, Google's Srinath Mandalapu traces one attention computation from Python down to TPU VLIW instructions. L12 covers the JAX ecosystem, the memory and compute units of TPU Ironwood, and how XLA compiles attention into three fused kernels. L13 covers what XLA cannot do: using Pallas to control movement between HBM and VMEM yourself, writing FlashAttention, then adding block sparsity to get Splash Attention. The ideas match the GPU version. The difference is that on TPU the compiler does most of the scheduling, and Pallas is how you take loops and block sizes back into your own hands.

Reading Transformers, T5, LLaMA, and GPT-3 Like a Systems Engineer | CMU 11-868 Lectures 6–7

11-868 spends only two lectures on the model itself. L06 breaks the Transformer into embeddings, multi-head attention, FFN, LayerNorm, and residuals; L07 uses T5, LLaMA, and GPT-3 to show what modern LLMs changed. For a systems engineer the point is to remember the shapes: GPT-3 175B has 96 layers, d_model 12288, a 2048-token context, and trained on 300B tokens; LLaMA 65B has 80 layers, d_model 8192, and trained on 1.4T tokens. Those numbers set the workload for every acceleration, parallelism, and serving lecture that follows.

How ZeRO Cuts Data-Parallel Memory — Optimizer State, Gradients, Then Parameters | CMU 11-868 Lecture 18

Data parallelism keeps a full copy of parameters, gradients, and optimizer state on every GPU. With Adam and mixed precision that is about 20 bytes per parameter, 16 of them optimizer-related, so LLaMA-3 8B already needs 160GB. CMU 11-868 L18 builds on the ZeRO paper and animates its three stages frame by frame: ZeRO-1 partitions optimizer state, ZeRO-2 also partitions gradients, and ZeRO-3 partitions parameters too. The slides conclude that the first two stages add no communication and save up to 8x memory; stage 3 makes per-GPU memory shrink with GPU count, at what the slides estimate as about 3x the communication.

From One Chip to a Whole Datacenter — Dataflow Hardware, Kernel Fusion, Parallelism Strategies, and the Memory Bottleneck | Stanford CS149 Lecture 12

L12 is about moving data. The first part uses the SambaNova SN40L to explain dataflow architecture and metapipelining: running Llama 3.1 8B, the slides say the RDU needs about 3 kernel calls per token versus about 800 on a GPU, because it can fuse an entire decoder into one kernel. The middle part scales up to the datacenter: which collective each of TP, PP, EP, and DP requires, and why overlapping compute with communication decides how well you scale. The last part returns to energy and DRAM: moving a byte costs far more than computing on it, and memory controllers, burst mode, and HBM all attack the same problem. There is no public video for this lecture; this post relies on the slides alone.

Cache Coherence: MSI, MESI, and False Sharing | Stanford CS149 Lecture 14

When every core has its own cache, one address can have several copies, and different cores can see different values. Locks can't fix this; the hardware created it by replicating data. CS149 L14 defines what coherent means, then takes apart the snooping MSI protocol: before writing, broadcast BusRdX so everyone else invalidates. MESI adds an E state that saves the second transaction in read-then-write, and directories replace broadcast with point-to-point messages. The practical consequence for programmers is false sharing: two threads write different variables, but because they share a cache line, the line bounces between cores. In the lecture's demo it made the program three times slower.

Parallel Computing: What the Course Covers and Two Reading Routes | Stanford CS149 Course Overview

CS149 is Stanford's parallel computing course, taught by Kayvon Fatahalian and Kunle Olukotun. It runs from multi-core CPUs and SIMD through GPUs, AI accelerators, and the datacenter, then returns to cache coherence and lock-free programming. For Fall 2025, all 18 slide decks, the starter code and READMEs for 5 programming assignments, and 4 written-assignment PDFs are public, so this series rates it A3 (self-study ready). There are four gaps: the Fall 2025 lecture videos are Canvas-only; PA1 is graded on Stanford's myth machines; PA4 needs a self-funded AWS Trainium2 instance and a private course AMI; PA5's H100 job queue and leaderboard require a SUNet ID. The public videos are from 2023, and this series treats them as a listening supplement only.

Data-Parallel Thinking, Replacing Locks with Map, Scan, and Sort | Stanford CS149 Lecture 8

Lecture 8 asks you to switch mental models: stop thinking about what each worker does and write algorithms as operations on sequences, such as map, fold, scan, segmented scan, gather/scatter, sort, and groupBy. These primitives have efficient parallel implementations, and they turn irregular parallelism into regular parallelism and fine-grained synchronization into coarse synchronization. The price is extra passes over the data, so they are bandwidth hungry.

Running DNNs Efficiently on GPUs — Conv as GEMM, Blocking, Fusion, and the Road to FlashAttention | Stanford CS149 Lecture 9

L9 opens with a claim: if you understand arithmetic intensity and the roofline, you know almost everything about software-side performance optimization for modern AI. It then shows three things. Fully connected layers, conv layers, and attention all reduce to matrix multiplication (GEMM). GEMM needs blocking so data stays in cache. Adjacent layers should be fused so intermediates never round-trip through DRAM. Softmax can be computed in chunks, which is why fused attention (the core idea behind FlashAttention) never has to store the N×N matrix.

Performance Optimization Beyond the Experts — Halide's Algorithm/Schedule Split, Autoschedulers, and LLM Agents | Stanford CS149 Lecture 13

L13 asks what to do when there are too few people who can write fast code. The slides offer three answers. First, raise the level of abstraction: Halide splits what to compute (the algorithm) from how to compute it (the schedule), so one line of schedule turns the same blur into a tiled, vectorized, multi-core version. Second, intelligent search: because the schedule space is well defined, search plus a learned cost model can generate schedules automatically. Third, the emerging option of LLM agents: have a model write CUDA, run it, read the profiler, reflect, and revise, and let it improve itself with a database of examples or prompt optimization. The last slide leaves you with a question: is the real value in DSL design or in the LLM agent?

Fine-Grained Locking and Lock-Free Programming, from Test-and-Set to the ABA Problem | Stanford CS149 Lecture 16

CS149 L16 has three parts. It first looks at lock implementations through the lens of cache coherence (test-and-set, test-and-test-and-set, ticket locks, CAS, LL/SC). It then takes a sorted linked list from one big lock to hand-over-hand fine-grained locking. Finally it introduces lock-free programming: a single-producer/single-consumer queue, a CAS-based stack, the ABA problem, and hazard pointers. The slides land on a practical conclusion: when your program has the machine to itself, well-written lock-based code is often just as fast and much simpler. Lock-free designs pay off in systems where threads can be preempted or page-fault inside a critical section.

GPU Architecture and CUDA Programming | Stanford CS149 Lecture 7

CUDA's grid, thread block, and CUDA thread are programming abstractions; the GPU implements them with SMs, warps, and a hardware block scheduler. The heart of the lecture is keeping two things apart: the system may run thread blocks in any order, but all threads in one block are guaranteed to be live at once. That is why a block can cooperate through shared memory and __syncthreads(), and why the number of blocks an SM can hold is set by registers and shared memory.

Why General-Purpose Processors Waste Energy — Hardware Specialization, Tensor Cores, TPU Systolic Arrays, and Dataflow Architectures | Stanford CS149 Lecture 10

L10 starts from one equation: when power is capped, performance can only improve by spending fewer joules per operation, and a general-purpose processor spends most of its energy fetching, decoding, and moving data rather than computing. The slides' rule of thumb is that GPUs give about 10x better perf/watt than CPUs and fixed-function ASICs can reach 100–1000x. The lecture then judges the H100's Tensor Cores and TMA, Google's TPU systolic array, and reconfigurable dataflow architectures against the same checklist: tiled tensors, asynchronous compute and memory, and compute units talking directly to each other.

Fast Processors, Slow Data — Latency vs. Bandwidth, and How ISPC Separates Abstraction from Implementation | Stanford CS149 Lecture 3

The first half of L3 uses a highway and a laundry room to pull latency and bandwidth apart, then does the math: element-wise vector multiply runs at under 1% efficiency on a V100 because memory cannot feed the ALUs fast enough. The second half is about abstraction vs. implementation. ISPC lets you think in SPMD terms (a gang of program instances, each doing its share), while the compiler implements that with SIMD instructions. Mixing up the two layers is the most common source of confusion in the course.

Locality, Communication, and Arithmetic Intensity: Why Moving Less Data Beats Adding Cores | Stanford CS149 Lecture 6

CS149 Lecture 6 asks you to read "communication" broadly: data moving between a processor and its cache, its memory, or another machine all counts. Modern parallel processors have far more compute than bandwidth, so arithmetic intensity (how much computation you do per unit of data moved) decides whether you can keep the hardware fed. The levers fall into three groups: change the assignment to cut inherent communication, use blocking and loop fusion to cut cache-induced communication, and spread out or stagger accesses to reduce contention.

Multi-Core, SIMD, and Hardware Multithreading, and the Problem Each One Solves | Stanford CS149 Lecture 2

The second lecture of CS149 Fall 2025 takes a loop that computes sin(x) and adds three ideas in turn: spend transistors on more cores (multi-core), let one instruction drive many ALUs (SIMD), and interleave several threads on one core to hide memory latency (hardware multithreading). The first two add compute; the third keeps that compute busy while waiting on memory. The conclusion is three requirements: enough parallel work, groups of work that run the same instructions, and more parallel work than ALUs so latency can be hidden.

Measuring Speedup on a Quad-Core CPU and Explaining Why It Isn't Linear | Stanford CS149 Programming Assignment 1 / Written Assignment 1

PA1 has little code and a lot of analysis. Its six programs cover work assignment across threads, SIMD masking, ISPC gangs and tasks, how input data shapes SIMD efficiency, a bandwidth-bound saxpy, and finding a K-Means hotspot with timers. Written 1 drills the same intuitions on paper: peak throughput, instruction dependencies, pipelining, latency hiding with multithreading, and SIMD divergence. Official grading uses Stanford's myth machines; you can run everything on your own hardware, but your numbers won't match the reference.

Building a Task Execution Library from Scratch — Thread Pools, Sleeping, and Task Graphs with Dependencies | Stanford CS149 Programming Assignment 2

CS149 PA2 has you write a C++ task execution library for a multi-core CPU, and write it four times: spawn threads on every run(), switch to a spinning thread pool, switch to a sleeping thread pool, and finally extend it to asynchronous task graphs with dependencies. Every step must be a fully correct system, and each is timed against the official reference implementation. Official grading runs on AWS c7g.4xlarge; you can work on your own multi-core machine outside Stanford, but your numbers won't be directly comparable to the official thresholds.

and Written 2: A CUDA Circle Renderer That Must Be Both Correctly Ordered and Fast | Stanford CS149 Programming Assignment 3 / Written Assignment 2

PA3 has three parts: port SAXPY to CUDA and time it two ways, implement find_repeats with an exclusive scan, and write a CUDA circle renderer that is both correct and fast (85 points). The hard part of the renderer is that blending semi-transparent circles doesn't commute, so every pixel must be updated in input order, and the starter code's one-thread-per-circle approach gets neither atomicity nor order right. Written 2 has five graded problems (fusion, SIMD utilization, a barrier instead of locks, data-parallel primitives on graphs, locks in a particle simulation) plus 14 practice problems. Outside Stanford you need your own NVIDIA GPU. No solutions here.

Moving Your Own Data on Trainium2 — NKI, SBUF/PSUM, and a Fused Conv+Maxpool | Stanford CS149 Programming Assignment 4 / Written Assignment 3

PA4 drops you onto a single NeuronCore of an AWS Trainium2 chip. No cache decides what stays on chip: you move data into SBUF (28 MiB) and PSUM (2 MiB) yourself with dma_copy, and the partition dimension tops out at 128. Part 1 teaches those limits and the cost of DMA through vector add and transpose. Part 2 asks you to rewrite convolution as a series of matmuls and fuse it with max pooling so nothing spills back to HBM. Written 3 drills the same idea with a line buffer, two back-to-back box blurs, softmax hardware, and metapipelining: keep intermediates on chip. The environment needs the course's private AMI and a paid capacity block, so for outside readers this assignment is effectively A2.