Warp specialization and TMA get credited every time a benchmark table shows FlashAttention-3 or FlashAttention-4 beating its predecessor, but most explanations stop at naming the terms. The GPU producer-consumer kernel pattern is the mechanism underneath both: split a kernel's warps into two roles, one that only moves data and one that only computes on it, and let a dedicated hardware copy engine keep the pipeline full so neither role ever waits idle on the other. This post works through that pattern from first principles: what a Tensor Memory Accelerator actually is, how warp specialization turns a synchronous stall into an overlapped pipeline, and why that overlap, not raw FLOPS, is the real reason FlashAttention-3 and FlashAttention-4 outperform FlashAttention-2. For the tiling and online-softmax math this pattern sits on top of, see our FlashAttention explainer; this post picks up where that one stops.
TL;DR: What Is the GPU Producer-Consumer Kernel Pattern?
- Definition: The GPU producer-consumer kernel pattern splits a thread block into producer warps that issue async copies and consumer warps that run tensor core math, overlapping the next tile's load with the current tile's compute.
- Hardware: Hopper's Tensor Memory Accelerator (TMA) makes the split possible: one thread issues an async copy of an entire tile, freeing every other thread to compute.
- Hopper: FlashAttention-3 reaches up to 740 TFLOPs/s in FP16 on H100, versus FlashAttention-2's roughly 35% GPU utilization on the same chip.
- Blackwell: FlashAttention-4 adds Tensor Memory (TMEM) and hits up to 1,613 TFLOPs/s in BF16 on B200.
- Spheron rents bare-metal H100 GPU instances with the performance-counter access this pattern needs to profile.
What a Tensor Memory Accelerator Actually Does (Async Copy, Not Compute)
A Tensor Memory Accelerator is not a compute unit. It's a dedicated hardware engine, introduced with Hopper's SM90 architecture, whose only job is moving tensor data between global memory and shared memory asynchronously, so the streaming multiprocessor's compute pipelines never have to stall on address calculations or a manual copy loop. NVIDIA's own CUDA Core Compute Libraries documentation is direct about the scope of the job: "The Tensor Memory Accelerator (TMA) is a hardware feature available on Hopper (SM90) and newer GPUs that enables efficient asynchronous memory copies of tensor data between global and (cluster) shared memory."
TMA operates on whole tensor tiles, not individual elements. It performs asynchronous copies of one-dimensional through five-dimensional tensors between global memory and shared memory, including cluster-shared memory across the multiple SMs in a Hopper thread block cluster. That range matters for attention kernels specifically, where the tiles being staged are the Q, K, and V blocks that later feed a matrix multiply, and getting all of a tile's dimensions moved by one hardware-issued copy is what lets the rest of the warpgroup skip per-element bookkeeping entirely.
One Thread Issues the Copy, 127 Others Keep Computing
The mechanic that makes TMA a genuine hardware shift, not just a faster instruction, is who issues it. As NVIDIA's Hopper architecture deep dive describes it, only one thread in a warp or cooperative thread array issues the TMA operation, passing it a tensor map descriptor that describes the tile's shape and location; the hardware then handles all the address generation and the actual data movement itself, freeing the other threads in that warp, and the other warps in the warpgroup, to keep computing or to wait on a barrier instead of participating in the copy at all. In a 128-thread Hopper warpgroup, that means one thread's instruction stages a tile for the other 127.
Completion isn't polled. As Colfax Research's Hopper TMA tutorial explains, TMA signals that a copy has landed through a shared-memory async barrier, an mbarrier object, which is what lets a consumer warp simply wait on the barrier for a tile to be ready instead of spinning a loop checking a flag. That barrier-based signaling is also what makes the producer-consumer split composable across multiple pipeline stages, since a producer can be several tiles ahead of the consumer it's feeding without either side needing to poll the other's progress directly.
Why TMA Replaces Manual cp.async and Per-Thread Address Math
Before TMA, the standard way to move a tile asynchronously (Ampere's cp.async) still required every participating thread to compute its own global memory address and issue its own copy instruction for its own slice of the tile. That per-thread address math costs registers, and register pressure is a direct input to how many warps an SM can keep resident at once. TMA collapses an entire tile's worth of per-thread address generation into one descriptor and one issuing thread, which is exactly the register headroom FlashAttention-3's producer warpgroup spends on nothing but issuing loads, and exactly what its consumer warpgroups reclaim to run bigger matrix multiplies instead. If you've hit an occupancy ceiling before, this is the same register-pressure trade-off, just moved from your own kernel code into hardware.
Warp Specialization: How CUDA Builds the GPU Producer-Consumer Kernel Pattern
Warp specialization means assigning different warps, or in Hopper's case different warpgroups of four warps each, a single fixed job for the entire kernel launch, instead of every warp running the identical load-then-compute sequence in lockstep. It's worth distinguishing this deliberately from warp divergence, which is a different phenomenon entirely: divergence is an unplanned cost that happens when threads inside one warp disagree about which branch to take and the hardware serializes both paths. Warp specialization is the opposite move on purpose, giving entire warpgroups different jobs by design so each one runs a single, uniform instruction stream with no branching at all. If you're writing a kernel from scratch, this is the design decision that turns a synchronous load-compute loop into the overlapped pipeline the rest of this post describes.
The GPU Producer-Consumer Kernel Pattern, in Warpgroup Terms
FlashAttention-3 implements the pattern directly in its kernel structure. One warpgroup is dedicated as the producer: it issues TMA loads of the Q, K, and V tiles and does no matrix multiplication at all, so it deallocates its own registers using setmaxnreg since it has no arithmetic to hold values for. The remaining warpgroups are consumers: they reallocate those freed registers to run WGMMA, Hopper's warpgroup-wide matrix multiply-accumulate instruction, on the tiles the producer already staged. The FlashAttention-3 paper frames the whole design choice as a hardware-driven one: "Asynchrony is a result of hardware specialization to accelerate the most important operations in a ML workload: specific hardware units performing matrix multiplication (Tensor Cores) or memory loading (Tensor Memory Accelerator - TMA)."
The producer and consumer don't work tile by tile in lockstep either. FlashAttention-3 uses a circular shared-memory buffer with multiple pipeline stages, so the producer warpgroup can already be loading stage j % s of the buffer while the consumer warpgroups are still processing an earlier stage, overlapping the copy for one tile with the compute on a previous one instead of serializing the two. That circular-buffer structure is close to what a hand-written CUTLASS kernel looks like under the hood; our CUTLASS CuTe DSL guide walks through a TMA-based GEMM built the same way, in Python.
Ping-Pong Scheduling: Hiding a 256x Throughput Gap Between Softmax and GEMM
The producer-consumer split alone doesn't fully hide FlashAttention's other bottleneck: softmax and GEMM run on entirely different parts of the SM, at wildly different speeds. The FlashAttention-3 paper puts numbers on the gap: on H100, the exponential inside softmax runs on the special function unit at roughly 3.9 TFLOPS, while the GEMM runs on tensor cores at roughly 989 TFLOPS in FP16, a gap of around 256x. Run those two operations back to back on the same warpgroup and the tensor cores sit idle for however long the exponential takes.
FlashAttention-3's answer is ping-pong scheduling: it alternates two consumer warpgroups between the GEMM role and the softmax role across pipeline stages, so one warpgroup's softmax work overlaps with the other warpgroup's GEMM instead of each warpgroup doing both in sequence. Combined with the producer-consumer split for memory, this is what closes most of the remaining gap between FlashAttention-3's measured throughput and H100's theoretical peak. Our FlashAttention-2 vs FlashAttention-3 benchmark guide has the full H100 and H200 throughput tables this scheduling produces at different context lengths.
Why This Is the Real Reason Hopper and Blackwell Beat Ampere on Attention
The headline TFLOPS number on a spec sheet describes what a GPU can theoretically do, not what a kernel actually gets out of it, and if you're choosing between generations on that number alone, you're comparing the wrong thing. The real story behind FlashAttention-3 and FlashAttention-4's generational gains is that they close the distance between those two numbers, and they close it almost entirely through overlap: keeping tensor cores fed continuously via an async producer-consumer pipeline, rather than through any change to the attention math itself.
FlashAttention-2's Ceiling: 35% GPU Utilization Without Async Hardware
FlashAttention-2 is a genuinely hand-tuned kernel, and the FlashAttention-3 paper reports it still reaches only about 35% GPU utilization on H100, because it has no async hardware to build a producer-consumer split on top of. Without TMA and WGMMA, a warp that needs a tile has to load it and wait, then compute, then load the next one; there's no cheap way to have one set of warps stay several tiles ahead of another. FlashAttention-3, running the same class of attention computation on the same H100 hardware but built around TMA, warp specialization, and ping-pong scheduling, reaches up to 740 TFLOPs/s in FP16 (about 75% of the GPU's theoretical peak) and roughly 1.2 PFLOPs/s in FP8, a 1.5-2.0x speedup over FlashAttention-2's forward pass. Same silicon, same attention formula; the difference is entirely in how continuously the tensor cores stay busy.
What Changes on Blackwell: Tensor Memory (TMEM) and FlashAttention-4
Blackwell moves the pattern one layer further into hardware. According to the FlashAttention-4 paper, on Blackwell tensor core MMA instructions write their output directly and asynchronously to a new on-chip memory called Tensor Memory (TMEM), 256KB per SM, instead of to the register file the way Hopper's WGMMA does. That removes accumulator values from the register-pressure budget entirely, which is a second, independent source of headroom on top of what the producer-consumer split already frees up.
Blackwell's MMA instruction also processes larger tiles, 128xN instead of Hopper's 64xN, and FlashAttention-4's forward pass takes advantage of that by using two 128-thread warpgroups where each thread handles an entire row of the tile, which removes the inter-warp shuffles FlashAttention-3 needed on Hopper to reassemble a full row's worth of data across warps. The same paper reports the combined effect on B200 is up to 1,613 TFLOPs/s in BF16 (71% of peak hardware utilization), 1.3x faster than cuDNN 9.13 and 2.7x faster than a Triton implementation of the same kernel. Our FlashAttention-4 Blackwell guide covers the full migration path, including which Blackwell SKUs support it and which fall back to FlashAttention-2.
Reproducing the Pattern: Profiling a Warp-Specialized Kernel Yourself
Seeing the producer-consumer split in a profiler is more convincing than reading about it. Nsight Compute (ncu) reports separate warp-state and instruction-mix breakdowns for a producer warpgroup versus a consumer warpgroup in the same kernel launch; a healthy producer shows almost nothing but memory-issue and barrier-wait states, while a healthy consumer shows tensor-core issue slots staying busy across the run instead of alternating with long stalls. That level of detail requires hardware performance counter access, which most shared or serverless GPU platforms restrict for tenant isolation, the same limitation that blocks profiling any kernel-level effect in detail on that kind of platform.
A multi-hour kernel-profiling and iteration session is also a bad fit for spot capacity, since a reclaimed instance mid-run loses whatever ncu state was in progress; an on-demand instance is the honest choice here, not spot. Spheron's pricing page states its instances bill "with per-minute billing granularity, so you only pay for the exact time you use," with "no minimum rental period." That combination, bare-metal counter access plus per-minute billing, means a single short profiling session on a Hopper card or a Blackwell B200 instance doesn't require committing to a monthly contract on either generation.
Profiling a warp-specialized attention kernel means holding privileged Nsight Compute counters for hours, not minutes, which per-minute bare-metal billing is built for rather than against.
Frequently Asked Questions
It's a kernel design where a CUDA thread block's warps are split into two fixed roles instead of every warp doing the same load-then-compute sequence. Producer warps only issue asynchronous memory copies (via the Tensor Memory Accelerator on Hopper and newer GPUs); consumer warps only run tensor core math on the tiles the producers already staged. Because the two roles run concurrently on the same SM, memory movement for tile N+1 overlaps with compute on tile N, instead of the kernel stalling on each load before it can compute.
TMA is a dedicated hardware copy engine, not a compute unit. A single thread issues a TMA instruction with a tensor descriptor describing a multidimensional tile; the hardware then generates every address and moves the entire tile between global memory (HBM) and shared memory asynchronously, signaling completion through a shared-memory barrier. This replaces the older pattern where every thread in a warp computed its own address and issued its own copy instruction.
Warp specialization means different warps (or warpgroups) inside the same thread block are assigned different, fixed jobs for the kernel's whole lifetime, rather than all warps executing identical code. In FlashAttention-3, one warpgroup is a dedicated producer that issues TMA loads and does no math, and the remaining warpgroups are consumers that run WGMMA matrix multiplies on the data the producer staged.
Softmax's exponential runs on the SM's low-throughput special function unit while the GEMM runs on tensor cores, and the two units have very different peak throughput on H100. Ping-pong scheduling alternates two consumer warpgroups between the GEMM and softmax roles across pipeline stages, so one warpgroup's softmax work overlaps with the other warpgroup's GEMM instead of the two executing back to back on an idle tensor core.
Blackwell's tensor cores write MMA output directly and asynchronously to a new on-chip memory, Tensor Memory (TMEM), instead of to the register file the way Hopper's WGMMA does. Blackwell's MMA instruction also operates on 128xN tiles instead of Hopper's 64xN, which lets FlashAttention-4 use two 128-thread warpgroups where each thread owns a full row, removing the inter-warp shuffles FlashAttention-3 needed on Hopper.
