Loom

Memory coalescing, one warp at a time

Every GPU performance story starts here.

A warp is a group of 32 threads that run together in lockstep. Each of the 32 loads a word (4 bytes) from global memory. The hardware does not fetch 32 little pieces. It fetches aligned 32-byte sectors and hands each thread its slice.

Read neighbouring words and a whole sector is useful. That is a coalesced load. Stride apart and each sector arrives carrying mostly data nobody asked for.

Start small. Set the warp to 1 thread, press play, then step up to 4 and then 32. Then flip to strided and predict the damage before you reveal it.

Why the strided case hurts

A coalesced 32-lane load needs exactly four sectors. 128 bytes in, 128 bytes used, 100% of the bus doing real work.

The strided load touches a fresh sector for every lane. That is 32 sectors, 1024 bytes moved to deliver the same 128 useful bytes.

Same data, eight times the traffic. Eight times the wait on the slowest memory in the machine.

That one idea, move the bytes you need and no more, is the seed of tiling, cp.async, TMA, and everything Loom builds next.

So who does the waiting?

A global load takes hundreds of cycles.

If a thread just sat and waited, the machine would idle most of the time. It does not.

On an A100 or H100, an SM keeps up to 64 warps resident at once. Its four warp schedulers each issue one instruction per cycle from an eligible warp, one that is not currently waiting on memory.

While one warp is stalled on memory, the scheduler selects another. That is latency hiding, and it is why GPUs want lots of threads.

Switching costs nothing. Every warp's registers stay live in hardware.

Drag the warp count up from 1 and watch the idle time vanish. Amber cells are warps stalled on memory. Teal is the one the scheduler selected this cycle. With too few warps it hits bubbles, cycles where nobody is eligible.

Your turn: kill the bank conflict

Shared memory is fast, but it is split into 32 banks, and a warp can read only one word per bank per cycle. Consecutive 4-byte words are spread across the banks in order, so word w sits in bank w mod 32, wrapping back to bank 0 every 32 words. Two threads that want different words in the same bank are a conflict, and the reads serialize.

Hit 32 different banks and the whole warp reads at once. Collide, and the read serializes into one transaction per pile-up.

Below, 32 threads read smem[thread × stride]. Find the stride that keeps all 32 banks distinct.

Follow the data: a tiled matmul

Now we put the pieces together on the job GPUs are built for: matrix multiply.

Each output number is a row times a column. Done naively, every thread streams a whole row and column from global memory, and neighbours re-read the exact same values.

That is the trap. The math is cheap, but the machine spends all its time waiting on memory. Matmul is memory-bound until you fix the reuse.

The fix is tiling. Load a small tile of A and a tile of B into shared memory once, then let every thread in the block (the group of warps that runs together on one SM and shares its fast memory) reuse them many times before moving on.

Click any output tile in C. It is a running sum. Sweep K and watch it pull one tile of A and one tile of B per step, add them in, and fill up.

Don't wait for the load. Prefetch it.

Tiling cut the traffic. But each tile still has to arrive before you can use it.

The naive loop stalls on every load: fetch a tile, sit idle, multiply, repeat. The slow memory shows up as dead time between every step.

cp.async changes that. It copies a tile from global straight into shared memory asynchronously, without parking the thread and without going through registers. An ordinary load routes each word global → register → shared, tying up the thread and a register the whole time; cp.async skips straight to shared, so the thread can keep working.

So while you multiply the current tile, you prefetch the next one. Two buffers, taking turns. The load hides behind the compute.

Run both schedules on one clock. Same work, same tiles. Watch the pipelined run cross the finish line while the synchronous one is still stalling.

Boss: the tile fights back

The tile is staged in shared memory. Now the threads read a column of it to do the multiply.

Remember the 32 banks from Phase 1? Stepping down a column jumps a full row, 32 elements, each time. With 32 banks, that stride sends every thread in a column to the same bank. That is a 32-way conflict, the worst case.

Fix it with the oldest trick in the book: a little padding.

That is the Ampere playbook: tile for reuse, pipeline with cp.async, and pad to keep shared memory conflict-free. Hopper and Blackwell rebuild this same dataflow with new hardware. That is where Loom goes next.

How many warps actually fit?

Phase 1 sold you on warps: the more you have, the better the scheduler hides latency.

So why not run the maximum every time? Because warps are not free.

Each SM has a fixed budget: one register file and one slab of shared memory. Every thread's registers come out of that file. Every block's shared memory comes out of that slab.

Ask for more registers per thread, and fewer warps fit. Ask for more shared memory per block, and fewer blocks fit. Occupancy is the fraction of the maximum warps you manage to keep resident.

Drag the sliders. Push registers per thread up and watch the warps vanish. The meter that turns amber is the resource that ran out first.

This is the hinge between the last two phases. Latency hiding (Phase 1) wants many warps. Tiling (Phase 2) spends shared memory to get reuse. Both pull on the same fixed budget, and occupancy is where you balance them.

Hopper: let one thread do the load

Phase 2's Ampere loop worked, but look at who was doing the loading. Every one of the block's threads computed its own addresses and fired a cp.async. That address arithmetic burns registers, the very budget Phase 3 showed is scarce.

Hopper adds a dedicated copy engine, the TMA (Tensor Memory Accelerator). Now a single thread hands it a descriptor, a small struct that says where the tensor lives, its shape and strides, the tile to grab, and how to swizzle it (permute the shared-memory addresses so the load lands conflict-free). The engine does all the address math.

When the tile lands, TMA flips an mbarrier, a hardware flag threads can wait on, to wake the waiting threads. The other 127 never touched an address. They are free to compute.

