Rendering in Retrospect · #12 of 12

My Next Study Plan: CUDA, Kernels, and Attention

December 23, 2026

Retrospective — written in 2026, looking back.

Starting Again, On Purpose

Every post in this series so far has been retrospective — knowledge I already have, organized. This last one is different. It is a study plan, written from a beginner's position, about the area I want to learn next: GPU compute kernels, in the direction of modern attention implementations.

Let me be plain about where I stand. I am a graphics programmer. I know GPU parallelism from the rendering side — shaders, dispatch shapes, barriers, memory spaces as they appear in graphics APIs. I do not yet know CUDA well. I have not optimized a GEMM beyond calling a library. I have not written anything serious in Triton. I could not implement FlashAttention today, and I only partly understand why it is fast. This post is the plan to change that, and an honest list of what I don't yet understand.

Why This Direction, for a Graphics Programmer

Rendering and machine learning compute look different at the surface and surprisingly similar underneath. Both are dominated by moving data through a memory hierarchy without stalling arithmetic units that are, on modern hardware, comparatively plentiful. The questions I have learned to ask in graphics — where does the data live, what is the access pattern, what is actually the bottleneck, bandwidth or math — are the same questions kernel optimization asks, with different vocabulary.

There is also a practical motivation. Attention is the defining workload of the current moment in computing, and it is, at its core, a sequence of matrix multiplications and a softmax, arranged cleverly. Understanding it from the kernel level up feels like the right way for someone with my background to understand what these systems actually do, rather than treating them as library calls. I want the mechanism, not the pitch — the same standard I tried to hold in the earlier posts in this series.

The Order: CUDA Basics, Then GEMM, Then Triton, Then Attention

The sequence matters, and here is my reasoning for it.

Stage one: CUDA basics. Before optimizing anything, I need the execution model in my hands: threads, blocks, grids, warps; global, shared, and register memory; what a coalesced access actually looks like in code; how occupancy is reasoned about rather than guessed at. The exercises here are deliberately small — vector addition, a stencil, a reduction — because each one isolates a concept. The canonical first kernel is a rite of passage:

// The textbook starting point — the first kernel in nearly every CUDA
// introduction. Writing it is stage one, not an achievement.
__global__ void VecAdd(const float* a, const float* b, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = a[i] + b[i];
}

What I expect to get wrong at this stage: launch configuration as superstition, and underestimating how much performance lives in memory access patterns rather than instruction choice. I also note, honestly, that my everyday machines complicate this stage: CUDA needs NVIDIA hardware, so some of this learning will happen on rented cloud GPUs or a borrowed machine, with the small local experiments done through alternative backends where possible. Pretending that constraint away would make the plan fiction.

Stage two: GEMM. Matrix multiplication is the atom of the whole field, and optimizing one is a complete education in miniature: tiling into shared memory, register blocking, double buffering, vectorized loads, and watching performance climb in steps as each bottleneck is removed and the next one reveals itself. My goal is modest and specific — write a naive GEMM, profile it, then improve it stage by stage toward something respectable compared to a library baseline, understanding each step's reason. I expect this to take the longest of any stage, and I expect the profiling discipline from the "measure first" principle to be the main transferable skill I bring with me.

Stage three: Triton. Triton promises GPU kernels written in a Python-embedded language at a higher level of abstraction — block-level rather than thread-level thinking. Learning it after CUDA is a deliberate choice: I want to know what the abstraction is abstracting. If I start with Triton, its tiling and masking model will feel like magic syntax. After hand-writing tiled CUDA, I hope it will feel like leverage. The project for this stage is to reimplement the stage-two exercises — especially a tiled matmul — in Triton and compare both the code size and the performance honestly.

Stage four: attention, and FlashAttention as the target text. Only at this point does attention itself enter. First a plain, naive attention implementation from the definition — scores, softmax, weighted sum — materializing the full score matrix and feeling why that hurts at long sequence lengths. Then FlashAttention as a paper-and-code study: the tiling of the computation so the score matrix never fully lands in high-bandwidth memory, the online softmax trick that makes that tiling possible, and recomputation in the backward pass as a deliberate trade of math for bandwidth.

What I Do Not Yet Understand

An honest study plan names its gaps. Mine, as of writing:

• I do not have intuition for how much of attention's cost at various sequence lengths is bandwidth versus compute, and I suspect my graphics-trained guesses are wrong in interesting ways.
• Online softmax makes sense to me algebraically on paper, but I cannot yet write it from memory or explain fluently why the rescaling it requires is numerically safe.
• I do not understand kernel fusion decisions — when combining operations into one kernel wins, and when the register pressure or occupancy cost makes it lose.
• Warp-level primitives, asynchronous copies, and the newer hardware features that recent FlashAttention versions exploit are vocabulary to me, not tools.
• I have no calibrated sense of what "fast" is for any of these kernels — no internal table of expected throughput to be surprised against, which in graphics I rely on constantly.

Each gap maps to a question the stages above are designed to answer by construction, not by reading alone.

The Learning Projects

Reading papers will not be enough; it never has been for me. The plan's projects, in order: (1) the small CUDA concept kernels — add, stencil, reduction — each profiled, not just run; (2) a GEMM improved in documented stages, with a short write-up per stage of what the bottleneck was and what evidence showed it; (3) the same matmul in Triton, compared directly; (4) a naive attention forward pass, correct against a reference implementation on small inputs; (5) a tiled attention attempt informed by the FlashAttention paper — success defined as correct and measurably less memory-hungry than my naive version, not as rivaling the real implementation, which is the work of specialists over years.

Each project also gets the treatment this series taught me: fixed test cases, recorded measurements with their conditions, and notes on where my mental model broke. If those notes are worth anything, they may become posts of their own someday — written, that time, with hindsight I do not currently have.

Closing the Series

"Rendering in Retrospect" began as a look back at lessons from 2021 onward: measure first, follow the data, budget everything, know where models break, and treat taste as a skill. It feels right to end it facing forward, as a beginner in an adjacent field, carrying those same five principles into territory where I do not yet know the edges of anything. The principles transfer. The expertise does not — that has to be earned again, one profiled kernel at a time. I am looking forward to it.