CUDA Occupancy: Filling the SM to Hide Latency The previous note showed how to use shared memory to fix uncoalesced access patterns. This note covers occupancy: how many warps can live on an SM at the same time, and why that number determines whether the GPU can hide memory latency. Occupancy is the first CUDA performance topic that is not about memory access at all. It is about keeping the warp schedulers busy. ...

CUDA Shared Memory Tiling The previous note showed that coalesced memory access is the key to memory-bound kernel performance. But some algorithms naturally want uncoalesced access patterns. The fix is shared-memory tiling: load data into fast on-chip SRAM in a coalesced way, rearrange it there, then write it back to global memory in a coalesced way. This note explains the tiling pattern, why __syncthreads() is essential, the bank-conflict problem, and how to choose tile sizes. ...

CUDA Memory Coalescing: Why Access Patterns Matter The previous note showed how to write and launch a CUDA kernel. This note answers the first performance question: why does the same kernel sometimes run 5× or 10× slower just because of how threads read memory? The answer is memory coalescing. Because GPU memory is optimized for bandwidth over latency, the hardware rewards access patterns where threads in a warp read or write contiguous addresses together. If they do not, the GPU wastes bandwidth and cycles fetching data that most threads ignore. ...

CUDA Programming Model: Writing and Launching Kernels The previous note covered the hierarchy: thread, warp, block, grid, and how those map to SMs. This note turns that hierarchy into actual code. By the end, you will have seen a complete CUDA program, understood every line, and know how a host program hands work to the GPU. CUDA is an extension of C/C++. Most of the code you write is ordinary C++. A small number of CUDA-specific pieces — kernel functions, the launch syntax, a few memory APIs, and built-in thread indices — turn a sequential program into a massively parallel one. ...

CUDA Thread Hierarchy: Grids, Blocks, Warps, and Threads The previous note mapped the hardware: SM, warp scheduler, register file, shared memory, L2, device memory. This note maps the software abstraction CUDA exposes on top of that hardware. The two maps fit together almost one-to-one, and once you see how, most of CUDA stops being arbitrary syntax and becomes named hardware concepts. CUDA organizes parallel work into four nested levels: Grid └── Block (threads in a block can share fast on-chip memory, and sync) └── Warp (32 threads, executed in lockstep on the hardware) └── Thread (your kernel code, from one thread's point of view) This note works from the bottom up: start with the thread (the thing you actually program), then warp, then block, then grid, then show how the whole tower maps onto the GPU die. ...

Anatomy of a GPU: Hardware Components The previous note covered why GPUs look the way they do — throughput over latency, thousands of simple cores instead of a few smart ones. This note zooms into the actual silicon: what physically sits on a GPU die, what each piece is called, and what job it does. The goal is to have concrete hardware nouns (SM, warp scheduler, register file, L2, …) in hand before those same nouns start showing up as CUDA concepts (threadIdx, __shared__, occupancy, …). ...

CPU vs GPU Architecture Why GPUs exist For decades, single-thread CPU performance improved “for free” — you write the same sequential code, and it runs faster on the next generation of chips, because clock speeds kept climbing. Around the mid-2000s, that stopped. Clock speeds hit a power wall — you can’t keep cranking frequency without the chip melting. So the industry pivoted from “make one core faster” to “put more cores on the chip.” This is the multicore/manycore shift, and it’s the reason parallel programming stopped being a niche HPC skill and became something every programmer eventually runs into. ...

Page Frame Reclaiming The Page Frame Reclaiming Algorithm 1. Unified Page Cache & Dual-Caching Edge Case Unified Page Cache: Linux no longer separates the “disk/buffer cache” from the “page cache.” All file and block I/O goes through a single unified Page Cache indexed by (address_space, offset). The Dual-Caching Exception: Reading via a regular path (e.g., /app/data.db) uses the file inode’s address_space. Reading via the raw disk node (e.g., /dev/sda1) uses the block device inode’s address_space. Result: The exact same physical 4 KB block on disk can reside twice in RAM under two different address_space objects. Practical Impact: Rare in normal operation because accessing raw block nodes requires root privileges, and system tools handle raw access using unbuffered I/O (O_DIRECT). ...

RISC-V (pronounced “risk-five”) is an open, royalty-free RISC ISA originally designed to support academic research and education, and now increasingly used in commercial silicon. Unlike MIPS, ARM, or x86, RISC-V is not owned by any single company — its specification is maintained by RISC-V International, and anyone is free to implement it without paying licensing fees. This openness, combined with a deliberately small and modular base ISA, has made it popular in everything from microcontrollers to research CPUs to (increasingly) high-performance application processors. ...

Pipelining — First Principles What is pipelining? Pipelining is the technique of overlapping the execution of multiple instructions. An instruction consists of multiple pieces of work. Instead of completing one instruction entirely before starting the next, we divide the work into stages and let different instructions occupy different stages simultaneously. For a simple example: IF ID EX MEM WB Instruction 1 → → → → → Instruction 2 → → → → Instruction 3 → → → Instruction 4 → → At any moment, several instructions are being worked on, but each is at a different stage. ...