Die WMMA-Fragment-API

To program Tensor Cores directly you use WMMA, the warp matrix multiply-accumulate API in the nvcuda::wmma namespace. 🧩

WMMA is warp-wide: all 32 threads in a warp work together on one tile. You think in tiles, not in single threads.

A fragment is the WMMA data type holding one warp's slice of a matrix tile. Each thread quietly owns a few of its elements.

You declare fragments tagged matrix_a, matrix_b, or accumulator. The tag tells WMMA how that tile feeds the multiply-accumulate.

Fragments use set tile sizes like 16 by 16 by 16. You pick a supported shape; the hardware only accepts those combinations.

First you call load_matrix_sync to pull a tile from memory into a fragment. This load spreads the data across the warp for you.

Before adding, you zero the result tile with fill_fragment. A clean accumulator means your sums start from a known value.

The core call is mma_sync. It runs D = A times B plus C on the loaded fragments using the Tensor Cores directly.

Finally store_matrix_sync writes the accumulator fragment back to memory. This store gathers each thread's pieces into one tile.

The _sync suffix is a reminder: every WMMA call is warp-synchronous. All 32 lanes must reach it together or behavior is undefined.

Loads and stores need the leading dimension, the stride between rows, plus a row- or column-major layout to read memory correctly.

Which WMMA call actually runs the matrix multiply-accumulate on the Tensor Cores?

You walked the WMMA flow: declare a fragment, load tiles, clear the accumulator, call mma_sync, then store. Every step is warp-synchronous. 🙌