Flip between Ampere and Hopper. Watch the block go from every thread grinding on address math to one thread firing a descriptor and the rest set free.

Split the warps: producer and consumer

With TMA doing the copy, you can give whole warps different jobs. This is warp specialization. One group of warps, the producer, does nothing but fire TMA loads into a multi-stage shared-memory buffer. Another group, the consumer, does nothing but run wgmma, the warpgroup matrix multiply, draining that buffer. A warpgroup is four warps, 128 threads, that together issue one instruction to the tensor cores: the dedicated units inside the SM that multiply a small matrix in a single step, far faster than threads doing it one multiply at a time. Every phase from here runs on them.

An mbarrier per stage hands each tile from producer to consumer. Because the buffer is a few stages deep, the producer can run several tiles ahead. Once it is primed, the tensor cores never wait.

Run the pipeline. The amber lane is TMA loading tiles; the teal lane is wgmma multiplying them. After the buffer fills, the teal lane runs solid to the end. That is a saturated tensor core.

Same GEMM, new machine. The bytes still flow global → shared → tensor core, but Hopper moves the address math into hardware and lets warps specialize. Blackwell goes one step further and moves the accumulator out of the registers too. That is next.

Blackwell: make the numbers smaller

Every phase so far has moved bytes faster. Blackwell also makes each number smaller. FP8 and FP4 are ordinary floating-point numbers squeezed into 8 or 4 bits instead of the usual 16. A matmul in FP8 moves half the bytes of FP16 and runs the tensor cores about twice as fast. FP4 halves it again.

The catch is accuracy: four bits can barely tell numbers apart. The fix is microscaling. Store a block of values in FP4, but share one small scale factor (an 8-bit number, FP8 in NVFP4) across every 16 of them, plus one full-precision scale for the whole tensor. NVIDIA's NVFP4 does exactly this.

Pick a format. Watch how many values fit in one 32-byte chunk and how the tensor-core peak grows as the numbers shrink. For FP4, the amber ticks mark each shared scale.

Move the accumulator out of the registers

Here is the surprise. Feed the tensor cores at FP4 rates and the bottleneck stops being memory. It becomes the register file itself. The accumulator, the running sum D += A × B, cannot be read and written fast enough from registers.

So Blackwell gives it a dedicated home: Tensor Memory, or TMEM. That is 256 KB per SM, with its own bandwidth. The accumulator lives there for the whole matmul instead of clogging the registers.

And the instruction that drives it, tcgen05.mma (an MMA, matrix multiply-accumulate, the tensor core's core operation D = A × B + C), is issued by a single thread on behalf of the entire block. On Hopper 128 threads cooperated on a wgmma; on Blackwell one thread fires the MMA and the tensor core reads straight from shared memory and TMEM. A pair of SMs can even share one MMA, each staging half the operands.

Step the precision down and watch the accumulator's bandwidth demand climb. Past FP8 it crosses the register-file limit, and the accumulator hops to TMEM.

That closes the hardware tour: coalescing, the scheduler, tiling, occupancy, TMA, and now tensor memory. One last idea ties the thread-to-data mapping together, the algebra the kernels are actually written in. That is the capstone: CuTe layouts.

Capstone: the layout is the whole trick

Look back at the journey. Coalescing was about which address each thread reads. Tiling was about where a tile sits in shared memory. Swizzling and padding reshuffled that placement to dodge bank conflicts. TMA carried a descriptor of it. Every one of those is the same thing: a mapping from a logical coordinate to a memory offset.

CuTe, the foundation of CUTLASS (NVIDIA's library for fast matmul kernels), writes that mapping as one object: Layout = Shape ⊗ Stride. The shape says how big each axis is. The stride says how far apart its elements sit. Feed it a coordinate (i, j) and it returns an offset, i·strideᵢ + j·strideⱼ. That is the entire idea.

The power is that shape and stride are independent. Keep the same logical grid and change only the strides, and the data is suddenly row-major, column-major, or padded, without a single element moving.

Flip the layout between row-major and column-major and watch every offset renumber. Then choose whether a warp reads a row or a column, and the coalescing verdict from Phase 0 drops straight out of the stride.

This is why real kernels are written in CuTe. Tensor cores demand exact, swizzled thread-to-data mappings for TMA and wgmma, and hand-deriving them is a footgun. The layout algebra composes and tiles them at compile time, portably across Ampere, Hopper, and Blackwell.

You now have the whole machine and the language its kernels speak. The next track flips from the hardware to the software: how production kernels like FlashAttention, DeepGEMM, and the MoE libraries put every one of these ideas to work. It starts with the one picture that governs all of them, the roofline.

The payoff: one picture for every kernel

You have the machine. Now the question every real kernel asks: am I limited by the math or by the memory? The roofline answers it with a single plot.

The x-axis is arithmetic intensity, how many FLOPs (floating-point operations, the multiplies and adds) you do per byte you move. The roof has two parts. A sloped bandwidth roof on the left: if you touch memory a lot per flop, bandwidth caps you. A flat compute roof on the right: once you reuse enough, the tensor cores cap you. The corner where they meet is the ridge.

Every optimization in Loom was a fight to move rightward. Coalescing and fusion cut the bytes. Tiling multiplied the reuse. Watch it happen: naive matmul starves down on the diagonal, and tiling lifts the very same operation onto the flat roof.

Drag the intensity slider, or jump to a named op. Left of the ridge is memory-bound and slow. The goal of nearly every kernel trick is to climb the slope and reach the flat.

Two of those moves opened up whole topics of their own. The narrowing numbers became a zoo of formats: FP8, FP4, NVFP4, MXFP4, and the GGUF quants that run models on your laptop. And the matrix unit became the engine behind real kernels like FlashAttention. Each has its own page.