GPU Glossary
/
Device Hardware·2 min read

Tensor Memory

A specialized, tightly restricted memory in the Streaming Multiprocessor of GPUs like the B200 that stores Tensor Core inputs and outputs.

Tensor Memory is a specialized memory in the Streaming Multiprocessor (SM) of certain GPUs, like the B200. It stores the inputs and outputs of Tensor Cores.

Restricted access#

Access to Tensor Memory is highly restricted. Data must be moved collectively by the four warps of a warpgroup, and they can only:

  • move data in specific patterns between Tensor Memory and registers,
  • write shared memory to Tensor Memory, or
  • issue matrix-multiply-accumulate (MMA) instructions to Tensor Cores that use Tensor Memory for specific operands.

So much for a "compute-unified" device architecture!

Where MMA operands live#

For a tcgen05.mma Parallel Thread eXecution instruction computing D += A @ B with Tensor Memory:

Matrix Location
D (accumulator) must be in Tensor Memory
A (left-hand) may be in Tensor Memory or shared memory
B (right-hand) must be in shared memory, not Tensor Memory

None of the matrices are in registers.

Key idea The rules are complex but not arbitrary. Accumulators are accessed more often than tiles during matmuls, so they benefit more from specialized hardware, e.g. shorter, simpler wiring between the Tensor Cores and Tensor Memory.

Not the Tensor Memory Accelerator#

Tensor Memory is not directly related to the Tensor Memory Accelerator, which loads into the L1 data cache instead.

Roughly speaking, data moves from that cache into Tensor Memory only as a result of a Tensor Core operation. It is then explicitly moved out for post-processing, e.g. the non-linearity after a matrix multiplication in a neural network.

For details on Tensor Memory and how to use it in matrix multiplications, see the Programming Blackwell Tensor Cores with CUTLASS talk from GTC 2025.