Tensor Cores: Keep the Accumulator in the Fragment

The starter issues real WMMA instructions, but after every 16-element K slice it stores the partial C tile to global memory and loads it again for the next slice. Keep the accumulator fragment live across the entire K loop and store the completed tile once.

Languages: CUDA C++, Rust

Lesson

Issuing Tensor instructions is only the first step. Tensor-pipe utilization falls when every mma_sync is separated by unnecessary global-memory traffic. A float accumulator fragment can remain live across the K loop, so partial sums do not need to leave registers until the complete output tile is ready.

The forced global accumulator round trip is a warm-up. Keeping fragments live can increase register pressure; inspect occupancy and timing alongside fewer sectors.

What you will learn

Recognize global accumulator spill and reload traffic between WMMA K slices.
Initialize the accumulator fragment once before the K loop.
Reuse the same accumulator fragment across every mma_sync operation.
Store one completed 16 by 16 tile after the K loop.
Use Tensor utilization together with DRAM traffic rather than treating Tensor activity alone as proof of efficiency.

Starter code

    using namespace nvcuda;

    int tileRow = static_cast<int>(blockIdx.y) * 16;
    int tileCol = static_cast<int>(blockIdx.x) * 16;

    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> aFrag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::row_major> bFrag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> cFrag;

    for (int k = 0; k < N; k += 16) {
        if (k == 0) {
            wmma::fill_fragment(cFrag, 0.0f);
        } else {
            wmma::load_matrix_sync(cFrag, C + tileRow * N + tileCol, N, wmma::mem_row_major);
        }
        wmma::load_matrix_sync(aFrag, A + tileRow * N + k, N);
        wmma::load_matrix_sync(bFrag, B + k * N + tileCol, N);
        wmma::mma_sync(cFrag, aFrag, bFrag, cFrag);
        wmma::store_matrix_sync(C + tileRow * N + tileCol, cFrag, N, wmma::mem_row_major);
        __syncwarp();
    }