The WMMA Fragment API
Load, mma_sync, and store fragments.
The WMMA Fragment API is a free CUDA Academy lesson on CoddyKit — lesson 3 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
The WMMA API
To program Tensor Cores directly you use WMMA, the warp matrix multiply-accumulate API in the nvcuda::wmma namespace. 🧩
A Whole Warp Cooperates
WMMA is warp-wide: all 32 threads in a warp work together on one tile. You think in tiles, not in single threads.
Meet the Fragment
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.
Three Fragment Roles
You declare fragments tagged matrix_a, matrix_b, or accumulator. The tag tells WMMA how that tile feeds the multiply-accumulate.
Tile Shapes Are Fixed
Fragments use set tile sizes like 16 by 16 by 16. You pick a supported shape; the hardware only accepts those combinations.
Step 1: Load
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.
wmma::load_matrix_sync(a_frag, ptr, ldm);Step 2: Clear the Accumulator
Before adding, you zero the result tile with fill_fragment. A clean accumulator means your sums start from a known value.
wmma::fill_fragment(c_frag, 0.0f);Step 3: Multiply-Accumulate
The core call is mma_sync. It runs D = A times B plus C on the loaded fragments using the Tensor Cores directly.
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);Step 4: Store
Finally store_matrix_sync writes the accumulator fragment back to memory. This store gathers each thread's pieces into one tile.
wmma::store_matrix_sync(out, c_frag, ldm, layout);Sync Means Warp-Synchronous
The _sync suffix is a reminder: every WMMA call is warp-synchronous. All 32 lanes must reach it together or behavior is undefined.
Layout and Leading Dimension
Loads and stores need the leading dimension, the stride between rows, plus a row- or column-major layout to read memory correctly.
Quick Check
Which WMMA call actually runs the matrix multiply-accumulate on the Tensor Cores?
Recap
You walked the WMMA flow: declare a fragment, load tiles, clear the accumulator, call mma_sync, then store. Every step is warp-synchronous. 🙌
Frequently asked questions
Is the “The WMMA Fragment API” lesson free?
Yes — the full text of “The WMMA Fragment API” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “The WMMA Fragment API”?
Load, mma_sync, and store fragments. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 3 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “The WMMA Fragment API” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.