CMU 11-868's three GPU lectures answer one question: why does a correct CUDA matmul use only 2.48% of an A100's FP32 compute? L02 covers SMs, warps, and the grid/block/thread hierarchy. L03 covers cudaMalloc, cudaMemcpy, and kernel indexing. L04 uses tiling, coalesced access, and bank-conflict avoidance to bring data closer than global memory, which sits about 500 cycles away.
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.
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.
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.
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?
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.
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.
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.
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.
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.
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.
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.
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.
The last programming assignment in CS149 Fall 2025 is open-ended. Pick at least one of five kernels (Histogram, a 1D occupancy decoder, FlashAttention, a 3D heat equation with RK4, SwiGLU) and make it faster than its PyTorch baseline on an H100. You can write CUDA, Triton, or TileLang, and you may use LLMs. There is no speed threshold. The grade depends on a work log that shows what you measured at each step, what hypothesis you formed, and why you stopped. The H100 job queue and leaderboard need a SUNet ID; outside Stanford you can only run eval.py on your own NVIDIA GPU.
L4 lays out a thought process for parallelizing code: decompose the problem to find independent work, assign that work to workers, orchestrate communication and synchronization, then map workers to hardware. Amdahl's Law reminds you that the sequential fraction caps speedup. The running example is a 2D grid solver whose original dependencies are hard to exploit; switching to a red-black update order makes it expressible in either a data-parallel or a shared-address-space model.
L11 asks what programmers pay once hardware specializes for AI. On the H100, saturating Tensor Cores means 16×16 tiles, TMA moving data asynchronously, and producer and consumer warps running as a pipeline. That's hard to write, which is why DSLs like ThunderKittens exist. The other route is a dataflow architecture (SambaNova SN40L): describe the computation with parallel patterns such as map, reduce, and zip, and let the compiler handle tiling, metapipelining, and placement. The slides say this can fuse an entire Llama 3.1 8B decoder layer into one kernel.
The first lecture of CS149 Fall 2025 defines speedup, then uses three classroom demos to show how communication and load imbalance eat into it. Next it explains why single-core performance stalled: superscalar execution runs out of instruction-level parallelism at about four instructions per clock, and clock frequency hits the power wall. So performance now has to come from more cores and specialized hardware. The last part turns to efficiency. A DRAM access takes about 60 times as long as an L1 cache hit, and moving 64 bits costs over a thousand times the energy of an integer op. Efficiency almost always comes down to accessing data efficiently.
Load balancing is hard because it pulls against scheduling cost: smaller tasks balance better, but every task grab pays a synchronization cost. CS149 Lecture 5 first lays the options out as a continuum from static to dynamic, then takes apart the Cilk runtime: one deque per worker, run the child at a spawn and leave the continuation for others to steal, and idle threads steal the biggest chunk of work from the top of someone else's deque.
This post covers Fall 2026 material, not the Fall 2024 edition the rest of the series follows. Fall 2026 replaced the pruning lab with "Efficient AI Fundamentals" (lab1_gpu_basics.zip). Part 1 has you hand-write a triple-loop GEMM and compute MAC, FLOPs, and I/O. Part 2 plots GEMM and GEMV rooflines. Part 3 works through a gemma-3-270m-it decoder layer, computing attention and MLP costs and comparing prefill with decode. Part 4 uses the PyTorch Profiler to inspect kernels, has you write GeLU to feel kernel fusion, then tries torch.compile and CUDA Graphs. Part 5 compares SDPA with FlashAttention. The core is 80 points plus 20 bonus, and all of Part 5 became bonus because Colab's T4 can't run it.
Once algorithms have shrunk the model, how much more can the system layer squeeze out? Lecture 11 uses a single matrix multiply to show it: loop reordering gives 12x, tiling 19x (on an Intel Xeon 4114), and a CUDA version runs 94x faster end to end on a 2080Ti. The second half covers TinyEngine's inference tricks: im2col; in-place depthwise, which cuts peak memory from 2×C×H×W to (1+C)×H×W; NHWC for pointwise and NCHW for depthwise; and Winograd, with 2.25x fewer multiplications.
Parsing a 150-page PDF via Vision API took 29 minutes (one page per request). Batching cut it to 1.5×, adding 5-way concurrency brought it down to 24 seconds. One week after launch, a customer uploaded 23 PDFs at once — 70 parallel Bedrock requests triggered full throttling: 90 pages skipped, 5 files failed, 8 stuck. Fixed with Redis-based cluster-wide slots + backoff retries.
Cloudflare Cache Rules are zone-level cache policy: request expressions decide what is eligible for cache, how Edge TTL and Browser TTL behave, what dimensions enter the cache key, and how stale content, ETags, and purge interact. Use them for CDN cache policy; use the Worker Cache API for programmatic caching.
Cloudflare Images has two paths: transform images stored in R2/S3/origin at the edge, or store images in Images and deliver named variants. The first is priced by unique transformations; the second also involves stored and delivered images.
Smart Shield is Cloudflare's origin protection bundle: Smart Tiered Cache, connection reuse, Argo Smart Routing, Regional Tiered Cache, Cache Reserve, Health Checks, and Dedicated CDN Egress IPs reduce requests and connections reaching your origin.
Lecture 5 explains GPUs through SMs, warps, and the memory hierarchy, then unifies common optimization under low precision, fusion, recomputation, coalescing, and tiling. FlashAttention combines those principles for attention.
Lecture 6 turns GPU principles into kernels: benchmark scaling across shapes, profile actual calls and time, then implement GeLU, softmax, reductions, and tiled matrix multiplication in Triton. Speed begins with measuring correctly.
Lecture 2 reduces model training to tensors, FLOPs, bytes, and time: use einops to track dimensions, arithmetic intensity and roofline analysis to identify bottlenecks, then trade compute for memory with gradient accumulation and activation checkpointing.
CS107 Lecture 25 builds the essential cache model from a concise deck: memory access costs are nonuniform, smaller and faster layers retain data likely to be reused, and temporal and spatial locality determine whether a program benefits.
RAG system costs come from LLM tokens, Embedding APIs, and vector search. Every stage has room for cost reduction, but you need to verify that optimizations don't sacrifice too much quality.
Caching doesn't have to match exact query strings -- semantically similar questions can hit the cache too, skipping the entire RAG pipeline execution.