Computational Graphs and Backpropagation This note explains how to compute gradients for any function by breaking it into a graph of simple operations. It is the bridge between the gradient-descent picture from the linear and logistic regression note and the layered functions we will later call neural networks. The ideas are: Draw the function as a graph of operations. Evaluate the graph from inputs to output: the forward pass. Use the chain rule to carry sensitivities from the output back to the inputs: the backward pass. We build this on one tiny example and walk through every step. ...

ML Refresher: Linear and Logistic Regression This is the first note in the ML → Deep Learning → Transformers → LLMs series. The goal is to rebuild working memory of the basics before we get to neural networks: what a model is, how a loss function measures error, and how gradient descent tunes parameters. We will implement linear regression and logistic regression from scratch in NumPy, then compare with scikit-learn. If the code and gradients feel obvious, you are ready for the next note (computational graphs and backprop). If not, this is exactly the foundation to lock down first. ...

CUDA Profiling and Performance Measurement The notes up to this point covered what makes CUDA kernels fast in theory: coalescing, tiling, occupancy, and asynchronous execution. This note covers how to measure whether a kernel is actually fast and how to find the bottleneck when it is not. Profiling is what separates guessing from optimization. Without it, you can apply every best practice and still be slower than a simpler implementation. ...

CUDA Unified Memory: One Pointer, Two Processors The previous note covered explicit CUDA streams and how to overlap host-device transfers with computation. This note covers Unified Memory, an alternative to explicit cudaMalloc / cudaMemcpy / cudaFree. It lets both CPU and GPU use a single pointer, at the cost of some performance complexity under the hood. Unified Memory is seductive: it removes the need to think about host and device pointers separately. But it does not remove the physical reality that CPU and GPU memory are separate. Understanding when it helps and when it hurts is essential. ...

CUDA Reduction: Parallel Sum, Max, and Other Tree Algorithms The previous note covered occupancy. This note covers reduction: taking a large array and producing a single value (a sum, max, min, product, or any associative binary operation). Reduction is one of the most important parallel algorithms on a GPU, and it teaches several key CUDA ideas at once: shared memory, thread cooperation, warp-level execution, and multi-kernel launches. This note builds the algorithm step by step, from a naive atomic version to a tree-based shared-memory reduction. ...

CUDA Streams and Host/Device Overlap The previous note covered reduction, the first parallel algorithm on the GPU side. This note moves to the host-device boundary: how to keep the GPU busy while the CPU prepares the next batch of data, and how to overlap data movement with computation. So far, every kernel launch and every cudaMemcpy has been synchronous. The host waits, the GPU works, the host waits again. CUDA streams let you break that pattern. ...

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