GPU Glossary
/
Device Software·2 min read

Warpgroup

A set of four contiguous warps (128 threads), aligned so the first warp's rank is a multiple of 4, used for warpgroup-level instructions.

A warpgroup is a set of four contiguous warps where the warp-rank of the first warp is a multiple of 4. Warpgroups let a single instruction drive larger matrix multiplications, which is key to keeping modern Tensor Cores busy.

Why warpgroups exist#

A warpgroup-level instruction coordinates 128 threads: 4 warps per warpgroup × 32 threads per warp. Working at this larger granularity removes the need for explicit synchronization between warps and lets each instruction handle a larger problem, especially larger matrix multiplications. Larger matrix multiplications more easily saturate the huge arithmetic bandwidth of the Tensor Cores in recent data center GPUs.

Warpgroups were introduced in NVIDIA's Hopper Streaming Multiprocessor architecture to support warpgroup-level matrix multiplication instructions such as wgmma.mma_async. This Colfax blog post is a deep dive. Warpgroups are prominent in how high-performance Hopper and Blackwell kernels organize their pipeline components, for example in Flash Attention 4.

Key idea A warpgroup is 128 threads issuing one matrix instruction together, sized to feed the Tensor Cores.

Computing warp-rank#

In the Parallel Thread Execution (PTX) IR, a warp's rank is:

int linearIdx = (%tid.x + %tid.y * %ntid.x  + %tid.z * %ntid.x * %ntid.y);
int warpRank = linearIdx / 32;

Here tid is the thread index, read from special PTX registers.

So for an 8-warp dispatch, the valid warpgroups are:

  • Warpgroup 0: warp-ranks 0, 1, 2 and 3
  • Warpgroup 1: warp-ranks 4, 5, 6 and 7

The alignment restriction#

To our knowledge, NVIDIA doesn't document why warpgroups must be aligned on warp-rank. However, the Streaming Multiprocessors of recent data center GPUs appear to contain four (unnamed) subunits, each with its own Warp Scheduler and Tensor Core.