FP16 Is Not Enough: Activate Tensor Cores

The starter already stores A and B as half precision, but each lane performs ordinary scalar multiply-add work, so the Tensor pipe remains idle. Make the full warp cooperate on 16 by 16 WMMA fragments and issue mma_sync while keeping float accumulation and the same matrix result.

Languages: CUDA C++, Rust

Lesson

Half precision is a storage and arithmetic type; it does not by itself select Tensor Cores. WMMA expresses a warp-wide matrix multiply-accumulate. All 32 lanes participate in fragment loads, mma_sync, and the fragment store. The harness uses aligned allocations and dimensions that are multiples of 16 so the lesson focuses on mapping work to the Tensor pipe rather than padding edge tiles.

The starter uses FP32 scalar arithmetic; the WMMA path converts inputs to half and accumulates in float. A WMMA API operation may lower to multiple hardware instructions depending on architecture/compiler.

What you will learn

Distinguish half-precision scalar arithmetic from Tensor Core matrix instructions.
Map one warp to one 16 by 16 output tile.
Load row-major A and B fragments and accumulate into float.
Keep every lane converged across load_matrix_sync, mma_sync, and store_matrix_sync.
Interpret GPUFlight Tensor Core utilization, where zero means the Tensor pipe issued no work.

Starter code

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

    // The data is half precision, but this is still scalar CUDA-core work.
    for (int element = lane; element < 16 * 16; element += 32) {
        int row = tileRow + element / 16;
        int col = tileCol + element % 16;
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) {
            sum += __half2float(A[row * N + k])
                 * __half2float(B[k * N + col]);
        }
        C[row * N + col] = sum;
    }